From d84915b10c97cdf14403406a64453001488dc6ce Mon Sep 17 00:00:00 2001 From: Witold Wojciechowski Date: Thu, 14 May 2026 21:42:39 +0200 Subject: [PATCH 1/7] Added ContextBase.h and refactored GpuContext.h, as child. --- hiprt/hiprt_cpu.h | 28 ++++ hiprt/hiprt_types.h | 11 +- hiprt/impl/BatchBuilder.cpp | 4 +- hiprt/impl/BatchBuilder.h | 12 +- hiprt/impl/BvhImporter.cpp | 16 +-- hiprt/impl/BvhImporter.h | 30 ++-- hiprt/impl/Compiler.cpp | 14 +- hiprt/impl/Compiler.h | 14 +- hiprt/impl/ContextBase.h | 96 +++++++++++++ hiprt/impl/CpuContext.cpp | 154 +++++++++++++++++++++ hiprt/impl/CpuContext.h | 115 +++++++++++++++ hiprt/impl/{Context.cpp => GpuContext.cpp} | 82 +++++------ hiprt/impl/{Context.h => GpuContext.h} | 62 +++++---- hiprt/impl/LbvhBuilder.cpp | 16 +-- hiprt/impl/LbvhBuilder.h | 26 ++-- hiprt/impl/PlocBuilder.cpp | 16 +-- hiprt/impl/PlocBuilder.h | 26 ++-- hiprt/impl/SbvhBuilder.cpp | 16 +-- hiprt/impl/SbvhBuilder.h | 30 ++-- hiprt/impl/Utility.h | 2 +- hiprt/impl/hiprt.cpp | 110 +++++++-------- 21 files changed, 643 insertions(+), 237 deletions(-) create mode 100644 hiprt/hiprt_cpu.h create mode 100644 hiprt/impl/ContextBase.h create mode 100644 hiprt/impl/CpuContext.cpp create mode 100644 hiprt/impl/CpuContext.h rename hiprt/impl/{Context.cpp => GpuContext.cpp} (92%) rename hiprt/impl/{Context.h => GpuContext.h} (80%) diff --git a/hiprt/hiprt_cpu.h b/hiprt/hiprt_cpu.h new file mode 100644 index 0000000..f6bbfa6 --- /dev/null +++ b/hiprt/hiprt_cpu.h @@ -0,0 +1,28 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + + + diff --git a/hiprt/hiprt_types.h b/hiprt/hiprt_types.h index 2e631c3..82f97a5 100644 --- a/hiprt/hiprt_types.h +++ b/hiprt/hiprt_types.h @@ -309,12 +309,11 @@ struct hiprtFuncNameSet /** \brief Device type. * */ -enum hiprtDeviceType +enum hiprtDeviceType : uint32_t { - /*!< AMD device */ - hiprtDeviceAMD, - /*!< Nvidia device */ - hiprtDeviceNVIDIA, + hiprtDeviceAMD = 1 << 0, + hiprtDeviceNVIDIA = 1 << 1, + hiprtDeviceCPU = 1 << 2, }; /** \brief Context creation input. @@ -328,6 +327,8 @@ struct hiprtContextCreationInput hiprtApiDevice device; /*!< HIPRT API device type */ hiprtDeviceType deviceType; + /*!< HIPRT API CPU Device type, number of threads */ + uint32_t numCpuThreads = 0; }; /** \brief Various flags controlling scene/geometry build process. diff --git a/hiprt/impl/BatchBuilder.cpp b/hiprt/impl/BatchBuilder.cpp index 528e8eb..35554f8 100644 --- a/hiprt/impl/BatchBuilder.cpp +++ b/hiprt/impl/BatchBuilder.cpp @@ -33,7 +33,7 @@ DECLARE_TYPE_TRAITS( hiprtSceneBuildInput ); #endif size_t BatchBuilder::getStorageBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); const size_t boxNodeCount = getMaxBoxNodeCount( buildInput, context.getRtip(), primCount ); @@ -42,7 +42,7 @@ size_t BatchBuilder::getStorageBufferSize( } size_t BatchBuilder::getStorageBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = buildInput.instanceCount; const size_t boxNodeCount = getMaxBoxNodeCount( buildInput, context.getRtip(), primCount ); diff --git a/hiprt/impl/BatchBuilder.h b/hiprt/impl/BatchBuilder.h index 25e3fbf..e624808 100644 --- a/hiprt/impl/BatchBuilder.h +++ b/hiprt/impl/BatchBuilder.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include #include #include #include @@ -56,21 +56,21 @@ class BatchBuilder template static size_t getTemporaryBufferSize( - [[maybe_unused]] Context& context, const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) + [[maybe_unused]] GpuContext& context, const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) { return RoundUp( sizeof( BuildInput ) * buildInputs.size(), DefaultAlignment ) + RoundUp( sizeof( hiprtDevicePtr ) * buildInputs.size(), DefaultAlignment ); } static size_t - getStorageBufferSize( Context& context, const hiprtGeometryBuildInput& buildInputs, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInputs, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtSceneBuildInput& buildInputs, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInputs, const hiprtBuildOptions buildOptions ); template static void build( - Context& context, + GpuContext& context, const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -80,7 +80,7 @@ class BatchBuilder template void BatchBuilder::build( - Context& context, + GpuContext& context, const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, diff --git a/hiprt/impl/BvhImporter.cpp b/hiprt/impl/BvhImporter.cpp index d244208..7650038 100644 --- a/hiprt/impl/BvhImporter.cpp +++ b/hiprt/impl/BvhImporter.cpp @@ -32,19 +32,19 @@ namespace hiprt { size_t BvhImporter::getTemporaryBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { return getTemporaryBufferSize( context, buildInput, buildOptions ); } size_t BvhImporter::getTemporaryBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { return getTemporaryBufferSize( context, buildInput, buildOptions ); } size_t BvhImporter::getStorageBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); const size_t maxReferenceCount = buildInput.nodeList.nodeCount; @@ -55,7 +55,7 @@ size_t BvhImporter::getStorageBufferSize( } size_t BvhImporter::getStorageBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = buildInput.instanceCount; const size_t maxReferenceCount = buildInput.nodeList.nodeCount; @@ -70,7 +70,7 @@ size_t BvhImporter::getStorageBufferSize( } void BvhImporter::build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -112,7 +112,7 @@ void BvhImporter::build( } void BvhImporter::build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -154,7 +154,7 @@ void BvhImporter::build( } void BvhImporter::update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, @@ -189,7 +189,7 @@ void BvhImporter::update( } } void BvhImporter::update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, diff --git a/hiprt/impl/BvhImporter.h b/hiprt/impl/BvhImporter.h index 1e2e5be..1e1ec69 100644 --- a/hiprt/impl/BvhImporter.h +++ b/hiprt/impl/BvhImporter.h @@ -25,7 +25,7 @@ #pragma once #include #include -#include +#include #include #include #include @@ -49,22 +49,22 @@ class BvhImporter template static size_t - getTemporaryBufferSize( Context& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static void build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -72,7 +72,7 @@ class BvhImporter hiprtDevicePtr buffer ); static void build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -81,7 +81,7 @@ class BvhImporter template static void build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, NodeList& nodes, const hiprtBuildOptions buildOptions, @@ -91,7 +91,7 @@ class BvhImporter MemoryArena& storageMemoryArena ); static void update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -99,7 +99,7 @@ class BvhImporter hiprtDevicePtr buffer ); static void update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -108,7 +108,7 @@ class BvhImporter template static void update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const NodeList& nodes, const hiprtBuildOptions buildOptions, @@ -118,7 +118,7 @@ class BvhImporter template size_t -BvhImporter::getTemporaryBufferSize( Context& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ) +BvhImporter::getTemporaryBufferSize( GpuContext& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ) { const bool kdops = context.getRtip() >= 31 && !( buildOptions.buildFlags & hiprtBuildFlagBitDisableOrientedBoundingBoxes ); const size_t count = buildInput.nodeList.nodeCount; @@ -133,7 +133,7 @@ BvhImporter::getTemporaryBufferSize( Context& context, const BuildInput& buildIn template void BvhImporter::build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, NodeList& nodes, const hiprtBuildOptions buildOptions, @@ -360,7 +360,7 @@ void BvhImporter::build( template void BvhImporter::update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const NodeList& nodes, const hiprtBuildOptions buildOptions, diff --git a/hiprt/impl/Compiler.cpp b/hiprt/impl/Compiler.cpp index 7fb18e9..d6e50a4 100644 --- a/hiprt/impl/Compiler.cpp +++ b/hiprt/impl/Compiler.cpp @@ -26,7 +26,7 @@ #include #include #include -#include +#include #include #if defined( HIPRT_BAKE_KERNEL_GENERATED ) #include @@ -110,7 +110,7 @@ Compiler::~Compiler() } Kernel Compiler::getKernel( - Context& context, + GpuContext& context, const std::filesystem::path& moduleName, const std::string& funcName, std::vector& options, @@ -230,7 +230,7 @@ void Compiler::buildProgram( } void Compiler::buildKernels( - Context& context, + GpuContext& context, const std::vector& funcNames, const std::string& src, const std::filesystem::path& moduleName, @@ -336,7 +336,7 @@ void Compiler::buildKernels( } void Compiler::buildKernelsFromBitcode( - Context& context, + GpuContext& context, const std::vector& funcNames, const std::filesystem::path& moduleName, const std::string_view bitcodeBinary, @@ -511,7 +511,7 @@ Compiler::readSourceCode( const std::filesystem::path& path, std::optional& opts, bool extended ) +void Compiler::addCommonOpts( GpuContext& context, std::vector& opts, bool extended ) { if ( !extended ) // do not use fast math for trace kernel (the fast math option can be passed from the application) { @@ -625,7 +625,7 @@ void Compiler::addCustomFuncsSwitchCase( } std::string Compiler::getCacheFilename( - Context& context, + GpuContext& context, const std::string& src, const std::filesystem::path& moduleName, std::optional> options, @@ -813,7 +813,7 @@ oroFunction Compiler::getFunctionFromPrecompiledBinary( const std::string& funcN } std::string Compiler::buildFunctionTableBitcode( - Context& context, uint32_t numGeomTypes, uint32_t numRayTypes, const std::vector& funcNameSets ) + GpuContext& context, uint32_t numGeomTypes, uint32_t numRayTypes, const std::vector& funcNameSets ) { if constexpr ( BakedCodeIsGenerated ) { diff --git a/hiprt/impl/Compiler.h b/hiprt/impl/Compiler.h index f9f2587..b49dc0c 100644 --- a/hiprt/impl/Compiler.h +++ b/hiprt/impl/Compiler.h @@ -31,7 +31,7 @@ namespace hiprt { -class Context; +class GpuContext; class Compiler { public: @@ -40,7 +40,7 @@ class Compiler void init(); Kernel getKernel( - Context& context, + GpuContext& context, const std::filesystem::path& moduleName, const std::string& funcName, std::vector& options, @@ -58,7 +58,7 @@ class Compiler orortcProgram& progOut ); void buildKernels( - Context& context, + GpuContext& context, const std::vector& funcNames, const std::string& src, const std::filesystem::path& moduleName, @@ -74,7 +74,7 @@ class Compiler bool cache ); void buildKernelsFromBitcode( - Context& context, + GpuContext& context, const std::vector& funcNames, const std::filesystem::path& moduleName, const std::string_view bitcodeBinary, @@ -105,10 +105,10 @@ class Compiler uint32_t numGeomTypes = 0, uint32_t numRayTypes = 1 ); - void addCommonOpts( Context& context, std::vector& opts, bool extended ); + void addCommonOpts( GpuContext& context, std::vector& opts, bool extended ); std::string getCacheFilename( - Context& context, + GpuContext& context, const std::string& src, const std::filesystem::path& moduleName, std::optional> options = std::nullopt, @@ -123,7 +123,7 @@ class Compiler oroFunction getFunctionFromPrecompiledBinary( const std::string& funcName ); std::string buildFunctionTableBitcode( - Context& context, uint32_t numGeomTypes, uint32_t numRayTypes, const std::vector& funcNameSets ); + GpuContext& context, uint32_t numGeomTypes, uint32_t numRayTypes, const std::vector& funcNameSets ); std::filesystem::path m_cacheDirectory = "cache"; diff --git a/hiprt/impl/ContextBase.h b/hiprt/impl/ContextBase.h new file mode 100644 index 0000000..5a74d7d --- /dev/null +++ b/hiprt/impl/ContextBase.h @@ -0,0 +1,96 @@ + +#pragma once + +#include +#include +#include +#include +#include + +#include +#include + +namespace hiprt +{ + + +class ContextBase +{ + public: + virtual ~ContextBase() = default; + + virtual std::vector + createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) = 0; + virtual void buildGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) = 0; + virtual void updateGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) = 0; + virtual void destroyGeometries( const std::vector& geometries ) = 0; + virtual size_t getGeometriesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) = 0; + virtual std::vector compactGeometries( const std::vector& geometries, oroStream stream ) = 0; + + virtual std::vector + createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) = 0; + virtual void buildScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) = 0; + virtual void updateScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) = 0; + virtual void destroyScenes( const std::vector& scenes ) = 0; + virtual size_t getScenesBuildTempBufferSize( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) = 0; + virtual std::vector compactScenes( const std::vector& scenes, oroStream stream ) = 0; + + virtual void saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) = 0; + virtual hiprtGeometry loadGeometry( const std::string& filename ) = 0; + virtual void saveScene( hiprtScene inScene, const std::string& filename ) = 0; + virtual hiprtScene loadScene( const std::string& filename ) = 0; + virtual void exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) = 0; + virtual void exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) = 0; + + virtual void buildKernels( + const std::vector& funcNames, + const std::string& src, + const std::filesystem::path& moduleName, + std::vector& headers, + std::vector& includeNames, + std::vector& options, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + oroModule& module, + bool cache ) = 0; + virtual void buildKernelsFromBitcode( + const std::vector& funcNames, + const std::filesystem::path& moduleName, + const std::string_view bitcodeBinary, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + bool cache ) = 0; + virtual void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) = 0; + virtual void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) = 0; + + virtual void setLogLevel( hiprtLogLevel level ) = 0; + + virtual struct CpuGeometryData* getCpuGeom( hiprtGeometry ) = 0; + virtual struct CpuSceneData* getCpuScene( hiprtScene ) = 0; +}; +} // namespace hiprt \ No newline at end of file diff --git a/hiprt/impl/CpuContext.cpp b/hiprt/impl/CpuContext.cpp new file mode 100644 index 0000000..6064cda --- /dev/null +++ b/hiprt/impl/CpuContext.cpp @@ -0,0 +1,154 @@ +#include "CpuContext.h" + +#include + +namespace hiprt +{ +namespace +{ +[[noreturn]] void throwNotImplemented( const char* what ) +{ + throw std::runtime_error( what ); +} +} // namespace + +std::vector CpuContext::createGeometries( + const std::vector&, const hiprtBuildOptions ) +{ + throwNotImplemented( "CpuContext::createGeometries is not implemented yet." ); +} + +void CpuContext::destroyGeometries( const std::vector& ) { throwNotImplemented( "CpuContext::destroyGeometries is not implemented yet." ); } + +void CpuContext::buildGeometries( + const std::vector&, + const hiprtBuildOptions, + hiprtDevicePtr, + oroStream, + std::vector& ) +{ + throwNotImplemented( "CpuContext::buildGeometries is not implemented yet." ); +} + +void CpuContext::updateGeometries( + const std::vector&, + const hiprtBuildOptions, + hiprtDevicePtr, + oroStream, + std::vector& ) +{ + throwNotImplemented( "CpuContext::updateGeometries is not implemented yet." ); +} + +size_t CpuContext::getGeometriesBuildTempBufferSize( const std::vector&, const hiprtBuildOptions ) +{ + throwNotImplemented( "CpuContext::getGeometriesBuildTempBufferSize is not implemented yet." ); +} + +std::vector CpuContext::compactGeometries( const std::vector&, oroStream ) +{ + throwNotImplemented( "CpuContext::compactGeometries is not implemented yet." ); +} + +std::vector CpuContext::createScenes( const std::vector&, const hiprtBuildOptions ) +{ + throwNotImplemented( "CpuContext::createScenes is not implemented yet." ); +} + +void CpuContext::destroyScenes( const std::vector& ) { throwNotImplemented( "CpuContext::destroyScenes is not implemented yet." ); } + +void CpuContext::buildScenes( + const std::vector&, + const hiprtBuildOptions, + hiprtDevicePtr, + oroStream, + std::vector& ) +{ + throwNotImplemented( "CpuContext::buildScenes is not implemented yet." ); +} + +void CpuContext::updateScenes( + const std::vector&, + const hiprtBuildOptions, + hiprtDevicePtr, + oroStream, + std::vector& ) +{ + throwNotImplemented( "CpuContext::updateScenes is not implemented yet." ); +} + +size_t CpuContext::getScenesBuildTempBufferSize( const std::vector&, const hiprtBuildOptions ) +{ + throwNotImplemented( "CpuContext::getScenesBuildTempBufferSize is not implemented yet." ); +} + +std::vector CpuContext::compactScenes( const std::vector&, oroStream ) +{ + throwNotImplemented( "CpuContext::compactScenes is not implemented yet." ); +} + +hiprtFuncTable CpuContext::createFuncTable( uint32_t, uint32_t ) +{ + throwNotImplemented( "CpuContext::createFuncTable is not implemented yet." ); +} + +void CpuContext::setFuncTable( hiprtFuncTable, uint32_t, uint32_t, hiprtFuncDataSet ) +{ + throwNotImplemented( "CpuContext::setFuncTable is not implemented yet." ); +} + +void CpuContext::destroyFuncTable( hiprtFuncTable ) { throwNotImplemented( "CpuContext::destroyFuncTable is not implemented yet." ); } + +void CpuContext::createGlobalStackBuffer( const hiprtGlobalStackBufferInput&, hiprtGlobalStackBuffer& ) +{ + throwNotImplemented( "CpuContext::createGlobalStackBuffer is not implemented yet." ); +} + +void CpuContext::destroyGlobalStackBuffer( hiprtGlobalStackBuffer ) { throwNotImplemented( "CpuContext::destroyGlobalStackBuffer is not implemented yet." ); } + +void CpuContext::saveGeometry( hiprtGeometry, const std::string& ) { throwNotImplemented( "CpuContext::saveGeometry is not implemented yet." ); } + +hiprtGeometry CpuContext::loadGeometry( const std::string& ) { throwNotImplemented( "CpuContext::loadGeometry is not implemented yet." ); } + +void CpuContext::saveScene( hiprtScene, const std::string& ) { throwNotImplemented( "CpuContext::saveScene is not implemented yet." ); } + +hiprtScene CpuContext::loadScene( const std::string& ) { throwNotImplemented( "CpuContext::loadScene is not implemented yet." ); } + +void CpuContext::exportGeometryAabb( hiprtGeometry, float3&, float3& ) +{ + throwNotImplemented( "CpuContext::exportGeometryAabb is not implemented yet." ); +} + +void CpuContext::exportSceneAabb( hiprtScene, float3&, float3& ) { throwNotImplemented( "CpuContext::exportSceneAabb is not implemented yet." ); } + +void CpuContext::buildKernels( + const std::vector&, + const std::string&, + const std::filesystem::path&, + std::vector&, + std::vector&, + std::vector&, + uint32_t, + uint32_t, + const std::vector&, + std::vector&, + oroModule&, + bool ) +{ + throwNotImplemented( "CpuContext::buildKernels is not implemented yet." ); +} + +void CpuContext::buildKernelsFromBitcode( + const std::vector&, + const std::filesystem::path&, + const std::string_view, + uint32_t, + uint32_t, + const std::vector&, + std::vector&, + bool ) +{ + throwNotImplemented( "CpuContext::buildKernelsFromBitcode is not implemented yet." ); +} + +} // namespace hiprt diff --git a/hiprt/impl/CpuContext.h b/hiprt/impl/CpuContext.h new file mode 100644 index 0000000..d3a0361 --- /dev/null +++ b/hiprt/impl/CpuContext.h @@ -0,0 +1,115 @@ +#pragma once +#include "ContextBase.h" + +#include +#include +#include +#include +#include +#include + + +namespace hiprt{ +class CpuContext : public ContextBase +{ + public: + std::vector + createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + void destroyGeometries( const std::vector& geometries ) override; + + void buildGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + void updateGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + size_t getGeometriesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + std::vector compactGeometries( const std::vector& geometries, oroStream stream ) override; + + std::vector + createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + void destroyScenes( const std::vector& scenes ) override; + + void buildScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + void updateScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + size_t getScenesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; + + hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ); + void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ); + void destroyFuncTable( hiprtFuncTable funcTable ); + + void setLogLevel( hiprtLogLevel level ) override { (void)level; } + + void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) override; + void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) override; + + void saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) override; + hiprtGeometry loadGeometry( const std::string& filename ) override; + + void saveScene( hiprtScene inScene, const std::string& filename ) override; + hiprtScene loadScene( const std::string& filename ) override; + + void exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) override; + void exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) override; + + void buildKernels( + const std::vector& funcNames, + const std::string& src, + const std::filesystem::path& moduleName, + std::vector& headers, + std::vector& includeNames, + std::vector& options, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + oroModule& module, + bool cache ) override; + + void buildKernelsFromBitcode( + const std::vector& funcNames, + const std::filesystem::path& moduleName, + const std::string_view bitcodeBinary, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + bool cache ) override; + + struct CpuGeometryData* getCpuGeom( hiprtGeometry ) override + { + return nullptr; + }; + struct CpuSceneData* getCpuScene( hiprtScene ) override + { + return nullptr; + }; +}; +}//namespace hiprt \ No newline at end of file diff --git a/hiprt/impl/Context.cpp b/hiprt/impl/GpuContext.cpp similarity index 92% rename from hiprt/impl/Context.cpp rename to hiprt/impl/GpuContext.cpp index 094cae8..2831861 100644 --- a/hiprt/impl/Context.cpp +++ b/hiprt/impl/GpuContext.cpp @@ -25,7 +25,7 @@ #include #include #include -#include +#include #include #include #include @@ -34,7 +34,7 @@ namespace hiprt { -Context::Context( const hiprtContextCreationInput& input ) +GpuContext::GpuContext( const hiprtContextCreationInput& input ) { oroApi api = ( input.deviceType == hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA; oroCtxCreateFromRaw( &m_ctxt, api, input.ctxt ); @@ -43,14 +43,14 @@ Context::Context( const hiprtContextCreationInput& input ) m_compiler.init(); } -Context::~Context() +GpuContext::~GpuContext() { m_oroutils.unloadKernelCache(); oroCtxCreateFromRawDestroy( m_ctxt ); } std::vector -Context::createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +GpuContext::createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -112,7 +112,7 @@ Context::createGeometries( const std::vector& buildInpu return geometries; } -void Context::destroyGeometries( const std::vector& geometries ) +void GpuContext::destroyGeometries( const std::vector& geometries ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -141,7 +141,7 @@ void Context::destroyGeometries( const std::vector& geometries ) } } -void Context::buildGeometries( +void GpuContext::buildGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -200,7 +200,7 @@ void Context::buildGeometries( } } -void Context::updateGeometries( +void GpuContext::updateGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -238,7 +238,7 @@ void Context::updateGeometries( } } -size_t Context::getGeometriesBuildTempBufferSize( +size_t GpuContext::getGeometriesBuildTempBufferSize( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) { std::vector batchInputs; @@ -289,7 +289,7 @@ size_t Context::getGeometriesBuildTempBufferSize( return size; } -std::vector Context::compactGeometries( const std::vector& geometriesIn, oroStream stream ) +std::vector GpuContext::compactGeometries( const std::vector& geometriesIn, oroStream stream ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -355,7 +355,7 @@ std::vector Context::compactGeometries( const std::vector -Context::createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +GpuContext::createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -426,7 +426,7 @@ Context::createScenes( const std::vector& buildInputs, con return scenes; } -void Context::destroyScenes( const std::vector& scenes ) +void GpuContext::destroyScenes( const std::vector& scenes ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -455,7 +455,7 @@ void Context::destroyScenes( const std::vector& scenes ) } } -void Context::buildScenes( +void GpuContext::buildScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -523,7 +523,7 @@ void Context::buildScenes( } } -void Context::updateScenes( +void GpuContext::updateScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -570,7 +570,7 @@ void Context::updateScenes( } } -size_t Context::getScenesBuildTempBufferSize( +size_t GpuContext::getScenesBuildTempBufferSize( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) { std::vector batchInputs; @@ -631,7 +631,7 @@ size_t Context::getScenesBuildTempBufferSize( return size; } -std::vector Context::compactScenes( const std::vector& scenesIn, oroStream stream ) +std::vector GpuContext::compactScenes( const std::vector& scenesIn, oroStream stream ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -711,7 +711,7 @@ std::vector Context::compactScenes( const std::vector& s return scenesOut; } -hiprtFuncTable Context::createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) +hiprtFuncTable GpuContext::createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -731,7 +731,7 @@ hiprtFuncTable Context::createFuncTable( uint32_t numGeomTypes, uint32_t numRayT return reinterpret_cast( ptr ); } -void Context::setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) +void GpuContext::setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -743,13 +743,13 @@ void Context::setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_ oroMemcpyHtoD( reinterpret_cast( &header.funcDataSets[index] ), &set, sizeof( hiprtFuncDataSet ) ) ); } -void Context::destroyFuncTable( hiprtFuncTable funcTable ) +void GpuContext::destroyFuncTable( hiprtFuncTable funcTable ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); checkOro( oroFree( reinterpret_cast( funcTable ) ) ); } -void Context::createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) +void GpuContext::createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -778,13 +778,13 @@ void Context::createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, } } -void Context::destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) +void GpuContext::destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); checkOro( oroFree( reinterpret_cast( stackBuffer.stackData ) ) ); } -void Context::saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) +void GpuContext::saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -809,7 +809,7 @@ void Context::saveGeometry( hiprtGeometry inGeometry, const std::string& filenam file.write( reinterpret_cast( buffer.data() ), header.m_size ); } -hiprtGeometry Context::loadGeometry( const std::string& filename ) +hiprtGeometry GpuContext::loadGeometry( const std::string& filename ) { std::ifstream file( filename, std::ios::in | std::ios::binary ); @@ -852,14 +852,14 @@ hiprtGeometry Context::loadGeometry( const std::string& filename ) return geometry; } -void Context::saveScene( [[maybe_unused]] hiprtScene inScene, [[maybe_unused]] const std::string& filename ) +void GpuContext::saveScene( [[maybe_unused]] hiprtScene inScene, [[maybe_unused]] const std::string& filename ) { throw std::runtime_error( "Not implemented" ); } -hiprtScene Context::loadScene( [[maybe_unused]] const std::string& filename ) { throw std::runtime_error( "Not implemented" ); } +hiprtScene GpuContext::loadScene( [[maybe_unused]] const std::string& filename ) { throw std::runtime_error( "Not implemented" ); } -void Context::exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) +void GpuContext::exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -876,7 +876,7 @@ void Context::exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, outAabbMax = box.m_max; } -void Context::exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) +void GpuContext::exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) { checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -893,7 +893,7 @@ void Context::exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& o outAabbMax = box.m_max; } -void Context::buildKernels( +void GpuContext::buildKernels( const std::vector& funcNames, const std::string& src, const std::filesystem::path& moduleName, @@ -925,7 +925,7 @@ void Context::buildKernels( cache ); } -void Context::buildKernelsFromBitcode( +void GpuContext::buildKernelsFromBitcode( const std::vector& funcNames, const std::filesystem::path& moduleName, const std::string_view bitcodeBinary, @@ -940,9 +940,9 @@ void Context::buildKernelsFromBitcode( *this, funcNames, moduleName, bitcodeBinary, numGeomTypes, numRayTypes, funcNameSets, functions, cache ); } -void Context::setCacheDir( const std::filesystem::path& path ) { m_compiler.setCacheDir( path ); } +void GpuContext::setCacheDir( const std::filesystem::path& path ) { m_compiler.setCacheDir( path ); } -uint32_t Context::getSMCount() const +uint32_t GpuContext::getSMCount() const { int smCount; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -950,7 +950,7 @@ uint32_t Context::getSMCount() const return smCount; } -uint32_t Context::getMaxBlockSize() const +uint32_t GpuContext::getMaxBlockSize() const { oroDeviceProp prop; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -958,7 +958,7 @@ uint32_t Context::getMaxBlockSize() const return prop.maxThreadsPerBlock; } -uint32_t Context::getMaxGridSize() const +uint32_t GpuContext::getMaxGridSize() const { oroDeviceProp prop; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -966,7 +966,7 @@ uint32_t Context::getMaxGridSize() const return prop.maxGridSize[0]; } -std::string Context::getDeviceName() const +std::string GpuContext::getDeviceName() const { oroDeviceProp prop; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -974,7 +974,7 @@ std::string Context::getDeviceName() const return std::string( prop.name ); } -std::string Context::getGcnArchName() const +std::string GpuContext::getGcnArchName() const { oroDeviceProp prop; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -982,7 +982,7 @@ std::string Context::getGcnArchName() const return std::string( prop.gcnArchName ); } -std::string Context::getDriverVersion() const +std::string GpuContext::getDriverVersion() const { int driverVersion; checkOro( oroCtxSetCurrent( m_ctxt ) ); @@ -990,7 +990,7 @@ std::string Context::getDriverVersion() const return std::to_string( driverVersion ); } -uint32_t Context::getRtip() const +uint32_t GpuContext::getRtip() const { std::string deviceName = getDeviceName(); std::string archName = getGcnArchName(); @@ -1043,13 +1043,13 @@ uint32_t Context::getRtip() const return rtip; } -uint32_t Context::getBranchingFactor() const +uint32_t GpuContext::getBranchingFactor() const { if ( getRtip() >= 31 ) return 8; return 4; } -uint32_t Context::getWarpSize() const +uint32_t GpuContext::getWarpSize() const { std::string deviceName = getDeviceName(); std::string archName = getGcnArchName(); @@ -1070,19 +1070,19 @@ uint32_t Context::getWarpSize() const return warpSize; } -size_t Context::getTriangleNodeSize() const +size_t GpuContext::getTriangleNodeSize() const { if ( getRtip() >= 31 ) return sizeof( TrianglePacketNode ); return sizeof( TrianglePairNode ); } -size_t Context::getBoxNodeSize() const +size_t GpuContext::getBoxNodeSize() const { if ( getRtip() >= 31 ) return sizeof( Box8Node ); return sizeof( Box4Node ); } -size_t Context::getInstanceNodeSize() const +size_t GpuContext::getInstanceNodeSize() const { if ( getRtip() >= 31 ) return sizeof( HwInstanceNode ); return sizeof( UserInstanceNode ); diff --git a/hiprt/impl/Context.h b/hiprt/impl/GpuContext.h similarity index 80% rename from hiprt/impl/Context.h rename to hiprt/impl/GpuContext.h index c078dbd..5970918 100644 --- a/hiprt/impl/Context.h +++ b/hiprt/impl/GpuContext.h @@ -23,6 +23,8 @@ ////////////////////////////////////////////////////////////////////////////////////////// #pragma once +#include "ContextBase.h" + #include #include #include @@ -32,75 +34,75 @@ namespace hiprt { -class Context +class GpuContext : public ContextBase { public: - Context( const hiprtContextCreationInput& input ); - ~Context(); + GpuContext( const hiprtContextCreationInput& input ); + ~GpuContext() override; std::vector - createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ); + createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; - void destroyGeometries( const std::vector& geometries ); + void destroyGeometries( const std::vector& geometries ) override; void buildGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, oroStream stream, - std::vector& buffers ); + std::vector& buffers ) override; void updateGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, oroStream stream, - std::vector& buffers ); + std::vector& buffers ) override; size_t getGeometriesBuildTempBufferSize( - const std::vector& buildInputs, const hiprtBuildOptions buildOptions ); + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; - std::vector compactGeometries( const std::vector& geometries, oroStream stream ); + std::vector compactGeometries( const std::vector& geometries, oroStream stream ) override; std::vector - createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ); + createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; - void destroyScenes( const std::vector& scenes ); + void destroyScenes( const std::vector& scenes ) override; void buildScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, oroStream stream, - std::vector& buffers ); + std::vector& buffers ) override; void updateScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, oroStream stream, - std::vector& buffers ); + std::vector& buffers ) override; size_t getScenesBuildTempBufferSize( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ); - std::vector compactScenes( const std::vector& scenes, oroStream stream ); + std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ); void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ); void destroyFuncTable( hiprtFuncTable funcTable ); - void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ); - void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ); + void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) override; + void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) override; - void saveGeometry( hiprtGeometry inGeometry, const std::string& filename ); - hiprtGeometry loadGeometry( const std::string& filename ); + void saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) override; + hiprtGeometry loadGeometry( const std::string& filename ) override; - void saveScene( hiprtScene inScene, const std::string& filename ); - hiprtScene loadScene( const std::string& filename ); + void saveScene( hiprtScene inScene, const std::string& filename ) override; + hiprtScene loadScene( const std::string& filename ) override; - void exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ); - void exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ); + void exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) override; + void exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) override; void buildKernels( const std::vector& funcNames, @@ -114,7 +116,7 @@ class Context const std::vector& funcNameSets, std::vector& functions, oroModule& module, - bool cache ); + bool cache ) override; void buildKernelsFromBitcode( const std::vector& funcNames, @@ -124,11 +126,20 @@ class Context uint32_t numRayTypes, const std::vector& funcNameSets, std::vector& functions, - bool cache ); + bool cache ) override; void setCacheDir( const std::filesystem::path& path ); - void setLogLevel( hiprtLogLevel level ) { m_logger.setLevel( level ); } + void setLogLevel( hiprtLogLevel level ) override { m_logger.setLevel( level ); } + + struct CpuGeometryData* getCpuGeom( hiprtGeometry ) override + { + return nullptr; + }; + struct CpuSceneData* getCpuScene( hiprtScene ) override + { + return nullptr; + }; template void logInfo( Args... args ) const @@ -148,6 +159,7 @@ class Context m_logger.print( hiprtLogLevelError, args... ); } + uint32_t getSMCount() const; uint32_t getMaxBlockSize() const; uint32_t getMaxGridSize() const; diff --git a/hiprt/impl/LbvhBuilder.cpp b/hiprt/impl/LbvhBuilder.cpp index e16b18c..2680a9e 100644 --- a/hiprt/impl/LbvhBuilder.cpp +++ b/hiprt/impl/LbvhBuilder.cpp @@ -42,7 +42,7 @@ size_t LbvhBuilder::getTemporaryBufferSize( const size_t count ) } size_t LbvhBuilder::getTemporaryBufferSize( - [[maybe_unused]] Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + [[maybe_unused]] GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); size_t size = getTemporaryBufferSize( primCount ); @@ -57,7 +57,7 @@ size_t LbvhBuilder::getTemporaryBufferSize( } size_t LbvhBuilder::getTemporaryBufferSize( - [[maybe_unused]] Context& context, + [[maybe_unused]] GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { @@ -65,7 +65,7 @@ size_t LbvhBuilder::getTemporaryBufferSize( } size_t LbvhBuilder::getStorageBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); const size_t primNodeCount = getMaxPrimNodeCount( buildInput, context.getRtip(), primCount ); @@ -75,7 +75,7 @@ size_t LbvhBuilder::getStorageBufferSize( } size_t LbvhBuilder::getStorageBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = buildInput.instanceCount; const size_t boxNodeCount = getMaxBoxNodeCount( buildInput, context.getRtip(), primCount ); @@ -84,7 +84,7 @@ size_t LbvhBuilder::getStorageBufferSize( } void LbvhBuilder::build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -124,7 +124,7 @@ void LbvhBuilder::build( } void LbvhBuilder::build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -164,7 +164,7 @@ void LbvhBuilder::build( } void LbvhBuilder::update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, @@ -198,7 +198,7 @@ void LbvhBuilder::update( } void LbvhBuilder::update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, diff --git a/hiprt/impl/LbvhBuilder.h b/hiprt/impl/LbvhBuilder.h index e8ff039..dd34119 100644 --- a/hiprt/impl/LbvhBuilder.h +++ b/hiprt/impl/LbvhBuilder.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include #include #include #include @@ -68,19 +68,19 @@ class LbvhBuilder static size_t getTemporaryBufferSize( const size_t count ); static size_t - getTemporaryBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static void build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -88,7 +88,7 @@ class LbvhBuilder hiprtDevicePtr buffer ); static void build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -97,7 +97,7 @@ class LbvhBuilder template static void build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, uint32_t geomType, @@ -106,7 +106,7 @@ class LbvhBuilder MemoryArena& storageMemoryArena ); static void update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -114,7 +114,7 @@ class LbvhBuilder hiprtDevicePtr buffer ); static void update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -123,7 +123,7 @@ class LbvhBuilder template static void update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, @@ -132,7 +132,7 @@ class LbvhBuilder template void LbvhBuilder::build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, uint32_t geomType, @@ -431,7 +431,7 @@ void LbvhBuilder::build( template void LbvhBuilder::update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, diff --git a/hiprt/impl/PlocBuilder.cpp b/hiprt/impl/PlocBuilder.cpp index 84e53f2..907919e 100644 --- a/hiprt/impl/PlocBuilder.cpp +++ b/hiprt/impl/PlocBuilder.cpp @@ -39,7 +39,7 @@ size_t PlocBuilder::getTemporaryBufferSize( const size_t count ) } size_t PlocBuilder::getTemporaryBufferSize( - [[maybe_unused]] Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + [[maybe_unused]] GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); size_t size = getTemporaryBufferSize( primCount ); @@ -54,7 +54,7 @@ size_t PlocBuilder::getTemporaryBufferSize( } size_t PlocBuilder::getTemporaryBufferSize( - [[maybe_unused]] Context& context, + [[maybe_unused]] GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { @@ -62,7 +62,7 @@ size_t PlocBuilder::getTemporaryBufferSize( } size_t PlocBuilder::getStorageBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = getPrimCount( buildInput ); const size_t primNodeCount = getMaxPrimNodeCount( buildInput, context.getRtip(), primCount ); @@ -72,7 +72,7 @@ size_t PlocBuilder::getStorageBufferSize( } size_t PlocBuilder::getStorageBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, [[maybe_unused]] const hiprtBuildOptions buildOptions ) { const size_t primCount = buildInput.instanceCount; const size_t boxNodeCount = getMaxBoxNodeCount( buildInput, context.getRtip(), primCount ); @@ -81,7 +81,7 @@ size_t PlocBuilder::getStorageBufferSize( } void PlocBuilder::build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -121,7 +121,7 @@ void PlocBuilder::build( } void PlocBuilder::build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -161,7 +161,7 @@ void PlocBuilder::build( } void PlocBuilder::update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, @@ -195,7 +195,7 @@ void PlocBuilder::update( } void PlocBuilder::update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, diff --git a/hiprt/impl/PlocBuilder.h b/hiprt/impl/PlocBuilder.h index 05eceab..f4118ad 100644 --- a/hiprt/impl/PlocBuilder.h +++ b/hiprt/impl/PlocBuilder.h @@ -26,7 +26,7 @@ #include #include #include -#include +#include #include #include #include @@ -70,19 +70,19 @@ class PlocBuilder static size_t getTemporaryBufferSize( const size_t count ); static size_t - getTemporaryBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static void build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -90,7 +90,7 @@ class PlocBuilder hiprtDevicePtr buffer ); static void build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -99,7 +99,7 @@ class PlocBuilder template static void build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, const uint32_t geomType, @@ -108,7 +108,7 @@ class PlocBuilder MemoryArena& storageMemoryArena ); static void update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -116,7 +116,7 @@ class PlocBuilder hiprtDevicePtr buffer ); static void update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -125,7 +125,7 @@ class PlocBuilder template static void update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, @@ -134,7 +134,7 @@ class PlocBuilder template void PlocBuilder::build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, uint32_t geomType, @@ -441,7 +441,7 @@ void PlocBuilder::build( template void PlocBuilder::update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, diff --git a/hiprt/impl/SbvhBuilder.cpp b/hiprt/impl/SbvhBuilder.cpp index a8b1b90..1f8913c 100644 --- a/hiprt/impl/SbvhBuilder.cpp +++ b/hiprt/impl/SbvhBuilder.cpp @@ -32,7 +32,7 @@ namespace hiprt { size_t SbvhBuilder::getTemporaryBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { size_t size = getTemporaryBufferSize( context, buildInput, buildOptions ); if ( buildInput.type == hiprtPrimitiveTypeTriangleMesh ) @@ -47,13 +47,13 @@ size_t SbvhBuilder::getTemporaryBufferSize( } size_t SbvhBuilder::getTemporaryBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { return getTemporaryBufferSize( context, buildInput, buildOptions ); } size_t SbvhBuilder::getStorageBufferSize( - Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { const float alpha = buildOptions.buildFlags & hiprtBuildFlagBitDisableSpatialSplits ? 1.0f : Alpha; const size_t primCount = getPrimCount( buildInput ); @@ -65,7 +65,7 @@ size_t SbvhBuilder::getStorageBufferSize( } size_t SbvhBuilder::getStorageBufferSize( - Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ) { const float alpha = buildOptions.buildFlags & hiprtBuildFlagBitDisableSpatialSplits ? 1.0f : Alpha; const size_t primCount = buildInput.instanceCount; @@ -81,7 +81,7 @@ size_t SbvhBuilder::getStorageBufferSize( } void SbvhBuilder::build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -119,7 +119,7 @@ void SbvhBuilder::build( } void SbvhBuilder::build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -159,7 +159,7 @@ void SbvhBuilder::build( } void SbvhBuilder::update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, @@ -191,7 +191,7 @@ void SbvhBuilder::update( } void SbvhBuilder::update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, [[maybe_unused]] hiprtDevicePtr temporaryBuffer, diff --git a/hiprt/impl/SbvhBuilder.h b/hiprt/impl/SbvhBuilder.h index 80906e4..1dd1cd7 100644 --- a/hiprt/impl/SbvhBuilder.h +++ b/hiprt/impl/SbvhBuilder.h @@ -27,7 +27,7 @@ #include #include #include -#include +#include #include #include #include @@ -78,22 +78,22 @@ class SbvhBuilder template static size_t - getTemporaryBufferSize( Context& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getTemporaryBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getTemporaryBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static size_t - getStorageBufferSize( Context& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); + getStorageBufferSize( GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions ); static void build( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -101,7 +101,7 @@ class SbvhBuilder hiprtDevicePtr buffer ); static void build( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -110,7 +110,7 @@ class SbvhBuilder template static void build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, uint32_t geomType, @@ -119,7 +119,7 @@ class SbvhBuilder MemoryArena& storageMemoryArena ); static void update( - Context& context, + GpuContext& context, const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -127,7 +127,7 @@ class SbvhBuilder hiprtDevicePtr buffer ); static void update( - Context& context, + GpuContext& context, const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions buildOptions, hiprtDevicePtr temporaryBuffer, @@ -136,7 +136,7 @@ class SbvhBuilder template static void update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, @@ -145,7 +145,7 @@ class SbvhBuilder template size_t -SbvhBuilder::getTemporaryBufferSize( Context& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ) +SbvhBuilder::getTemporaryBufferSize( GpuContext& context, const BuildInput& buildInput, const hiprtBuildOptions buildOptions ) { size_t count{}; if constexpr ( std::is_same::value ) @@ -172,7 +172,7 @@ SbvhBuilder::getTemporaryBufferSize( Context& context, const BuildInput& buildIn template void SbvhBuilder::build( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, uint32_t geomType, @@ -642,7 +642,7 @@ void SbvhBuilder::build( template void SbvhBuilder::update( - Context& context, + GpuContext& context, PrimitiveContainer& primitives, const hiprtBuildOptions buildOptions, oroStream stream, diff --git a/hiprt/impl/Utility.h b/hiprt/impl/Utility.h index 4c781ed..d0c7a48 100644 --- a/hiprt/impl/Utility.h +++ b/hiprt/impl/Utility.h @@ -29,7 +29,7 @@ namespace hiprt { -class Context; +class GpuContext; class Utility { diff --git a/hiprt/impl/hiprt.cpp b/hiprt/impl/hiprt.cpp index 7059b41..9f0d794 100644 --- a/hiprt/impl/hiprt.cpp +++ b/hiprt/impl/hiprt.cpp @@ -25,7 +25,7 @@ #include #include #include -#include +#include #include #include @@ -38,7 +38,7 @@ hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreat try { - Context* ctxt = new Context( input ); + GpuContext* ctxt = new GpuContext( input ); contextOut = reinterpret_cast( ctxt ); } catch ( std::exception& e ) @@ -53,7 +53,7 @@ hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreat hiprtError hiprtDestroyContext( hiprtContext context ) { if ( !context ) return hiprtErrorInvalidParameter; - delete reinterpret_cast( context ); + delete reinterpret_cast( context ); return hiprtSuccess; } @@ -88,13 +88,13 @@ hiprtError hiprtCreateGeometries( try { std::vector geometries = - reinterpret_cast( context )->createGeometries( buildInputs, buildOptions ); + reinterpret_cast( context )->createGeometries( buildInputs, buildOptions ); for ( uint32_t i = 0; i < numGeometries; ++i ) *geometriesOut[i] = geometries[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -119,11 +119,11 @@ hiprtError hiprtDestroyGeometries( hiprtContext context, uint32_t numGeometries, try { - reinterpret_cast( context )->destroyGeometries( geometries ); + reinterpret_cast( context )->destroyGeometries( geometries ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -169,12 +169,12 @@ hiprtError hiprtBuildGeometries( switch ( buildOperation ) { case hiprtBuildOperationBuild: { - reinterpret_cast( context )->buildGeometries( + reinterpret_cast( context )->buildGeometries( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } case hiprtBuildOperationUpdate: { - reinterpret_cast( context )->updateGeometries( + reinterpret_cast( context )->updateGeometries( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } @@ -182,7 +182,7 @@ hiprtError hiprtBuildGeometries( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -211,11 +211,11 @@ hiprtError hiprtGetGeometriesBuildTemporaryBufferSize( try { - sizeOut = reinterpret_cast( context )->getGeometriesBuildTempBufferSize( buildInputs, buildOptions ); + sizeOut = reinterpret_cast( context )->getGeometriesBuildTempBufferSize( buildInputs, buildOptions ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -250,13 +250,13 @@ hiprtError hiprtCompactGeometries( try { std::vector compactedGeometries = - reinterpret_cast( context )->compactGeometries( geometries, reinterpret_cast( stream ) ); + reinterpret_cast( context )->compactGeometries( geometries, reinterpret_cast( stream ) ); for ( uint32_t i = 0; i < numGeometries; ++i ) *geometriesOut[i] = compactedGeometries[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -288,13 +288,13 @@ hiprtError hiprtCreateScenes( try { - std::vector scenes = reinterpret_cast( context )->createScenes( buildInputs, buildOptions ); + std::vector scenes = reinterpret_cast( context )->createScenes( buildInputs, buildOptions ); for ( uint32_t i = 0; i < numScenes; ++i ) *scenesOut[i] = scenes[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -316,11 +316,11 @@ hiprtError hiprtDestroyScenes( hiprtContext context, uint32_t numScenes, hiprtSc try { - reinterpret_cast( context )->destroyScenes( scenes ); + reinterpret_cast( context )->destroyScenes( scenes ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -365,12 +365,12 @@ hiprtError hiprtBuildScenes( switch ( buildOperation ) { case hiprtBuildOperationBuild: { - reinterpret_cast( context )->buildScenes( + reinterpret_cast( context )->buildScenes( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } case hiprtBuildOperationUpdate: { - reinterpret_cast( context )->updateScenes( + reinterpret_cast( context )->updateScenes( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } @@ -378,7 +378,7 @@ hiprtError hiprtBuildScenes( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -407,11 +407,11 @@ hiprtError hiprtGetScenesBuildTemporaryBufferSize( try { - sizeOut = reinterpret_cast( context )->getScenesBuildTempBufferSize( buildInputs, buildOptions ); + sizeOut = reinterpret_cast( context )->getScenesBuildTempBufferSize( buildInputs, buildOptions ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -440,13 +440,13 @@ hiprtError hiprtCompactScenes( try { std::vector compactedScenes = - reinterpret_cast( context )->compactScenes( scenes, reinterpret_cast( stream ) ); + reinterpret_cast( context )->compactScenes( scenes, reinterpret_cast( stream ) ); for ( uint32_t i = 0; i < numScenes; ++i ) *scenesOut[i] = compactedScenes[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -458,11 +458,11 @@ hiprtCreateFuncTable( hiprtContext context, uint32_t numGeomTypes, uint32_t numR if ( !context ) return hiprtErrorInvalidParameter; try { - funcTableOut = reinterpret_cast( context )->createFuncTable( numGeomTypes, numRayTypes ); + funcTableOut = reinterpret_cast( context )->createFuncTable( numGeomTypes, numRayTypes ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -474,11 +474,11 @@ hiprtSetFuncTable( hiprtContext context, hiprtFuncTable funcTable, uint32_t geom if ( !context || !funcTable ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setFuncTable( funcTable, geomType, rayType, set ); + reinterpret_cast( context )->setFuncTable( funcTable, geomType, rayType, set ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -489,11 +489,11 @@ hiprtError hiprtDestroyFuncTable( hiprtContext context, hiprtFuncTable funcTable if ( !context || !funcTable ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->destroyFuncTable( funcTable ); + reinterpret_cast( context )->destroyFuncTable( funcTable ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -505,11 +505,11 @@ hiprtError hiprtCreateGlobalStackBuffer( if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->createGlobalStackBuffer( input, stackBufferOut ); + reinterpret_cast( context )->createGlobalStackBuffer( input, stackBufferOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -520,11 +520,11 @@ hiprtError hiprtDestroyGlobalStackBuffer( hiprtContext context, hiprtGlobalStack if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->destroyGlobalStackBuffer( stackBuffer ); + reinterpret_cast( context )->destroyGlobalStackBuffer( stackBuffer ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -535,11 +535,11 @@ hiprtError hiprtSaveGeometry( hiprtContext context, hiprtGeometry geometry, cons if ( !context || !geometry || !filename ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->saveGeometry( geometry, filename ); + reinterpret_cast( context )->saveGeometry( geometry, filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -550,11 +550,11 @@ hiprtError hiprtLoadGeometry( hiprtContext context, hiprtGeometry& geometryOut, if ( !context || !filename ) return hiprtErrorInvalidParameter; try { - geometryOut = reinterpret_cast( context )->loadGeometry( filename ); + geometryOut = reinterpret_cast( context )->loadGeometry( filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -565,11 +565,11 @@ hiprtError hiprtSaveScene( hiprtContext context, hiprtScene scene, const char* f if ( !context || !scene || !filename ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->saveScene( scene, filename ); + reinterpret_cast( context )->saveScene( scene, filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -580,11 +580,11 @@ hiprtError hiprtLoadScene( hiprtContext context, hiprtScene& sceneOut, const cha if ( !context || !sceneOut || !filename ) return hiprtErrorInvalidParameter; try { - sceneOut = reinterpret_cast( context )->loadScene( filename ); + sceneOut = reinterpret_cast( context )->loadScene( filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -595,11 +595,11 @@ hiprtError hiprtExportGeometryAabb( hiprtContext context, hiprtGeometry geometry if ( !context || !geometry ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->exportGeometryAabb( geometry, aabbMinOut, aabbMaxOut ); + reinterpret_cast( context )->exportGeometryAabb( geometry, aabbMinOut, aabbMaxOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -610,11 +610,11 @@ hiprtError hiprtExportSceneAabb( hiprtContext context, hiprtScene scene, float3& if ( !context || !scene ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->exportSceneAabb( scene, aabbMinOut, aabbMaxOut ); + reinterpret_cast( context )->exportSceneAabb( scene, aabbMinOut, aabbMaxOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -670,7 +670,7 @@ hiprtError hiprtBuildTraceKernels( std::vector functions; oroModule module = nullptr; - reinterpret_cast( context )->buildKernels( + reinterpret_cast( context )->buildKernels( funcNames, src, moduleName, @@ -691,7 +691,7 @@ hiprtError hiprtBuildTraceKernels( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -730,7 +730,7 @@ hiprtError hiprtBuildTraceKernelsFromBitcode( } std::string_view binary( bitcodeBinary, bitcodeBinarySize ); std::vector functions; - reinterpret_cast( context )->buildKernelsFromBitcode( + reinterpret_cast( context )->buildKernelsFromBitcode( funcNames, moduleName, binary, numGeomTypes, numRayTypes, funcNameSets, functions, cache ); for ( uint32_t i = 0; i < numFunctions; ++i ) @@ -738,7 +738,7 @@ hiprtError hiprtBuildTraceKernelsFromBitcode( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } @@ -750,11 +750,11 @@ hiprtError hiprtSetCacheDirPath( hiprtContext context, const char* path ) if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setCacheDir( path ); + reinterpret_cast( context )->setCacheDir( path ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; @@ -765,11 +765,11 @@ hiprtError hiprtSetLogLevel( hiprtContext context, hiprtLogLevel level ) if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setLogLevel( level ); + reinterpret_cast( context )->setLogLevel( level ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + reinterpret_cast( context )->logError( e.what() ); return hiprtErrorInternal; } return hiprtSuccess; From 1f7427f51a7e20606b4aeda5b3edb8b47662d8bd Mon Sep 17 00:00:00 2001 From: Witold Wojciechowski Date: Thu, 14 May 2026 22:48:43 +0200 Subject: [PATCH 2/7] Added CpuContext.h, and simple smoke test --- CMakeLists.txt | 20 ++++++ hiprt/impl/ContextBase.h | 6 ++ hiprt/impl/CpuContext.cpp | 8 +++ hiprt/impl/CpuContext.h | 13 ++-- hiprt/impl/GpuContext.h | 12 ++-- hiprt/impl/hiprt.cpp | 142 +++++++++++++++++++++++--------------- 6 files changed, 135 insertions(+), 66 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 625ea77..5e1c24c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -720,3 +720,23 @@ if(HIPRTEW) target_compile_definitions(hiprtewtest PRIVATE GTEST_HAS_TR1_TUPLE=0) endif() + + +# Project: CPU dispatch smoke test +# Tiny standalone program that proves hiprtCreateContext(deviceType=hiprtDeviceCPU) +# routes through CpuContext (not GpuContext). It does not need a GPU. +add_executable(cpu_smoke ${CMAKE_CURRENT_SOURCE_DIR}/test/cpu_smoke.cpp) + +target_include_directories(cpu_smoke PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR} + ${CMAKE_CURRENT_SOURCE_DIR}/contrib/Orochi) + +target_link_libraries(cpu_smoke PRIVATE ${HIPRT_NAME}) + +if(WIN32) + target_link_libraries(cpu_smoke PRIVATE version) +endif() + +if( CMAKE_BUILD_TYPE STREQUAL "Debug" ) + set_target_properties(cpu_smoke PROPERTIES OUTPUT_NAME "cpu_smokeD") +endif() diff --git a/hiprt/impl/ContextBase.h b/hiprt/impl/ContextBase.h index 5a74d7d..841aa15 100644 --- a/hiprt/impl/ContextBase.h +++ b/hiprt/impl/ContextBase.h @@ -88,6 +88,12 @@ class ContextBase virtual void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) = 0; virtual void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) = 0; + virtual hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) = 0; + virtual void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) = 0; + virtual void destroyFuncTable( hiprtFuncTable funcTable ) = 0; + + virtual void setCacheDir( const std::filesystem::path& path ) = 0; + virtual void setLogLevel( hiprtLogLevel level ) = 0; virtual struct CpuGeometryData* getCpuGeom( hiprtGeometry ) = 0; diff --git a/hiprt/impl/CpuContext.cpp b/hiprt/impl/CpuContext.cpp index 6064cda..317d592 100644 --- a/hiprt/impl/CpuContext.cpp +++ b/hiprt/impl/CpuContext.cpp @@ -1,5 +1,6 @@ #include "CpuContext.h" +#include #include namespace hiprt @@ -8,10 +9,17 @@ namespace { [[noreturn]] void throwNotImplemented( const char* what ) { + std::cerr << "[CpuContext] not implemented: " << what << std::endl; throw std::runtime_error( what ); } } // namespace +CpuContext::CpuContext( const hiprtContextCreationInput& input ) +{ + std::cerr << "[CpuContext] constructed (numCpuThreads=" << input.numCpuThreads << ")" << std::endl; + // TODO: Initialize Embree RTCDevice here, honoring input.numCpuThreads. +} + std::vector CpuContext::createGeometries( const std::vector&, const hiprtBuildOptions ) { diff --git a/hiprt/impl/CpuContext.h b/hiprt/impl/CpuContext.h index d3a0361..adf8079 100644 --- a/hiprt/impl/CpuContext.h +++ b/hiprt/impl/CpuContext.h @@ -13,7 +13,10 @@ namespace hiprt{ class CpuContext : public ContextBase { public: - std::vector + CpuContext( const hiprtContextCreationInput& input ); + ~CpuContext() override = default; + + std::vector createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; void destroyGeometries( const std::vector& geometries ) override; @@ -61,9 +64,11 @@ class CpuContext : public ContextBase std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; - hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ); - void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ); - void destroyFuncTable( hiprtFuncTable funcTable ); + hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) override; + void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) override; + void destroyFuncTable( hiprtFuncTable funcTable ) override; + + void setCacheDir( const std::filesystem::path& path ) override { (void)path; } void setLogLevel( hiprtLogLevel level ) override { (void)level; } diff --git a/hiprt/impl/GpuContext.h b/hiprt/impl/GpuContext.h index 5970918..7549193 100644 --- a/hiprt/impl/GpuContext.h +++ b/hiprt/impl/GpuContext.h @@ -83,14 +83,14 @@ class GpuContext : public ContextBase oroStream stream, std::vector& buffers ) override; - size_t - getScenesBuildTempBufferSize( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ); + size_t getScenesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; - hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ); - void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ); - void destroyFuncTable( hiprtFuncTable funcTable ); + hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) override; + void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) override; + void destroyFuncTable( hiprtFuncTable funcTable ) override; void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) override; void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) override; @@ -128,7 +128,7 @@ class GpuContext : public ContextBase std::vector& functions, bool cache ) override; - void setCacheDir( const std::filesystem::path& path ); + void setCacheDir( const std::filesystem::path& path ) override; void setLogLevel( hiprtLogLevel level ) override { m_logger.setLevel( level ); } diff --git a/hiprt/impl/hiprt.cpp b/hiprt/impl/hiprt.cpp index 9f0d794..770f7a2 100644 --- a/hiprt/impl/hiprt.cpp +++ b/hiprt/impl/hiprt.cpp @@ -24,6 +24,8 @@ #include #include +#include +#include #include #include #include @@ -31,15 +33,43 @@ using namespace hiprt; +namespace +{ +// Helper: every public API entry point gets the active backend through this +// cast. Using ContextBase* (not GpuContext*) means virtual dispatch routes the +// call to GpuContext or CpuContext depending on what was created. +inline ContextBase* asContext( hiprtContext context ) noexcept +{ + return reinterpret_cast( context ); +} +} // namespace + hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreationInput& input, hiprtContext& contextOut ) { - oroInitialize( ( input.deviceType == hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA, 0, g_hip_paths, g_hiprtc_paths ); if ( hiprtApiVersion != HIPRT_API_VERSION ) return hiprtErrorInvalidApiVersion; try { - GpuContext* ctxt = new GpuContext( input ); - contextOut = reinterpret_cast( ctxt ); + ContextBase* ctxt = nullptr; + + if ( input.deviceType == hiprtDeviceCPU ) + { + // Pure CPU context: no Orochi/HIP/CUDA init needed. + ctxt = new CpuContext( input ); + } + else + { + // GPU (and, later, GPU+CPU hybrid). Picking the underlying API: + // AMD bit set => HIP, otherwise treat as NVIDIA/CUDA. + oroInitialize( + ( input.deviceType & hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA, + 0, + g_hip_paths, + g_hiprtc_paths ); + ctxt = new GpuContext( input ); + } + + contextOut = reinterpret_cast( ctxt ); } catch ( std::exception& e ) { @@ -53,7 +83,7 @@ hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreat hiprtError hiprtDestroyContext( hiprtContext context ) { if ( !context ) return hiprtErrorInvalidParameter; - delete reinterpret_cast( context ); + delete asContext( context ); return hiprtSuccess; } @@ -88,13 +118,13 @@ hiprtError hiprtCreateGeometries( try { std::vector geometries = - reinterpret_cast( context )->createGeometries( buildInputs, buildOptions ); + asContext( context )->createGeometries( buildInputs, buildOptions ); for ( uint32_t i = 0; i < numGeometries; ++i ) *geometriesOut[i] = geometries[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -119,11 +149,11 @@ hiprtError hiprtDestroyGeometries( hiprtContext context, uint32_t numGeometries, try { - reinterpret_cast( context )->destroyGeometries( geometries ); + asContext( context )->destroyGeometries( geometries ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -169,12 +199,12 @@ hiprtError hiprtBuildGeometries( switch ( buildOperation ) { case hiprtBuildOperationBuild: { - reinterpret_cast( context )->buildGeometries( + asContext( context )->buildGeometries( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } case hiprtBuildOperationUpdate: { - reinterpret_cast( context )->updateGeometries( + asContext( context )->updateGeometries( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } @@ -182,7 +212,7 @@ hiprtError hiprtBuildGeometries( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -211,11 +241,11 @@ hiprtError hiprtGetGeometriesBuildTemporaryBufferSize( try { - sizeOut = reinterpret_cast( context )->getGeometriesBuildTempBufferSize( buildInputs, buildOptions ); + sizeOut = asContext( context )->getGeometriesBuildTempBufferSize( buildInputs, buildOptions ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -250,13 +280,13 @@ hiprtError hiprtCompactGeometries( try { std::vector compactedGeometries = - reinterpret_cast( context )->compactGeometries( geometries, reinterpret_cast( stream ) ); + asContext( context )->compactGeometries( geometries, reinterpret_cast( stream ) ); for ( uint32_t i = 0; i < numGeometries; ++i ) *geometriesOut[i] = compactedGeometries[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -288,13 +318,13 @@ hiprtError hiprtCreateScenes( try { - std::vector scenes = reinterpret_cast( context )->createScenes( buildInputs, buildOptions ); + std::vector scenes = asContext( context )->createScenes( buildInputs, buildOptions ); for ( uint32_t i = 0; i < numScenes; ++i ) *scenesOut[i] = scenes[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -316,11 +346,11 @@ hiprtError hiprtDestroyScenes( hiprtContext context, uint32_t numScenes, hiprtSc try { - reinterpret_cast( context )->destroyScenes( scenes ); + asContext( context )->destroyScenes( scenes ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -365,12 +395,12 @@ hiprtError hiprtBuildScenes( switch ( buildOperation ) { case hiprtBuildOperationBuild: { - reinterpret_cast( context )->buildScenes( + asContext( context )->buildScenes( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } case hiprtBuildOperationUpdate: { - reinterpret_cast( context )->updateScenes( + asContext( context )->updateScenes( buildInputs, buildOptions, temporaryBuffer, reinterpret_cast( stream ), buffers ); break; } @@ -378,7 +408,7 @@ hiprtError hiprtBuildScenes( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -407,11 +437,11 @@ hiprtError hiprtGetScenesBuildTemporaryBufferSize( try { - sizeOut = reinterpret_cast( context )->getScenesBuildTempBufferSize( buildInputs, buildOptions ); + sizeOut = asContext( context )->getScenesBuildTempBufferSize( buildInputs, buildOptions ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -440,13 +470,13 @@ hiprtError hiprtCompactScenes( try { std::vector compactedScenes = - reinterpret_cast( context )->compactScenes( scenes, reinterpret_cast( stream ) ); + asContext( context )->compactScenes( scenes, reinterpret_cast( stream ) ); for ( uint32_t i = 0; i < numScenes; ++i ) *scenesOut[i] = compactedScenes[i]; } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -458,11 +488,11 @@ hiprtCreateFuncTable( hiprtContext context, uint32_t numGeomTypes, uint32_t numR if ( !context ) return hiprtErrorInvalidParameter; try { - funcTableOut = reinterpret_cast( context )->createFuncTable( numGeomTypes, numRayTypes ); + funcTableOut = asContext( context )->createFuncTable( numGeomTypes, numRayTypes ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -474,11 +504,11 @@ hiprtSetFuncTable( hiprtContext context, hiprtFuncTable funcTable, uint32_t geom if ( !context || !funcTable ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setFuncTable( funcTable, geomType, rayType, set ); + asContext( context )->setFuncTable( funcTable, geomType, rayType, set ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -489,11 +519,11 @@ hiprtError hiprtDestroyFuncTable( hiprtContext context, hiprtFuncTable funcTable if ( !context || !funcTable ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->destroyFuncTable( funcTable ); + asContext( context )->destroyFuncTable( funcTable ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -505,11 +535,11 @@ hiprtError hiprtCreateGlobalStackBuffer( if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->createGlobalStackBuffer( input, stackBufferOut ); + asContext( context )->createGlobalStackBuffer( input, stackBufferOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -520,11 +550,11 @@ hiprtError hiprtDestroyGlobalStackBuffer( hiprtContext context, hiprtGlobalStack if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->destroyGlobalStackBuffer( stackBuffer ); + asContext( context )->destroyGlobalStackBuffer( stackBuffer ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -535,11 +565,11 @@ hiprtError hiprtSaveGeometry( hiprtContext context, hiprtGeometry geometry, cons if ( !context || !geometry || !filename ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->saveGeometry( geometry, filename ); + asContext( context )->saveGeometry( geometry, filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -550,11 +580,11 @@ hiprtError hiprtLoadGeometry( hiprtContext context, hiprtGeometry& geometryOut, if ( !context || !filename ) return hiprtErrorInvalidParameter; try { - geometryOut = reinterpret_cast( context )->loadGeometry( filename ); + geometryOut = asContext( context )->loadGeometry( filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -565,11 +595,11 @@ hiprtError hiprtSaveScene( hiprtContext context, hiprtScene scene, const char* f if ( !context || !scene || !filename ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->saveScene( scene, filename ); + asContext( context )->saveScene( scene, filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -580,11 +610,11 @@ hiprtError hiprtLoadScene( hiprtContext context, hiprtScene& sceneOut, const cha if ( !context || !sceneOut || !filename ) return hiprtErrorInvalidParameter; try { - sceneOut = reinterpret_cast( context )->loadScene( filename ); + sceneOut = asContext( context )->loadScene( filename ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -595,11 +625,11 @@ hiprtError hiprtExportGeometryAabb( hiprtContext context, hiprtGeometry geometry if ( !context || !geometry ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->exportGeometryAabb( geometry, aabbMinOut, aabbMaxOut ); + asContext( context )->exportGeometryAabb( geometry, aabbMinOut, aabbMaxOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -610,11 +640,11 @@ hiprtError hiprtExportSceneAabb( hiprtContext context, hiprtScene scene, float3& if ( !context || !scene ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->exportSceneAabb( scene, aabbMinOut, aabbMaxOut ); + asContext( context )->exportSceneAabb( scene, aabbMinOut, aabbMaxOut ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -670,7 +700,7 @@ hiprtError hiprtBuildTraceKernels( std::vector functions; oroModule module = nullptr; - reinterpret_cast( context )->buildKernels( + asContext( context )->buildKernels( funcNames, src, moduleName, @@ -691,7 +721,7 @@ hiprtError hiprtBuildTraceKernels( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -730,7 +760,7 @@ hiprtError hiprtBuildTraceKernelsFromBitcode( } std::string_view binary( bitcodeBinary, bitcodeBinarySize ); std::vector functions; - reinterpret_cast( context )->buildKernelsFromBitcode( + asContext( context )->buildKernelsFromBitcode( funcNames, moduleName, binary, numGeomTypes, numRayTypes, funcNameSets, functions, cache ); for ( uint32_t i = 0; i < numFunctions; ++i ) @@ -738,7 +768,7 @@ hiprtError hiprtBuildTraceKernelsFromBitcode( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -750,11 +780,11 @@ hiprtError hiprtSetCacheDirPath( hiprtContext context, const char* path ) if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setCacheDir( path ); + asContext( context )->setCacheDir( path ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; @@ -765,11 +795,11 @@ hiprtError hiprtSetLogLevel( hiprtContext context, hiprtLogLevel level ) if ( !context ) return hiprtErrorInvalidParameter; try { - reinterpret_cast( context )->setLogLevel( level ); + asContext( context )->setLogLevel( level ); } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } return hiprtSuccess; From c46a4a5767c64f58751f7f2b987dc04408488b4a Mon Sep 17 00:00:00 2001 From: unknown Date: Fri, 15 May 2026 05:36:06 +0200 Subject: [PATCH 3/7] cpu: basic geometry build and update with Embree and GTest setup --- CMakeLists.txt | 44 ++++++--- hiprt/impl/CpuContext.cpp | 185 ++++++++++++++++++++++++++++++++++---- hiprt/impl/CpuContext.h | 14 ++- hiprt/impl/CpuTypes.h | 20 +++++ test/CpuGeometryTest.cpp | 161 +++++++++++++++++++++++++++++++++ test/CpuGeometryTest.h | 19 ++++ 6 files changed, 411 insertions(+), 32 deletions(-) create mode 100644 hiprt/impl/CpuTypes.h create mode 100644 test/CpuGeometryTest.cpp create mode 100644 test/CpuGeometryTest.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 5e1c24c..608da9d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -594,6 +594,17 @@ endif() target_include_directories(${HIPRT_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}) target_include_directories(${HIPRT_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/contrib/Orochi) +target_include_directories(${HIPRT_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/contrib/embree/include) + +if(WIN32) + target_link_directories(${HIPRT_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/contrib/embree/win) + target_link_libraries(${HIPRT_NAME} ${CMAKE_CURRENT_SOURCE_DIR}/contrib/embree/win/embree4.lib) +endif() + +if(UNIX) + target_link_directories(${HIPRT_NAME} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/contrib/embree/linux) + target_link_libraries(${HIPRT_NAME} embree4 tbb) +endif() file(GLOB_RECURSE hiprt_sources "${CMAKE_CURRENT_SOURCE_DIR}/hiprt/*.h" "${CMAKE_CURRENT_SOURCE_DIR}/hiprt/*.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/hiprt/*.inl") list(FILTER hiprt_sources EXCLUDE REGEX "hiprt/bitcodes/.*") @@ -664,7 +675,16 @@ if(NOT NO_UNITTEST) target_link_libraries(unittest PRIVATE pthread dl) endif() - file(GLOB_RECURSE unittest_sources "${CMAKE_CURRENT_SOURCE_DIR}/test/hiprtT*.h" "${CMAKE_CURRENT_SOURCE_DIR}/test/hiprtT*.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/test/shared.h" "${CMAKE_CURRENT_SOURCE_DIR}/test/main.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/test/CornellBox.h" "${CMAKE_CURRENT_SOURCE_DIR}/test/kernels/*.h" "${CMAKE_CURRENT_SOURCE_DIR}/contrib/gtest-1.6.0/gtest-all.cc") + file(GLOB_RECURSE unittest_sources + "${CMAKE_CURRENT_SOURCE_DIR}/test/hiprtT*.h" + "${CMAKE_CURRENT_SOURCE_DIR}/test/hiprtT*.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/test/Cpu*Test.h" + "${CMAKE_CURRENT_SOURCE_DIR}/test/Cpu*Test.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/test/shared.h" + "${CMAKE_CURRENT_SOURCE_DIR}/test/main.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/test/CornellBox.h" + "${CMAKE_CURRENT_SOURCE_DIR}/test/kernels/*.h" + "${CMAKE_CURRENT_SOURCE_DIR}/contrib/gtest-1.6.0/gtest-all.cc") target_sources(unittest PRIVATE ${unittest_sources} ${orochi_sources}) @@ -725,18 +745,18 @@ endif() # Project: CPU dispatch smoke test # Tiny standalone program that proves hiprtCreateContext(deviceType=hiprtDeviceCPU) # routes through CpuContext (not GpuContext). It does not need a GPU. -add_executable(cpu_smoke ${CMAKE_CURRENT_SOURCE_DIR}/test/cpu_smoke.cpp) +# add_executable(cpu_smoke ${CMAKE_CURRENT_SOURCE_DIR}/test/cpu_smoke.cpp) -target_include_directories(cpu_smoke PRIVATE - ${CMAKE_CURRENT_SOURCE_DIR} - ${CMAKE_CURRENT_SOURCE_DIR}/contrib/Orochi) +# target_include_directories(cpu_smoke PRIVATE +# ${CMAKE_CURRENT_SOURCE_DIR} +# ${CMAKE_CURRENT_SOURCE_DIR}/contrib/Orochi) -target_link_libraries(cpu_smoke PRIVATE ${HIPRT_NAME}) +# target_link_libraries(cpu_smoke PRIVATE ${HIPRT_NAME}) -if(WIN32) - target_link_libraries(cpu_smoke PRIVATE version) -endif() +# if(WIN32) +# target_link_libraries(cpu_smoke PRIVATE version) +# endif() -if( CMAKE_BUILD_TYPE STREQUAL "Debug" ) - set_target_properties(cpu_smoke PROPERTIES OUTPUT_NAME "cpu_smokeD") -endif() +# if( CMAKE_BUILD_TYPE STREQUAL "Debug" ) +# set_target_properties(cpu_smoke PROPERTIES OUTPUT_NAME "cpu_smokeD") +# endif() diff --git a/hiprt/impl/CpuContext.cpp b/hiprt/impl/CpuContext.cpp index 317d592..b253e08 100644 --- a/hiprt/impl/CpuContext.cpp +++ b/hiprt/impl/CpuContext.cpp @@ -16,36 +16,187 @@ namespace CpuContext::CpuContext( const hiprtContextCreationInput& input ) { - std::cerr << "[CpuContext] constructed (numCpuThreads=" << input.numCpuThreads << ")" << std::endl; - // TODO: Initialize Embree RTCDevice here, honoring input.numCpuThreads. + std::string config; + if ( input.numCpuThreads != 0 ) + config = "threads=" + std::to_string( input.numCpuThreads ); + + m_rtcDevice = rtcNewDevice( config.empty() ? nullptr : config.c_str() ); + + if ( m_rtcDevice == nullptr ) + { + const RTCError err = rtcGetDeviceError( nullptr ); + std::cerr << "[CpuContext] rtcNewDevice failed, RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewDevice failed" ); + } + + std::cerr << "[CpuContext] constructed (numCpuThreads=" << input.numCpuThreads + << ( input.numCpuThreads == 0 ? " -> all available)" : ")" ) << std::endl; +} + +CpuContext::~CpuContext() +{ + if ( m_rtcDevice != nullptr ) + { + rtcReleaseDevice( m_rtcDevice ); + m_rtcDevice = nullptr; + } } std::vector CpuContext::createGeometries( - const std::vector&, const hiprtBuildOptions ) + const std::vector& buildInputs, const hiprtBuildOptions ) +{ + std::vector geometries; + geometries.reserve( buildInputs.size() ); + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + CpuGeometryData* data = new CpuGeometryData(); + data->rtcDevice = m_rtcDevice; + data->rtcScene = rtcNewScene( m_rtcDevice ); + + if ( data->rtcScene == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + delete data; + std::cerr << "[CpuContext] rtcNewScene failed at index " << i + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewScene failed" ); + } + + geometries.push_back( reinterpret_cast( data ) ); + } + + return geometries; +} + +void CpuContext::destroyGeometries( const std::vector& geometries ) { - throwNotImplemented( "CpuContext::createGeometries is not implemented yet." ); + for ( hiprtGeometry geometry : geometries ) + { + CpuGeometryData* data = reinterpret_cast( geometry ); + if ( data == nullptr ) continue; + + if ( data->rtcScene != nullptr ) + { + rtcReleaseScene( data->rtcScene ); + data->rtcScene = nullptr; + } + + delete data; + } } -void CpuContext::destroyGeometries( const std::vector& ) { throwNotImplemented( "CpuContext::destroyGeometries is not implemented yet." ); } +void CpuContext::buildGeometryEntry( CpuGeometryData* data, const hiprtGeometryBuildInput& input, size_t index ) +{ + if ( data == nullptr || data->rtcScene == nullptr ) + throw std::runtime_error( "CpuContext::buildGeometryEntry: invalid geometry handle" ); + + if ( input.type == hiprtPrimitiveTypeTriangleMesh ) + { + const hiprtTriangleMeshPrimitive& mesh = input.primitive.triangleMesh; + + RTCGeometry geom = rtcNewGeometry( m_rtcDevice, RTC_GEOMETRY_TYPE_TRIANGLE ); + if ( geom == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + std::cerr << "[CpuContext] rtcNewGeometry failed at index " << index + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewGeometry failed" ); + } + + rtcSetSharedGeometryBuffer( + geom, RTC_BUFFER_TYPE_VERTEX, 0, RTC_FORMAT_FLOAT3, + mesh.vertices, 0, mesh.vertexStride, mesh.vertexCount ); + + if ( mesh.triangleIndices != nullptr ) + { + rtcSetSharedGeometryBuffer( + geom, RTC_BUFFER_TYPE_INDEX, 0, RTC_FORMAT_UINT3, + mesh.triangleIndices, 0, mesh.triangleStride, mesh.triangleCount ); + } + else + { + const uint32_t triCount = mesh.triangleCount != 0 ? mesh.triangleCount : ( mesh.vertexCount / 3 ); + uint32_t* indices = static_cast( rtcSetNewGeometryBuffer( + geom, RTC_BUFFER_TYPE_INDEX, 0, RTC_FORMAT_UINT3, 3 * sizeof( uint32_t ), triCount ) ); + if ( indices == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + rtcReleaseGeometry( geom ); + std::cerr << "[CpuContext] rtcSetNewGeometryBuffer failed at index " << index + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcSetNewGeometryBuffer failed" ); + } + for ( uint32_t t = 0; t < triCount; ++t ) + { + indices[3 * t + 0] = 3 * t + 0; + indices[3 * t + 1] = 3 * t + 1; + indices[3 * t + 2] = 3 * t + 2; + } + } + + rtcCommitGeometry( geom ); + rtcAttachGeometry( data->rtcScene, geom ); + rtcReleaseGeometry( geom ); + } + else if ( input.type == hiprtPrimitiveTypeAABBList ) + { + std::cerr << "[CpuContext] Warning: AABB geometry build is not implemented yet." << std::endl; + } + else + { + throw std::runtime_error( "CpuContext::buildGeometryEntry: unknown hiprtPrimitiveType" ); + } + + rtcCommitScene( data->rtcScene ); +} void CpuContext::buildGeometries( - const std::vector&, - const hiprtBuildOptions, - hiprtDevicePtr, - oroStream, - std::vector& ) + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) { - throwNotImplemented( "CpuContext::buildGeometries is not implemented yet." ); + (void)buildOptions; + (void)temporaryBuffer; + (void)stream; + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + buildGeometryEntry( getCpuGeom( reinterpret_cast( buffers[i] ) ), buildInputs[i], i ); } void CpuContext::updateGeometries( - const std::vector&, - const hiprtBuildOptions, - hiprtDevicePtr, - oroStream, - std::vector& ) + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) { - throwNotImplemented( "CpuContext::updateGeometries is not implemented yet." ); + (void)buildOptions; + (void)temporaryBuffer; + (void)stream; + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + CpuGeometryData* data = getCpuGeom( reinterpret_cast( buffers[i] ) ); + if ( data == nullptr ) + throw std::runtime_error( "CpuContext::updateGeometries: invalid geometry handle" ); + + if ( data->rtcScene != nullptr ) + rtcReleaseScene( data->rtcScene ); + + data->rtcScene = rtcNewScene( m_rtcDevice ); + if ( data->rtcScene == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + std::cerr << "[CpuContext] rtcNewScene failed at index " << i + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewScene failed" ); + } + + buildGeometryEntry( data, buildInputs[i], i ); + } } size_t CpuContext::getGeometriesBuildTempBufferSize( const std::vector&, const hiprtBuildOptions ) diff --git a/hiprt/impl/CpuContext.h b/hiprt/impl/CpuContext.h index adf8079..66698af 100644 --- a/hiprt/impl/CpuContext.h +++ b/hiprt/impl/CpuContext.h @@ -1,5 +1,6 @@ #pragma once #include "ContextBase.h" +#include "CpuTypes.h" #include #include @@ -8,13 +9,15 @@ #include #include +#include + namespace hiprt{ class CpuContext : public ContextBase { public: CpuContext( const hiprtContextCreationInput& input ); - ~CpuContext() override = default; + ~CpuContext() override; std::vector createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; @@ -108,13 +111,18 @@ class CpuContext : public ContextBase std::vector& functions, bool cache ) override; - struct CpuGeometryData* getCpuGeom( hiprtGeometry ) override + struct CpuGeometryData* getCpuGeom( hiprtGeometry geometry ) override { - return nullptr; + return reinterpret_cast( geometry ); }; struct CpuSceneData* getCpuScene( hiprtScene ) override { return nullptr; }; + + private: + void buildGeometryEntry( CpuGeometryData* data, const hiprtGeometryBuildInput& input, size_t index ); + + RTCDevice m_rtcDevice = nullptr; }; }//namespace hiprt \ No newline at end of file diff --git a/hiprt/impl/CpuTypes.h b/hiprt/impl/CpuTypes.h new file mode 100644 index 0000000..30e152c --- /dev/null +++ b/hiprt/impl/CpuTypes.h @@ -0,0 +1,20 @@ +#pragma once + +#include + +namespace hiprt +{ + +struct CpuGeometryData +{ + RTCDevice rtcDevice = nullptr; + RTCScene rtcScene = nullptr; +}; + +struct CpuSceneData +{ + RTCDevice rtcDevice = nullptr; + RTCScene rtcScene = nullptr; +}; + +} // namespace hiprt diff --git a/test/CpuGeometryTest.cpp b/test/CpuGeometryTest.cpp new file mode 100644 index 0000000..603c11d --- /dev/null +++ b/test/CpuGeometryTest.cpp @@ -0,0 +1,161 @@ +#include "CpuGeometryTest.h" + +#include +#include +#include + +void CpuGeometryTest::SetUp() +{ + hiprtContextCreationInput ctxInput{}; + ctxInput.deviceType = hiprtDeviceCPU; + ctxInput.numCpuThreads = 1; + + ASSERT_EQ( hiprtCreateContext( HIPRT_API_VERSION, ctxInput, m_context ), hiprtSuccess ) + << "Failed to create CPU context. Ensure hiprt.cpp routing is working."; + ASSERT_TRUE( m_context != nullptr ) << "Context pointer is null after creation."; +} + +void CpuGeometryTest::TearDown() +{ + if ( m_context != nullptr ) + EXPECT_EQ( hiprtDestroyContext( m_context ), hiprtSuccess ) + << "Failed to cleanly destroy CPU context."; +} + +hiprtGeometry CpuGeometryTest::buildSampleTriangle( + std::vector& vertices, + std::vector& indices, + hiprtGeometryBuildInput& geomInputOut ) +{ + vertices = { + { 0.0f, 0.0f, 0.0f }, + { 1.0f, 0.0f, 0.0f }, + { 0.0f, 1.0f, 0.0f } }; + indices = { 0, 1, 2 }; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = vertices.data(); + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = indices.data(); + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + geomInputOut = {}; + geomInputOut.type = hiprtPrimitiveTypeTriangleMesh; + geomInputOut.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + EXPECT_EQ( hiprtCreateGeometry( m_context, geomInputOut, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed."; + EXPECT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + EXPECT_EQ( + hiprtBuildGeometry( m_context, hiprtBuildOperationBuild, geomInputOut, opts, nullptr, nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry (Build) failed."; + return geom; +} + +namespace +{ +bool traceHit( hiprtGeometry geom, float ox, float oy, float oz, float dx, float dy, float dz ) +{ + auto* data = reinterpret_cast( geom ); + if ( data == nullptr || data->rtcScene == nullptr ) + return false; + + RTCRayHit rh{}; + rh.ray.org_x = ox; + rh.ray.org_y = oy; + rh.ray.org_z = oz; + rh.ray.dir_x = dx; + rh.ray.dir_y = dy; + rh.ray.dir_z = dz; + rh.ray.tnear = 0.0f; + rh.ray.tfar = 1e30f; + rh.ray.mask = static_cast( -1 ); + rh.ray.flags = 0; + rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; + + rtcIntersect1( data->rtcScene, &rh ); + return rh.hit.geomID != RTC_INVALID_GEOMETRY_ID; +} +} // namespace + +TEST_F( CpuGeometryTest, CreateAndDestroyGeometry ) +{ + std::vector vertices = { + { 0.0f, 0.0f, 0.0f }, + { 1.0f, 0.0f, 0.0f }, + { 0.0f, 1.0f, 0.0f } }; + std::vector indices = { 0, 1, 2 }; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = vertices.data(); + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = indices.data(); + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + hiprtGeometryBuildInput in{}; + in.type = hiprtPrimitiveTypeTriangleMesh; + in.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + ASSERT_EQ( hiprtCreateGeometry( m_context, in, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed on CPU context."; + ASSERT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ) + << "hiprtDestroyGeometry failed to release Embree scene memory."; +} + +TEST_F( CpuGeometryTest, BuildGeometryProducesHittableBvh ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometryBuildInput in{}; + hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); + ASSERT_TRUE( geom != nullptr ); + + EXPECT_TRUE( traceHit( geom, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Promien ze srodka trojkata powinien trafic w zbudowane BVH."; + EXPECT_FALSE( traceHit( geom, 2.0f, 2.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Promien poza trojkatem nie powinien trafiac."; + + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( CpuGeometryTest, UpdateGeometryRebuildsBvh ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometryBuildInput in{}; + hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); + ASSERT_TRUE( geom != nullptr ); + + EXPECT_FALSE( traceHit( geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Przed update BVH nie powinno obejmowac y=1.5."; + + vertices[2] = { 0.0f, 2.0f, 0.0f }; + in.primitive.triangleMesh.vertices = vertices.data(); + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + ASSERT_EQ( + hiprtBuildGeometry( m_context, hiprtBuildOperationUpdate, in, opts, nullptr, nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry (Update) failed to rebuild the scene."; + + EXPECT_TRUE( traceHit( geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Po update BVH powinno obejmowac y=1.5 (nowy wierzcholek y=2)."; + + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} diff --git a/test/CpuGeometryTest.h b/test/CpuGeometryTest.h new file mode 100644 index 0000000..9b2c907 --- /dev/null +++ b/test/CpuGeometryTest.h @@ -0,0 +1,19 @@ +#pragma once + +#include +#include +#include + +class CpuGeometryTest : public ::testing::Test +{ + protected: + void SetUp() override; + void TearDown() override; + + hiprtGeometry buildSampleTriangle( + std::vector& vertices, + std::vector& indices, + hiprtGeometryBuildInput& geomInputOut ); + + hiprtContext m_context = nullptr; +}; From 5d9f2f28455b509bd27484eaf0e5b5c878016cfe Mon Sep 17 00:00:00 2001 From: unknown Date: Thu, 11 Jun 2026 14:57:52 +0200 Subject: [PATCH 4/7] cpu: basic ray traversal implementation with Embree and GTest integration --- hiprt/hiprt_cpu.h | 107 ++++++++++++++++++ hiprt/impl/CpuContext.cpp | 191 +++++++++++++++++++++++++++++--- hiprt/impl/CpuContext.h | 5 +- test/CpuGeometryTest.cpp | 32 ++---- test/CpuSceneTest.cpp | 223 ++++++++++++++++++++++++++++++++++++++ test/CpuSceneTest.h | 18 +++ 6 files changed, 536 insertions(+), 40 deletions(-) create mode 100644 test/CpuSceneTest.cpp create mode 100644 test/CpuSceneTest.h diff --git a/hiprt/hiprt_cpu.h b/hiprt/hiprt_cpu.h index f6bbfa6..5f2f9e5 100644 --- a/hiprt/hiprt_cpu.h +++ b/hiprt/hiprt_cpu.h @@ -24,5 +24,112 @@ #pragma once +#include +#include +#include +#include +// --------------------------------------------------------------------------- +// Internal helpers – not for external consumption +// --------------------------------------------------------------------------- +namespace hiprt_cpu_detail +{ +inline RTCRayHit makeRTCRayHit( const hiprtRay& ray ) noexcept +{ + RTCRayHit rh{}; + rh.ray.org_x = ray.origin.x; + rh.ray.org_y = ray.origin.y; + rh.ray.org_z = ray.origin.z; + rh.ray.dir_x = ray.direction.x; + rh.ray.dir_y = ray.direction.y; + rh.ray.dir_z = ray.direction.z; + rh.ray.tnear = ray.minT; + rh.ray.tfar = ray.maxT; + rh.ray.mask = static_cast( -1 ); + rh.ray.flags = 0; + rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; + rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; + return rh; +} + +inline hiprtHit makeHiprtHit( const RTCRayHit& rh ) noexcept +{ + hiprtHit hit{}; + if ( rh.hit.geomID == RTC_INVALID_GEOMETRY_ID ) + return hit; // primID stays hiprtInvalidValue -> hasHit() == false + + hit.primID = rh.hit.primID; + hit.instanceID = ( rh.hit.instID[0] != RTC_INVALID_GEOMETRY_ID ) + ? rh.hit.instID[0] + : hiprtInvalidValue; + hit.uv.x = rh.hit.u; + hit.uv.y = rh.hit.v; + hit.normal.x = rh.hit.Ng_x; + hit.normal.y = rh.hit.Ng_y; + hit.normal.z = rh.hit.Ng_z; + hit.t = rh.ray.tfar; // Embree writes the hit distance into tfar + return hit; +} + +} // namespace hiprt_cpu_detail + +// --------------------------------------------------------------------------- +// Public CPU traversal objects +// --------------------------------------------------------------------------- + +/** \brief Host-side traversal that finds the closest hit in a hiprtGeometry. + * + * Mirrors the GPU class hiprtGeomTraversalClosest. Constructing the object + * immediately casts the handle to CpuGeometryData, retrieves the underlying + * RTCScene built by CpuContext, and calls rtcIntersect1. + * Call getNextHit() to obtain the result. + */ +class hiprtGeomTraversalClosestCPU +{ + public: + hiprtGeomTraversalClosestCPU( hiprtGeometry geom, const hiprtRay& ray ) noexcept + { + m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); + auto* data = reinterpret_cast( geom ); + if ( data != nullptr && data->rtcScene != nullptr ) + rtcIntersect1( data->rtcScene, &m_rayHit ); + } + + hiprtHit getNextHit() const noexcept + { + return hiprt_cpu_detail::makeHiprtHit( m_rayHit ); + } + + private: + RTCRayHit m_rayHit{}; +}; + +/** \brief Host-side traversal that finds the closest hit in a hiprtScene. + * + * Mirrors the GPU class hiprtSceneTraversalClosest. Constructing the object + * immediately casts the handle to CpuSceneData, retrieves the underlying + * RTCScene built by CpuContext (which already encodes the full instance + * hierarchy via RTC_GEOMETRY_TYPE_INSTANCE), and calls rtcIntersect1. + * Call getNextHit() to obtain the result, including instanceID populated + * from instID[0]. + */ +class hiprtSceneTraversalClosestCPU +{ + public: + hiprtSceneTraversalClosestCPU( hiprtScene scene, const hiprtRay& ray ) noexcept + { + m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); + auto* data = reinterpret_cast( scene ); + if ( data != nullptr && data->rtcScene != nullptr ) + rtcIntersect1( data->rtcScene, &m_rayHit ); + } + + hiprtHit getNextHit() const noexcept + { + return hiprt_cpu_detail::makeHiprtHit( m_rayHit ); + } + + private: + RTCRayHit m_rayHit{}; +}; diff --git a/hiprt/impl/CpuContext.cpp b/hiprt/impl/CpuContext.cpp index b253e08..82be7ea 100644 --- a/hiprt/impl/CpuContext.cpp +++ b/hiprt/impl/CpuContext.cpp @@ -1,5 +1,7 @@ #include "CpuContext.h" +#include + #include #include @@ -12,6 +14,23 @@ namespace std::cerr << "[CpuContext] not implemented: " << what << std::endl; throw std::runtime_error( what ); } + +void srtFrameToMatrix3x4( const hiprtFrameSRT& f, float ( &m )[3][4] ) +{ + const float4 q = qtFromAxisAngle( f.rotation ); + float Q[3][3]; + qtToRotationMatrix( q, Q ); + const float s[3] = { f.scale.x, f.scale.y, f.scale.z }; + for ( int i = 0; i < 3; ++i ) + { + m[i][0] = Q[i][0] * s[0]; + m[i][1] = Q[i][1] * s[1]; + m[i][2] = Q[i][2] * s[2]; + } + m[0][3] = f.translation.x; + m[1][3] = f.translation.y; + m[2][3] = f.translation.z; +} } // namespace CpuContext::CpuContext( const hiprtContextCreationInput& input ) @@ -209,31 +228,173 @@ std::vector CpuContext::compactGeometries( const std::vector CpuContext::createScenes( const std::vector&, const hiprtBuildOptions ) +std::vector CpuContext::createScenes( + const std::vector& buildInputs, const hiprtBuildOptions ) { - throwNotImplemented( "CpuContext::createScenes is not implemented yet." ); + std::vector scenes; + scenes.reserve( buildInputs.size() ); + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + CpuSceneData* data = new CpuSceneData(); + data->rtcDevice = m_rtcDevice; + data->rtcScene = rtcNewScene( m_rtcDevice ); + + if ( data->rtcScene == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + delete data; + std::cerr << "[CpuContext] rtcNewScene failed at index " << i + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewScene failed" ); + } + + scenes.push_back( reinterpret_cast( data ) ); + } + return scenes; } -void CpuContext::destroyScenes( const std::vector& ) { throwNotImplemented( "CpuContext::destroyScenes is not implemented yet." ); } +void CpuContext::destroyScenes( const std::vector& scenes ) +{ + for ( hiprtScene scene : scenes ) + { + CpuSceneData* data = getCpuScene( scene ); + if ( data == nullptr ) continue; + + if ( data->rtcScene != nullptr ) + { + rtcReleaseScene( data->rtcScene ); + data->rtcScene = nullptr; + } + delete data; + } +} + +void CpuContext::buildSceneEntry( CpuSceneData* data, const hiprtSceneBuildInput& input, size_t index ) +{ + if ( data == nullptr || data->rtcScene == nullptr ) + throw std::runtime_error( "CpuContext::buildSceneEntry: invalid scene handle" ); + + const hiprtInstance* instances = + reinterpret_cast( input.instances ); + const hiprtTransformHeader* headers = + reinterpret_cast( input.instanceTransformHeaders ); + + if ( instances == nullptr && input.instanceCount != 0 ) + throw std::runtime_error( "CpuContext::buildSceneEntry: null instances array" ); + + for ( uint32_t i = 0; i < input.instanceCount; ++i ) + { + const hiprtInstance& inst = instances[i]; + + RTCScene child = nullptr; + if ( inst.type == hiprtInstanceTypeGeometry ) + { + CpuGeometryData* g = getCpuGeom( inst.geometry ); + child = ( g != nullptr ) ? g->rtcScene : nullptr; + } + else + { + CpuSceneData* s = getCpuScene( inst.scene ); + child = ( s != nullptr ) ? s->rtcScene : nullptr; + } + + if ( child == nullptr ) + throw std::runtime_error( "CpuContext::buildSceneEntry: instance has null child scene" ); + + RTCGeometry instGeom = rtcNewGeometry( m_rtcDevice, RTC_GEOMETRY_TYPE_INSTANCE ); + if ( instGeom == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + std::cerr << "[CpuContext] rtcNewGeometry(INSTANCE) failed at scene " << index + << " instance " << i << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewGeometry(INSTANCE) failed" ); + } + + rtcSetGeometryInstancedScene( instGeom, child ); + + float xfm[3][4]; + if ( input.instanceFrames == nullptr ) + { + xfm[0][0] = 1.0f; xfm[0][1] = 0.0f; xfm[0][2] = 0.0f; xfm[0][3] = 0.0f; + xfm[1][0] = 0.0f; xfm[1][1] = 1.0f; xfm[1][2] = 0.0f; xfm[1][3] = 0.0f; + xfm[2][0] = 0.0f; xfm[2][1] = 0.0f; xfm[2][2] = 1.0f; xfm[2][3] = 0.0f; + } + else + { + const uint32_t frameIndex = ( headers != nullptr ) ? headers[i].frameIndex : i; + + if ( input.frameType == hiprtFrameTypeMatrix ) + { + const hiprtFrameMatrix* mats = + reinterpret_cast( input.instanceFrames ); + for ( int r = 0; r < 3; ++r ) + for ( int c = 0; c < 4; ++c ) + xfm[r][c] = mats[frameIndex].matrix[r][c]; + } + else + { + const hiprtFrameSRT* srts = + reinterpret_cast( input.instanceFrames ); + srtFrameToMatrix3x4( srts[frameIndex], xfm ); + } + } + + rtcSetGeometryTransform( instGeom, 0, RTC_FORMAT_FLOAT3X4_ROW_MAJOR, &xfm[0][0] ); + rtcCommitGeometry( instGeom ); + rtcAttachGeometry( data->rtcScene, instGeom ); + rtcReleaseGeometry( instGeom ); + } + + rtcCommitScene( data->rtcScene ); +} void CpuContext::buildScenes( - const std::vector&, - const hiprtBuildOptions, - hiprtDevicePtr, - oroStream, - std::vector& ) + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) { - throwNotImplemented( "CpuContext::buildScenes is not implemented yet." ); + (void)buildOptions; + (void)temporaryBuffer; + (void)stream; + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + buildSceneEntry( getCpuScene( reinterpret_cast( buffers[i] ) ), buildInputs[i], i ); } void CpuContext::updateScenes( - const std::vector&, - const hiprtBuildOptions, - hiprtDevicePtr, - oroStream, - std::vector& ) + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) { - throwNotImplemented( "CpuContext::updateScenes is not implemented yet." ); + (void)buildOptions; + (void)temporaryBuffer; + (void)stream; + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + CpuSceneData* data = getCpuScene( reinterpret_cast( buffers[i] ) ); + if ( data == nullptr ) + throw std::runtime_error( "CpuContext::updateScenes: invalid scene handle" ); + + if ( data->rtcScene != nullptr ) + rtcReleaseScene( data->rtcScene ); + + data->rtcScene = rtcNewScene( m_rtcDevice ); + if ( data->rtcScene == nullptr ) + { + const RTCError err = rtcGetDeviceError( m_rtcDevice ); + std::cerr << "[CpuContext] rtcNewScene failed at index " << i + << ", RTCError=" << static_cast( err ) << std::endl; + throw std::runtime_error( "rtcNewScene failed" ); + } + + buildSceneEntry( data, buildInputs[i], i ); + } } size_t CpuContext::getScenesBuildTempBufferSize( const std::vector&, const hiprtBuildOptions ) diff --git a/hiprt/impl/CpuContext.h b/hiprt/impl/CpuContext.h index 66698af..9a86e34 100644 --- a/hiprt/impl/CpuContext.h +++ b/hiprt/impl/CpuContext.h @@ -115,13 +115,14 @@ class CpuContext : public ContextBase { return reinterpret_cast( geometry ); }; - struct CpuSceneData* getCpuScene( hiprtScene ) override + struct CpuSceneData* getCpuScene( hiprtScene scene ) override { - return nullptr; + return reinterpret_cast( scene ); }; private: void buildGeometryEntry( CpuGeometryData* data, const hiprtGeometryBuildInput& input, size_t index ); + void buildSceneEntry( CpuSceneData* data, const hiprtSceneBuildInput& input, size_t index ); RTCDevice m_rtcDevice = nullptr; }; diff --git a/test/CpuGeometryTest.cpp b/test/CpuGeometryTest.cpp index 603c11d..3b94de4 100644 --- a/test/CpuGeometryTest.cpp +++ b/test/CpuGeometryTest.cpp @@ -1,8 +1,6 @@ #include "CpuGeometryTest.h" -#include -#include -#include +#include void CpuGeometryTest::SetUp() { @@ -63,26 +61,14 @@ namespace { bool traceHit( hiprtGeometry geom, float ox, float oy, float oz, float dx, float dy, float dz ) { - auto* data = reinterpret_cast( geom ); - if ( data == nullptr || data->rtcScene == nullptr ) - return false; - - RTCRayHit rh{}; - rh.ray.org_x = ox; - rh.ray.org_y = oy; - rh.ray.org_z = oz; - rh.ray.dir_x = dx; - rh.ray.dir_y = dy; - rh.ray.dir_z = dz; - rh.ray.tnear = 0.0f; - rh.ray.tfar = 1e30f; - rh.ray.mask = static_cast( -1 ); - rh.ray.flags = 0; - rh.hit.geomID = RTC_INVALID_GEOMETRY_ID; - rh.hit.instID[0] = RTC_INVALID_GEOMETRY_ID; - - rtcIntersect1( data->rtcScene, &rh ); - return rh.hit.geomID != RTC_INVALID_GEOMETRY_ID; + hiprtRay ray{}; + ray.origin = { ox, oy, oz }; + ray.minT = 0.0f; + ray.direction = { dx, dy, dz }; + ray.maxT = 1e30f; + + hiprtGeomTraversalClosestCPU tr( geom, ray ); + return tr.getNextHit().hasHit(); } } // namespace diff --git a/test/CpuSceneTest.cpp b/test/CpuSceneTest.cpp new file mode 100644 index 0000000..df36e67 --- /dev/null +++ b/test/CpuSceneTest.cpp @@ -0,0 +1,223 @@ +#include "CpuSceneTest.h" + +#include + +void CpuSceneTest::SetUp() +{ + hiprtContextCreationInput ctxInput{}; + ctxInput.deviceType = hiprtDeviceCPU; + ctxInput.numCpuThreads = 1; + ASSERT_EQ( hiprtCreateContext( HIPRT_API_VERSION, ctxInput, m_context ), hiprtSuccess ) + << "Failed to create CPU context."; + ASSERT_TRUE( m_context != nullptr ) << "Context pointer is null after creation."; +} + +void CpuSceneTest::TearDown() +{ + if ( m_context != nullptr ) + EXPECT_EQ( hiprtDestroyContext( m_context ), hiprtSuccess ) + << "Failed to cleanly destroy CPU context."; +} + +hiprtGeometry CpuSceneTest::buildTriangleGeometry( + std::vector& vertices, + std::vector& indices ) +{ + vertices = { + { 0.0f, 0.0f, 0.0f }, + { 1.0f, 0.0f, 0.0f }, + { 0.0f, 1.0f, 0.0f } }; + indices = { 0, 1, 2 }; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = vertices.data(); + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = indices.data(); + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + hiprtGeometryBuildInput in{}; + in.type = hiprtPrimitiveTypeTriangleMesh; + in.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + EXPECT_EQ( hiprtCreateGeometry( m_context, in, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed."; + EXPECT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + EXPECT_EQ( + hiprtBuildGeometry( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry failed."; + return geom; +} + +namespace +{ +hiprtFrameSRT makeSrt( float tx, float ty, float tz ) +{ + hiprtFrameSRT f{}; + f.rotation = { 0.0f, 0.0f, 0.0f, 0.0f }; + f.scale = { 1.0f, 1.0f, 1.0f }; + f.translation = { tx, ty, tz }; + f.time = 0.0f; + return f; +} + +bool traceScene( hiprtScene scene, float ox, float oy, float oz, float dx, float dy, float dz ) +{ + hiprtRay ray{}; + ray.origin = { ox, oy, oz }; + ray.minT = 0.0f; + ray.direction = { dx, dy, dz }; + ray.maxT = 1e30f; + + hiprtSceneTraversalClosestCPU tr( scene, ray ); + const hiprtHit hit = tr.getNextHit(); + return hit.instanceID != hiprtInvalidValue; +} +} // namespace + +TEST_F( CpuSceneTest, CreateAndDestroyScene ) +{ + hiprtSceneBuildInput in{}; + in.instanceCount = 0; + in.frameType = hiprtFrameTypeSRT; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ) + << "hiprtCreateScene failed on CPU context."; + ASSERT_TRUE( scene != nullptr ) << "Returned scene handle is null."; + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ) + << "hiprtDestroyScene failed to release Embree scene memory."; +} + +TEST_F( CpuSceneTest, BuildSceneWithSingleInstanceIsHittable ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 0.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + in.instances = instances.data(); + in.instanceCount = 1; + in.instanceFrames = frames.data(); + in.frameCount = 1; + in.frameType = hiprtFrameTypeSRT; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Promien w obszar instancji powinien trafic."; + EXPECT_FALSE( traceScene( scene, 5.0f, 5.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Promien poza instancja nie powinien trafic."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( CpuSceneTest, BuildSceneWithTranslatedInstance ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 3.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + in.instances = instances.data(); + in.instanceCount = 1; + in.instanceFrames = frames.data(); + in.frameCount = 1; + in.frameType = hiprtFrameTypeSRT; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Po translacji +3X instancja powinna byc trafiona w x=3.25."; + EXPECT_FALSE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Oryginalna pozycja x=0.25 nie powinna byc juz trafiona."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( CpuSceneTest, UpdateSceneRebuildsInstances ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 0.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + in.instances = instances.data(); + in.instanceCount = 1; + in.instanceFrames = frames.data(); + in.frameCount = 1; + in.frameType = hiprtFrameTypeSRT; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Przed update trafienie w x=0.25."; + + frames[0] = makeSrt( 3.0f, 0.0f, 0.0f ); + in.instanceFrames = frames.data(); + ASSERT_EQ( + hiprtBuildScene( m_context, hiprtBuildOperationUpdate, in, opts, nullptr, nullptr, scene ), + hiprtSuccess ) + << "hiprtBuildScene (Update) failed."; + + EXPECT_FALSE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Po update instancja przesunieta, stara pozycja pusta."; + EXPECT_TRUE( traceScene( scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Po update trafienie w nowej pozycji x=3.25."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} diff --git a/test/CpuSceneTest.h b/test/CpuSceneTest.h new file mode 100644 index 0000000..290ab57 --- /dev/null +++ b/test/CpuSceneTest.h @@ -0,0 +1,18 @@ +#pragma once + +#include +#include +#include + +class CpuSceneTest : public ::testing::Test +{ + protected: + void SetUp() override; + void TearDown() override; + + hiprtGeometry buildTriangleGeometry( + std::vector& vertices, + std::vector& indices ); + + hiprtContext m_context = nullptr; +}; From a79d354b00a1a3b9f1020641e9f0291f0d81dfcb Mon Sep 17 00:00:00 2001 From: unknown Date: Sun, 14 Jun 2026 15:03:14 +0200 Subject: [PATCH 5/7] hybrid: HybridContext implementation and global handle registry for seamless CPU/GPU traversal --- hiprt/hiprt.h.in | 22 +++ hiprt/hiprt_cpu.h | 9 +- hiprt/hiprt_types.h | 1 + hiprt/impl/CpuContext.cpp | 13 +- hiprt/impl/CpuDataRegistry.h | 29 +++ hiprt/impl/HybridContext.cpp | 373 +++++++++++++++++++++++++++++++++++ hiprt/impl/HybridContext.h | 148 ++++++++++++++ hiprt/impl/hiprt.cpp | 77 +++++++- 8 files changed, 665 insertions(+), 7 deletions(-) create mode 100644 hiprt/impl/CpuDataRegistry.h create mode 100644 hiprt/impl/HybridContext.cpp create mode 100644 hiprt/impl/HybridContext.h diff --git a/hiprt/hiprt.h.in b/hiprt/hiprt.h.in index b83ae53..392e14d 100644 --- a/hiprt/hiprt.h.in +++ b/hiprt/hiprt.h.in @@ -600,6 +600,28 @@ HIPRT_API hiprtError hiprtSetCacheDirPath( hiprtContext context, const char* pat */ HIPRT_API hiprtError hiprtSetLogLevel( hiprtContext context, hiprtLogLevel level ); +/** \brief Returns the Embree CpuGeometryData pointer associated with a geometry handle. + * + * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceHybrid) contexts. + * Returns nullptr if the handle has no CPU mirror (pure GPU context). + * The returned pointer is opaque; cast it to hiprt::CpuGeometryData* in hiprt_cpu.h. + * + * \param geometry The geometry handle. + * \return Pointer to the internal CpuGeometryData, or nullptr. + */ +HIPRT_API void* hiprtGetInternalCpuDataFromGeometry( hiprtGeometry geometry ); + +/** \brief Returns the Embree CpuSceneData pointer associated with a scene handle. + * + * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceHybrid) contexts. + * Returns nullptr if the handle has no CPU mirror (pure GPU context). + * The returned pointer is opaque; cast it to hiprt::CpuSceneData* in hiprt_cpu.h. + * + * \param scene The scene handle. + * \return Pointer to the internal CpuSceneData, or nullptr. + */ +HIPRT_API void* hiprtGetInternalCpuDataFromScene( hiprtScene scene ); + #ifdef __cplusplus } #endif diff --git a/hiprt/hiprt_cpu.h b/hiprt/hiprt_cpu.h index 5f2f9e5..263516e 100644 --- a/hiprt/hiprt_cpu.h +++ b/hiprt/hiprt_cpu.h @@ -24,6 +24,7 @@ #pragma once +#include // HIPRT_API + hiprtGetInternalCpuDataFromGeometry/Scene #include #include #include @@ -90,8 +91,8 @@ class hiprtGeomTraversalClosestCPU public: hiprtGeomTraversalClosestCPU( hiprtGeometry geom, const hiprtRay& ray ) noexcept { - m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); - auto* data = reinterpret_cast( geom ); + m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromGeometry( geom ) ); if ( data != nullptr && data->rtcScene != nullptr ) rtcIntersect1( data->rtcScene, &m_rayHit ); } @@ -119,8 +120,8 @@ class hiprtSceneTraversalClosestCPU public: hiprtSceneTraversalClosestCPU( hiprtScene scene, const hiprtRay& ray ) noexcept { - m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); - auto* data = reinterpret_cast( scene ); + m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromScene( scene ) ); if ( data != nullptr && data->rtcScene != nullptr ) rtcIntersect1( data->rtcScene, &m_rayHit ); } diff --git a/hiprt/hiprt_types.h b/hiprt/hiprt_types.h index 82f97a5..b70b220 100644 --- a/hiprt/hiprt_types.h +++ b/hiprt/hiprt_types.h @@ -314,6 +314,7 @@ enum hiprtDeviceType : uint32_t hiprtDeviceAMD = 1 << 0, hiprtDeviceNVIDIA = 1 << 1, hiprtDeviceCPU = 1 << 2, + hiprtDeviceHybrid = 1 << 3, }; /** \brief Context creation input. diff --git a/hiprt/impl/CpuContext.cpp b/hiprt/impl/CpuContext.cpp index 82be7ea..d2f71cd 100644 --- a/hiprt/impl/CpuContext.cpp +++ b/hiprt/impl/CpuContext.cpp @@ -1,4 +1,5 @@ #include "CpuContext.h" +#include #include @@ -82,7 +83,9 @@ std::vector CpuContext::createGeometries( throw std::runtime_error( "rtcNewScene failed" ); } - geometries.push_back( reinterpret_cast( data ) ); + hiprtGeometry handle = reinterpret_cast( data ); + registerCpuGeom( handle, data ); + geometries.push_back( handle ); } return geometries; @@ -95,6 +98,8 @@ void CpuContext::destroyGeometries( const std::vector& geometries CpuGeometryData* data = reinterpret_cast( geometry ); if ( data == nullptr ) continue; + unregisterCpuGeom( geometry ); + if ( data->rtcScene != nullptr ) { rtcReleaseScene( data->rtcScene ); @@ -249,7 +254,9 @@ std::vector CpuContext::createScenes( throw std::runtime_error( "rtcNewScene failed" ); } - scenes.push_back( reinterpret_cast( data ) ); + hiprtScene handle = reinterpret_cast( data ); + registerCpuScene( handle, data ); + scenes.push_back( handle ); } return scenes; } @@ -261,6 +268,8 @@ void CpuContext::destroyScenes( const std::vector& scenes ) CpuSceneData* data = getCpuScene( scene ); if ( data == nullptr ) continue; + unregisterCpuScene( scene ); + if ( data->rtcScene != nullptr ) { rtcReleaseScene( data->rtcScene ); diff --git a/hiprt/impl/CpuDataRegistry.h b/hiprt/impl/CpuDataRegistry.h new file mode 100644 index 0000000..b87aa7e --- /dev/null +++ b/hiprt/impl/CpuDataRegistry.h @@ -0,0 +1,29 @@ +#pragma once + +#include + +namespace hiprt +{ +struct CpuGeometryData; +struct CpuSceneData; + +/** + * Global CPU-data registry. + * + * GPU pointers (oroDeviceptr cast to hiprtGeometry/hiprtScene) are unique + * process-wide, so a single flat map is sufficient. Pure-CPU handles are the + * CpuGeometryData/CpuSceneData pointers themselves — they are also unique. + * Both can therefore share the same table without ambiguity. + * + * All functions are thread-safe via an internal mutex. + */ + +void registerCpuGeom( hiprtGeometry handle, CpuGeometryData* data ); +void unregisterCpuGeom( hiprtGeometry handle ); +CpuGeometryData* lookupCpuGeom( hiprtGeometry handle ); + +void registerCpuScene( hiprtScene handle, CpuSceneData* data ); +void unregisterCpuScene( hiprtScene handle ); +CpuSceneData* lookupCpuScene( hiprtScene handle ); + +} // namespace hiprt diff --git a/hiprt/impl/HybridContext.cpp b/hiprt/impl/HybridContext.cpp new file mode 100644 index 0000000..e0670d3 --- /dev/null +++ b/hiprt/impl/HybridContext.cpp @@ -0,0 +1,373 @@ +#include +#include + +#include +#include +#include + +namespace hiprt +{ + +// --------------------------------------------------------------------------- +// Construction / destruction +// --------------------------------------------------------------------------- + +HybridContext::HybridContext( const hiprtContextCreationInput& input ) + : m_gpu( input ), m_cpu( input ) +{ + std::cerr << "[HybridContext] constructed — GPU BVH + Embree CPU mirror active.\n"; +} + +HybridContext::~HybridContext() +{ + // Both sub-contexts clean up their own resources. + // Maps are empty at this point (destroy calls already handled them). +} + +// --------------------------------------------------------------------------- +// Geometry +// --------------------------------------------------------------------------- + +std::vector HybridContext::createGeometries( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +{ + std::vector gpuHandles = m_gpu.createGeometries( buildInputs, buildOptions ); + std::vector cpuHandles = m_cpu.createGeometries( buildInputs, buildOptions ); + + for ( size_t i = 0; i < gpuHandles.size(); ++i ) + { + auto* cpuData = reinterpret_cast( cpuHandles[i] ); + m_gpuToCpuGeom[gpuHandles[i]] = cpuHandles[i]; + registerCpuGeom( gpuHandles[i], cpuData ); + } + + return gpuHandles; // return GPU handles; CPU side is hidden in the maps. +} + +void HybridContext::destroyGeometries( const std::vector& gpuGeometries ) +{ + std::vector cpuHandles; + cpuHandles.reserve( gpuGeometries.size() ); + + for ( hiprtGeometry g : gpuGeometries ) + { + auto it = m_gpuToCpuGeom.find( g ); + if ( it != m_gpuToCpuGeom.end() ) + { + cpuHandles.push_back( it->second ); + unregisterCpuGeom( g ); + m_gpuToCpuGeom.erase( it ); + } + } + + m_gpu.destroyGeometries( gpuGeometries ); + if ( !cpuHandles.empty() ) + m_cpu.destroyGeometries( cpuHandles ); +} + +void HybridContext::dispatchCpuGeomBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ) +{ + // Translate GPU buffers -> CPU buffers for CpuContext + std::vector cpuBuffers; + cpuBuffers.reserve( gpuBuffers.size() ); + for ( auto gp : gpuBuffers ) + { + auto it = m_gpuToCpuGeom.find( reinterpret_cast( gp ) ); + if ( it == m_gpuToCpuGeom.end() ) + throw std::runtime_error( "HybridContext: GPU geometry buffer not found in map" ); + cpuBuffers.push_back( reinterpret_cast( it->second ) ); + } + + if ( isUpdate ) + m_cpu.updateGeometries( buildInputs, {}, nullptr, nullptr, cpuBuffers ); + else + m_cpu.buildGeometries( buildInputs, {}, nullptr, nullptr, cpuBuffers ); +} + +void HybridContext::buildGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + m_gpu.buildGeometries( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + dispatchCpuGeomBuild( buildInputs, buffers, false ); +} + +void HybridContext::updateGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + m_gpu.updateGeometries( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + dispatchCpuGeomBuild( buildInputs, buffers, true ); +} + +size_t HybridContext::getGeometriesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +{ + return m_gpu.getGeometriesBuildTempBufferSize( buildInputs, buildOptions ); +} + +std::vector HybridContext::compactGeometries( + const std::vector& geometries, oroStream stream ) +{ + // GPU compaction only in this pass; CPU mirror is unaffected. + return m_gpu.compactGeometries( geometries, stream ); +} + +// --------------------------------------------------------------------------- +// Scene +// --------------------------------------------------------------------------- + +std::vector HybridContext::createScenes( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +{ + std::vector gpuHandles = m_gpu.createScenes( buildInputs, buildOptions ); + std::vector cpuHandles = m_cpu.createScenes( buildInputs, buildOptions ); + + for ( size_t i = 0; i < gpuHandles.size(); ++i ) + { + auto* cpuData = reinterpret_cast( cpuHandles[i] ); + m_gpuToCpuScene[gpuHandles[i]] = cpuHandles[i]; + registerCpuScene( gpuHandles[i], cpuData ); + } + + return gpuHandles; +} + +void HybridContext::destroyScenes( const std::vector& gpuScenes ) +{ + std::vector cpuHandles; + cpuHandles.reserve( gpuScenes.size() ); + + for ( hiprtScene s : gpuScenes ) + { + auto it = m_gpuToCpuScene.find( s ); + if ( it != m_gpuToCpuScene.end() ) + { + cpuHandles.push_back( it->second ); + unregisterCpuScene( s ); + m_gpuToCpuScene.erase( it ); + } + } + + m_gpu.destroyScenes( gpuScenes ); + if ( !cpuHandles.empty() ) + m_cpu.destroyScenes( cpuHandles ); +} + +void HybridContext::dispatchCpuSceneBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ) +{ + // Build CPU-side buffers and remap instance handles (GPU->CPU). + // We deep-copy each hiprtSceneBuildInput and patch its instances array. + std::vector cpuBuffers; + std::vector cpuInputs; + // Storage for the remapped instances arrays (keep alive during the call). + std::vector> remappedInstances( buildInputs.size() ); + + cpuBuffers.reserve( gpuBuffers.size() ); + cpuInputs.reserve( buildInputs.size() ); + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + // Translate GPU buffer to CPU buffer. + auto it = m_gpuToCpuScene.find( reinterpret_cast( gpuBuffers[i] ) ); + if ( it == m_gpuToCpuScene.end() ) + throw std::runtime_error( "HybridContext: GPU scene buffer not found in map" ); + cpuBuffers.push_back( reinterpret_cast( it->second ) ); + + // Copy the build input and remap instances. + hiprtSceneBuildInput cpuIn = buildInputs[i]; + + if ( buildInputs[i].instanceCount > 0 && buildInputs[i].instances != nullptr ) + { + const auto* srcInst = reinterpret_cast( buildInputs[i].instances ); + auto& dst = remappedInstances[i]; + dst.resize( buildInputs[i].instanceCount ); + + for ( uint32_t j = 0; j < buildInputs[i].instanceCount; ++j ) + { + dst[j] = srcInst[j]; + if ( srcInst[j].type == hiprtInstanceTypeGeometry ) + { + auto gIt = m_gpuToCpuGeom.find( srcInst[j].geometry ); + if ( gIt != m_gpuToCpuGeom.end() ) + dst[j].geometry = gIt->second; + } + else + { + auto sIt = m_gpuToCpuScene.find( srcInst[j].scene ); + if ( sIt != m_gpuToCpuScene.end() ) + dst[j].scene = sIt->second; + } + } + cpuIn.instances = dst.data(); + } + + cpuInputs.push_back( cpuIn ); + } + + if ( isUpdate ) + m_cpu.updateScenes( cpuInputs, {}, nullptr, nullptr, cpuBuffers ); + else + m_cpu.buildScenes( cpuInputs, {}, nullptr, nullptr, cpuBuffers ); +} + +void HybridContext::buildScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + m_gpu.buildScenes( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + dispatchCpuSceneBuild( buildInputs, buffers, false ); +} + +void HybridContext::updateScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + m_gpu.updateScenes( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + dispatchCpuSceneBuild( buildInputs, buffers, true ); +} + +size_t HybridContext::getScenesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) +{ + return m_gpu.getScenesBuildTempBufferSize( buildInputs, buildOptions ); +} + +std::vector HybridContext::compactScenes( + const std::vector& scenes, oroStream stream ) +{ + return m_gpu.compactScenes( scenes, stream ); +} + +// --------------------------------------------------------------------------- +// GPU-authoritative delegators +// --------------------------------------------------------------------------- + +hiprtFuncTable HybridContext::createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) +{ + return m_gpu.createFuncTable( numGeomTypes, numRayTypes ); +} + +void HybridContext::setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) +{ + m_gpu.setFuncTable( funcTable, geomType, rayType, set ); +} + +void HybridContext::destroyFuncTable( hiprtFuncTable funcTable ) { m_gpu.destroyFuncTable( funcTable ); } + +void HybridContext::createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) +{ + m_gpu.createGlobalStackBuffer( input, stackBufferOut ); +} + +void HybridContext::destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) +{ + m_gpu.destroyGlobalStackBuffer( stackBuffer ); +} + +void HybridContext::saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) +{ + m_gpu.saveGeometry( inGeometry, filename ); +} + +hiprtGeometry HybridContext::loadGeometry( const std::string& filename ) +{ + return m_gpu.loadGeometry( filename ); +} + +void HybridContext::saveScene( hiprtScene inScene, const std::string& filename ) +{ + m_gpu.saveScene( inScene, filename ); +} + +hiprtScene HybridContext::loadScene( const std::string& filename ) +{ + return m_gpu.loadScene( filename ); +} + +void HybridContext::exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) +{ + m_gpu.exportGeometryAabb( inGeometry, outAabbMin, outAabbMax ); +} + +void HybridContext::exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) +{ + m_gpu.exportSceneAabb( inScene, outAabbMin, outAabbMax ); +} + +void HybridContext::buildKernels( + const std::vector& funcNames, + const std::string& src, + const std::filesystem::path& moduleName, + std::vector& headers, + std::vector& includeNames, + std::vector& options, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + oroModule& module, + bool cache ) +{ + m_gpu.buildKernels( + funcNames, src, moduleName, headers, includeNames, options, + numGeomTypes, numRayTypes, funcNameSets, functions, module, cache ); +} + +void HybridContext::buildKernelsFromBitcode( + const std::vector& funcNames, + const std::filesystem::path& moduleName, + const std::string_view bitcodeBinary, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + bool cache ) +{ + m_gpu.buildKernelsFromBitcode( funcNames, moduleName, bitcodeBinary, numGeomTypes, numRayTypes, funcNameSets, functions, cache ); +} + +void HybridContext::setCacheDir( const std::filesystem::path& path ) { m_gpu.setCacheDir( path ); } + +void HybridContext::setLogLevel( hiprtLogLevel level ) +{ + m_gpu.setLogLevel( level ); +} + +// --------------------------------------------------------------------------- +// CPU data accessors +// --------------------------------------------------------------------------- + +CpuGeometryData* HybridContext::getCpuGeom( hiprtGeometry geometry ) +{ + auto it = m_gpuToCpuGeom.find( geometry ); + if ( it == m_gpuToCpuGeom.end() ) return nullptr; + return reinterpret_cast( it->second ); +} + +CpuSceneData* HybridContext::getCpuScene( hiprtScene scene ) +{ + auto it = m_gpuToCpuScene.find( scene ); + if ( it == m_gpuToCpuScene.end() ) return nullptr; + return reinterpret_cast( it->second ); +} + +} // namespace hiprt diff --git a/hiprt/impl/HybridContext.h b/hiprt/impl/HybridContext.h new file mode 100644 index 0000000..5206b33 --- /dev/null +++ b/hiprt/impl/HybridContext.h @@ -0,0 +1,148 @@ +#pragma once + +#include +#include +#include +#include + +#include + +namespace hiprt +{ + +/** + * HybridContext — composes GpuContext + CpuContext. + * + * Every build operation is performed twice: + * 1. On the GPU (via GpuContext) — the GPU handle is returned to the caller. + * 2. On the CPU (via CpuContext / Embree) — the CpuGeometryData / CpuSceneData ptr + * is stored in the instance maps and registered in the global CpuDataRegistry + * so that hiprt_cpu.h traversal classes can resolve it by the GPU handle. + * + * This preserves 100 % binary compatibility with existing GPU-only code: the + * opaque handles visible to the application remain GPU device pointers. + */ +class HybridContext : public ContextBase +{ + public: + explicit HybridContext( const hiprtContextCreationInput& input ); + ~HybridContext() override; + + // ----- Geometry -------------------------------------------------------- + std::vector + createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + void destroyGeometries( const std::vector& geometries ) override; + + void buildGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + void updateGeometries( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + size_t getGeometriesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + std::vector compactGeometries( const std::vector& geometries, oroStream stream ) override; + + // ----- Scene ----------------------------------------------------------- + std::vector + createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + void destroyScenes( const std::vector& scenes ) override; + + void buildScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + void updateScenes( + const std::vector& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) override; + + size_t getScenesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; + + std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; + + // ----- GPU-authoritative operations (delegate to GpuContext) ----------- + hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) override; + void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) override; + void destroyFuncTable( hiprtFuncTable funcTable ) override; + + void createGlobalStackBuffer( const hiprtGlobalStackBufferInput& input, hiprtGlobalStackBuffer& stackBufferOut ) override; + void destroyGlobalStackBuffer( hiprtGlobalStackBuffer stackBuffer ) override; + + void saveGeometry( hiprtGeometry inGeometry, const std::string& filename ) override; + hiprtGeometry loadGeometry( const std::string& filename ) override; + void saveScene( hiprtScene inScene, const std::string& filename ) override; + hiprtScene loadScene( const std::string& filename ) override; + + void exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) override; + void exportSceneAabb( hiprtScene inScene, float3& outAabbMin, float3& outAabbMax ) override; + + void buildKernels( + const std::vector& funcNames, + const std::string& src, + const std::filesystem::path& moduleName, + std::vector& headers, + std::vector& includeNames, + std::vector& options, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + oroModule& module, + bool cache ) override; + + void buildKernelsFromBitcode( + const std::vector& funcNames, + const std::filesystem::path& moduleName, + const std::string_view bitcodeBinary, + uint32_t numGeomTypes, + uint32_t numRayTypes, + const std::vector& funcNameSets, + std::vector& functions, + bool cache ) override; + + void setCacheDir( const std::filesystem::path& path ) override; + void setLogLevel( hiprtLogLevel level ) override; + + // ----- CPU-data accessors (used by hiprt_cpu.h via registry) ----------- + CpuGeometryData* getCpuGeom( hiprtGeometry geometry ) override; + CpuSceneData* getCpuScene( hiprtScene scene ) override; + + private: + // Translate GPU geometry buffers to their CPU counterparts and call fn. + void dispatchCpuGeomBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ); + + // Translate GPU scene buffers (remapping instance handles) and call fn. + void dispatchCpuSceneBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ); + + GpuContext m_gpu; + CpuContext m_cpu; + + std::unordered_map m_gpuToCpuGeom; // GPU handle -> CPU handle + std::unordered_map m_gpuToCpuScene; // GPU handle -> CPU handle +}; + +} // namespace hiprt diff --git a/hiprt/impl/hiprt.cpp b/hiprt/impl/hiprt.cpp index 770f7a2..b68fd88 100644 --- a/hiprt/impl/hiprt.cpp +++ b/hiprt/impl/hiprt.cpp @@ -26,13 +26,71 @@ #include #include #include +#include +#include #include #include #include +#include #include +#include +#include + using namespace hiprt; +// --------------------------------------------------------------------------- +// Global CPU-data registry implementation +// --------------------------------------------------------------------------- +namespace hiprt +{ +namespace +{ +std::mutex g_cpuGeomMutex; +std::unordered_map g_cpuGeomMap; +std::mutex g_cpuSceneMutex; +std::unordered_map g_cpuSceneMap; +} // namespace + +void registerCpuGeom( hiprtGeometry handle, CpuGeometryData* data ) +{ + std::lock_guard lk( g_cpuGeomMutex ); + g_cpuGeomMap[static_cast( handle )] = static_cast( data ); +} + +void unregisterCpuGeom( hiprtGeometry handle ) +{ + std::lock_guard lk( g_cpuGeomMutex ); + g_cpuGeomMap.erase( static_cast( handle ) ); +} + +CpuGeometryData* lookupCpuGeom( hiprtGeometry handle ) +{ + std::lock_guard lk( g_cpuGeomMutex ); + auto it = g_cpuGeomMap.find( static_cast( handle ) ); + return ( it != g_cpuGeomMap.end() ) ? static_cast( it->second ) : nullptr; +} + +void registerCpuScene( hiprtScene handle, CpuSceneData* data ) +{ + std::lock_guard lk( g_cpuSceneMutex ); + g_cpuSceneMap[static_cast( handle )] = static_cast( data ); +} + +void unregisterCpuScene( hiprtScene handle ) +{ + std::lock_guard lk( g_cpuSceneMutex ); + g_cpuSceneMap.erase( static_cast( handle ) ); +} + +CpuSceneData* lookupCpuScene( hiprtScene handle ) +{ + std::lock_guard lk( g_cpuSceneMutex ); + auto it = g_cpuSceneMap.find( static_cast( handle ) ); + return ( it != g_cpuSceneMap.end() ) ? static_cast( it->second ) : nullptr; +} +} // namespace hiprt + namespace { // Helper: every public API entry point gets the active backend through this @@ -57,9 +115,16 @@ hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreat // Pure CPU context: no Orochi/HIP/CUDA init needed. ctxt = new CpuContext( input ); } + else if ( input.deviceType == hiprtDeviceHybrid ) + { + // Hybrid context: GPU BVH via hiprt + CPU Embree mirror in parallel. + // Default to HIP; can be extended later with a secondary GPU-type flag. + oroInitialize( ORO_API_HIP, 0, g_hip_paths, g_hiprtc_paths ); + ctxt = new HybridContext( input ); + } else { - // GPU (and, later, GPU+CPU hybrid). Picking the underlying API: + // Pure GPU. Picking the underlying API: // AMD bit set => HIP, otherwise treat as NVIDIA/CUDA. oroInitialize( ( input.deviceType & hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA, @@ -804,3 +869,13 @@ hiprtError hiprtSetLogLevel( hiprtContext context, hiprtLogLevel level ) } return hiprtSuccess; } + +void* hiprtGetInternalCpuDataFromGeometry( hiprtGeometry geometry ) +{ + return static_cast( hiprt::lookupCpuGeom( geometry ) ); +} + +void* hiprtGetInternalCpuDataFromScene( hiprtScene scene ) +{ + return static_cast( hiprt::lookupCpuScene( scene ) ); +} From 3f3f329979b4af25ec680ee21427f9e0b624eac9 Mon Sep 17 00:00:00 2001 From: Krzysztof Date: Thu, 18 Jun 2026 17:31:48 +0200 Subject: [PATCH 6/7] Basic hybrid rendering helper --- CMakeLists.txt | 2 + hiprt/hiprt.h.in | 4 +- hiprt/hiprt_cpu.h | 91 +++--- hiprt/hiprt_hybrid.h | 76 +++++ hiprt/hiprt_types.h | 7 +- hiprt/impl/HybridContext.cpp | 16 +- hiprt/impl/hiprt.cpp | 14 +- hiprt/impl/hiprt_hybrid.cpp | 126 ++++++++ hiprt/kernels/HybridSceneTraceBatchKernel.h | 47 +++ premake5.lua | 15 +- test/CpuGeometryTest.cpp | 16 +- test/CpuSceneTest.cpp | 21 +- test/HybridGeometryTest.cpp | 236 ++++++++++++++ test/HybridGeometryTest.h | 30 ++ test/HybridMixedTracingTest.cpp | 142 ++++++++ test/HybridMixedTracingTest.h | 7 + test/HybridSceneTest.cpp | 339 ++++++++++++++++++++ test/HybridSceneTest.h | 34 ++ test/HybridTraceBatchTest.cpp | 209 ++++++++++++ test/HybridTraceBatchTest.h | 9 + test/kernels/HybridRowsKernel.h | 43 +++ 21 files changed, 1410 insertions(+), 74 deletions(-) create mode 100644 hiprt/hiprt_hybrid.h create mode 100644 hiprt/impl/hiprt_hybrid.cpp create mode 100644 hiprt/kernels/HybridSceneTraceBatchKernel.h create mode 100644 test/HybridGeometryTest.cpp create mode 100644 test/HybridGeometryTest.h create mode 100644 test/HybridMixedTracingTest.cpp create mode 100644 test/HybridMixedTracingTest.h create mode 100644 test/HybridSceneTest.cpp create mode 100644 test/HybridSceneTest.h create mode 100644 test/HybridTraceBatchTest.cpp create mode 100644 test/HybridTraceBatchTest.h create mode 100644 test/kernels/HybridRowsKernel.h diff --git a/CMakeLists.txt b/CMakeLists.txt index 608da9d..651e607 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -680,6 +680,8 @@ if(NOT NO_UNITTEST) "${CMAKE_CURRENT_SOURCE_DIR}/test/hiprtT*.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/test/Cpu*Test.h" "${CMAKE_CURRENT_SOURCE_DIR}/test/Cpu*Test.cpp" + "${CMAKE_CURRENT_SOURCE_DIR}/test/Hybrid*Test.h" + "${CMAKE_CURRENT_SOURCE_DIR}/test/Hybrid*Test.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/test/shared.h" "${CMAKE_CURRENT_SOURCE_DIR}/test/main.cpp" "${CMAKE_CURRENT_SOURCE_DIR}/test/CornellBox.h" diff --git a/hiprt/hiprt.h.in b/hiprt/hiprt.h.in index 392e14d..4dc1d46 100644 --- a/hiprt/hiprt.h.in +++ b/hiprt/hiprt.h.in @@ -602,7 +602,7 @@ HIPRT_API hiprtError hiprtSetLogLevel( hiprtContext context, hiprtLogLevel level /** \brief Returns the Embree CpuGeometryData pointer associated with a geometry handle. * - * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceHybrid) contexts. + * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceAMD | hiprtDeviceCPU) contexts. * Returns nullptr if the handle has no CPU mirror (pure GPU context). * The returned pointer is opaque; cast it to hiprt::CpuGeometryData* in hiprt_cpu.h. * @@ -613,7 +613,7 @@ HIPRT_API void* hiprtGetInternalCpuDataFromGeometry( hiprtGeometry geometry ); /** \brief Returns the Embree CpuSceneData pointer associated with a scene handle. * - * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceHybrid) contexts. + * Works for both pure-CPU (hiprtDeviceCPU) and hybrid (hiprtDeviceAMD | hiprtDeviceCPU) contexts. * Returns nullptr if the handle has no CPU mirror (pure GPU context). * The returned pointer is opaque; cast it to hiprt::CpuSceneData* in hiprt_cpu.h. * diff --git a/hiprt/hiprt_cpu.h b/hiprt/hiprt_cpu.h index 263516e..5c7678c 100644 --- a/hiprt/hiprt_cpu.h +++ b/hiprt/hiprt_cpu.h @@ -24,7 +24,7 @@ #pragma once -#include // HIPRT_API + hiprtGetInternalCpuDataFromGeometry/Scene +#include #include #include #include @@ -58,7 +58,7 @@ inline hiprtHit makeHiprtHit( const RTCRayHit& rh ) noexcept { hiprtHit hit{}; if ( rh.hit.geomID == RTC_INVALID_GEOMETRY_ID ) - return hit; // primID stays hiprtInvalidValue -> hasHit() == false + return hit; hit.primID = rh.hit.primID; hit.instanceID = ( rh.hit.instID[0] != RTC_INVALID_GEOMETRY_ID ) @@ -69,68 +69,67 @@ inline hiprtHit makeHiprtHit( const RTCRayHit& rh ) noexcept hit.normal.x = rh.hit.Ng_x; hit.normal.y = rh.hit.Ng_y; hit.normal.z = rh.hit.Ng_z; - hit.t = rh.ray.tfar; // Embree writes the hit distance into tfar + hit.t = rh.ray.tfar; return hit; } -} // namespace hiprt_cpu_detail - -// --------------------------------------------------------------------------- -// Public CPU traversal objects -// --------------------------------------------------------------------------- - -/** \brief Host-side traversal that finds the closest hit in a hiprtGeometry. - * - * Mirrors the GPU class hiprtGeomTraversalClosest. Constructing the object - * immediately casts the handle to CpuGeometryData, retrieves the underlying - * RTCScene built by CpuContext, and calls rtcIntersect1. - * Call getNextHit() to obtain the result. - */ -class hiprtGeomTraversalClosestCPU +inline void traceGeometryBatch( + RTCScene scene, const hiprtRay* rays, hiprtHit* hits, uint32_t count ) noexcept { - public: - hiprtGeomTraversalClosestCPU( hiprtGeometry geom, const hiprtRay& ray ) noexcept + if ( scene == nullptr ) { - m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); - auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromGeometry( geom ) ); - if ( data != nullptr && data->rtcScene != nullptr ) - rtcIntersect1( data->rtcScene, &m_rayHit ); + for ( uint32_t i = 0; i < count; ++i ) + hits[i] = hiprtHit{}; + return; } - hiprtHit getNextHit() const noexcept + for ( uint32_t i = 0; i < count; ++i ) { - return hiprt_cpu_detail::makeHiprtHit( m_rayHit ); + RTCRayHit rh = makeRTCRayHit( rays[i] ); + rtcIntersect1( scene, &rh ); + hits[i] = makeHiprtHit( rh ); } +} - private: - RTCRayHit m_rayHit{}; -}; +} // namespace hiprt_cpu_detail + +// --------------------------------------------------------------------------- +// Public CPU traversal API +// --------------------------------------------------------------------------- -/** \brief Host-side traversal that finds the closest hit in a hiprtScene. +/** \brief Host-side batch traversal that finds the closest hit. * - * Mirrors the GPU class hiprtSceneTraversalClosest. Constructing the object - * immediately casts the handle to CpuSceneData, retrieves the underlying - * RTCScene built by CpuContext (which already encodes the full instance - * hierarchy via RTC_GEOMETRY_TYPE_INSTANCE), and calls rtcIntersect1. - * Call getNextHit() to obtain the result, including instanceID populated - * from instID[0]. + * In CPU-only mode the handle is a direct CpuGeometryData / CpuSceneData pointer. + * In hybrid mode (hiprtDeviceAMD | hiprtDeviceCPU or hiprtDeviceNVIDIA | hiprtDeviceCPU) + * geometry and scene handles are GPU handles; the CPU Embree mirror is resolved + * through the context side table. */ -class hiprtSceneTraversalClosestCPU +class hiprtGeomTraversalClosestCPU { public: - hiprtSceneTraversalClosestCPU( hiprtScene scene, const hiprtRay& ray ) noexcept + static void traceBatch( + hiprtContext ctx, + hiprtGeometry geom, + const hiprtRay* rays, + hiprtHit* hits, + uint32_t count ) noexcept { - m_rayHit = hiprt_cpu_detail::makeRTCRayHit( ray ); - auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromScene( scene ) ); - if ( data != nullptr && data->rtcScene != nullptr ) - rtcIntersect1( data->rtcScene, &m_rayHit ); + (void)ctx; + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromGeometry( geom ) ); + hiprt_cpu_detail::traceGeometryBatch( + ( data != nullptr ) ? data->rtcScene : nullptr, rays, hits, count ); } - hiprtHit getNextHit() const noexcept + static void traceBatch( + hiprtContext ctx, + hiprtScene scene, + const hiprtRay* rays, + hiprtHit* hits, + uint32_t count ) noexcept { - return hiprt_cpu_detail::makeHiprtHit( m_rayHit ); + (void)ctx; + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromScene( scene ) ); + hiprt_cpu_detail::traceGeometryBatch( + ( data != nullptr ) ? data->rtcScene : nullptr, rays, hits, count ); } - - private: - RTCRayHit m_rayHit{}; }; diff --git a/hiprt/hiprt_hybrid.h b/hiprt/hiprt_hybrid.h new file mode 100644 index 0000000..d2de7f7 --- /dev/null +++ b/hiprt/hiprt_hybrid.h @@ -0,0 +1,76 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include + +/** \brief Scheduling parameters for hybrid closest-hit tracing. + */ +struct hiprtHybridTraceConfig +{ + /*!< Fraction of rays dispatched to the GPU in [0, 1]. */ + float gpuFraction = 0.8f; + /*!< Minimum number of rays kept for the CPU path when both backends are used. */ + uint32_t minCpuBatch = 1024u; +}; + +/** \brief GPU resources required for hybrid scene tracing. + * + * Compile hiprt/kernels/HybridSceneTraceBatchKernel.h with hiprtBuildTraceKernels + * and pass the resulting function handle here. globalStack is unused by the + * default batch kernel but reserved for advanced kernels such as TraceKernel. + */ +struct hiprtHybridTraceGpuInput +{ + hiprtApiFunction traceKernel = nullptr; + hiprtGlobalStackBuffer globalStack = {}; +}; + +/** \brief Trace a batch of closest-hit scene rays on GPU and CPU concurrently. + * + * Requires a hybrid context (hiprtDeviceAMD | hiprtDeviceCPU or + * hiprtDeviceNVIDIA | hiprtDeviceCPU). The first gpuCount rays are traced on + * the GPU; the remainder on the CPU Embree mirror. rays and hits must be + * host-visible (managed/UVA or pinned mapped) so both backends can access them. + * + * \param context Hybrid HIPRT context. + * \param scene GPU scene handle with a CPU mirror. + * \param config Split policy between GPU and CPU. + * \param gpuInput Precompiled GPU trace kernel (and optional stack buffer). + * \param rays Host-visible ray array of length rayCount. + * \param hits Host-visible hit array of length rayCount. + * \param rayCount Number of rays to trace. + * \param stream GPU stream for the device launch (0 for default stream). + * \return hiprtSuccess or an error code. + */ +HIPRT_API hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); diff --git a/hiprt/hiprt_types.h b/hiprt/hiprt_types.h index b70b220..555b3d1 100644 --- a/hiprt/hiprt_types.h +++ b/hiprt/hiprt_types.h @@ -314,11 +314,16 @@ enum hiprtDeviceType : uint32_t hiprtDeviceAMD = 1 << 0, hiprtDeviceNVIDIA = 1 << 1, hiprtDeviceCPU = 1 << 2, - hiprtDeviceHybrid = 1 << 3, }; /** \brief Context creation input. * + * Hybrid mode is selected with hiprtDeviceAMD | hiprtDeviceCPU or + * hiprtDeviceNVIDIA | hiprtDeviceCPU. + * + * For hiprtDeviceCPU and hybrid modes, all geometry/scene build inputs + * (vertices, indices, instances, transforms, etc.) must be host-readable. + * Use managed/UVA or pinned mapped allocations so Embree can read them on CPU. */ struct hiprtContextCreationInput { diff --git a/hiprt/impl/HybridContext.cpp b/hiprt/impl/HybridContext.cpp index e0670d3..32c536c 100644 --- a/hiprt/impl/HybridContext.cpp +++ b/hiprt/impl/HybridContext.cpp @@ -1,12 +1,22 @@ #include #include +#include #include #include #include namespace hiprt { +namespace +{ +void syncGpu( oroStream stream ) +{ + const oroError err = ( stream != nullptr ) ? oroStreamSynchronize( stream ) : oroDeviceSynchronize(); + if ( err != oroSuccess ) + throw std::runtime_error( "HybridContext: GPU synchronization failed before CPU mirror build" ); +} +} // namespace // --------------------------------------------------------------------------- // Construction / destruction @@ -15,7 +25,7 @@ namespace hiprt HybridContext::HybridContext( const hiprtContextCreationInput& input ) : m_gpu( input ), m_cpu( input ) { - std::cerr << "[HybridContext] constructed — GPU BVH + Embree CPU mirror active.\n"; + std::cerr << "[HybridContext] constructed - GPU BVH + Embree CPU mirror active.\n"; } HybridContext::~HybridContext() @@ -95,6 +105,7 @@ void HybridContext::buildGeometries( std::vector& buffers ) { m_gpu.buildGeometries( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + syncGpu( stream ); dispatchCpuGeomBuild( buildInputs, buffers, false ); } @@ -106,6 +117,7 @@ void HybridContext::updateGeometries( std::vector& buffers ) { m_gpu.updateGeometries( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + syncGpu( stream ); dispatchCpuGeomBuild( buildInputs, buffers, true ); } @@ -231,6 +243,7 @@ void HybridContext::buildScenes( std::vector& buffers ) { m_gpu.buildScenes( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + syncGpu( stream ); dispatchCpuSceneBuild( buildInputs, buffers, false ); } @@ -242,6 +255,7 @@ void HybridContext::updateScenes( std::vector& buffers ) { m_gpu.updateScenes( buildInputs, buildOptions, temporaryBuffer, stream, buffers ); + syncGpu( stream ); dispatchCpuSceneBuild( buildInputs, buffers, true ); } diff --git a/hiprt/impl/hiprt.cpp b/hiprt/impl/hiprt.cpp index b68fd88..cfa70df 100644 --- a/hiprt/impl/hiprt.cpp +++ b/hiprt/impl/hiprt.cpp @@ -115,17 +115,19 @@ hiprtError hiprtCreateContext( uint32_t hiprtApiVersion, const hiprtContextCreat // Pure CPU context: no Orochi/HIP/CUDA init needed. ctxt = new CpuContext( input ); } - else if ( input.deviceType == hiprtDeviceHybrid ) + else if ( input.deviceType & hiprtDeviceCPU ) { - // Hybrid context: GPU BVH via hiprt + CPU Embree mirror in parallel. - // Default to HIP; can be extended later with a secondary GPU-type flag. - oroInitialize( ORO_API_HIP, 0, g_hip_paths, g_hiprtc_paths ); + // Hybrid: hiprtDeviceAMD | hiprtDeviceCPU or hiprtDeviceNVIDIA | hiprtDeviceCPU. + oroInitialize( + ( input.deviceType & hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA, + 0, + g_hip_paths, + g_hiprtc_paths ); ctxt = new HybridContext( input ); } else { - // Pure GPU. Picking the underlying API: - // AMD bit set => HIP, otherwise treat as NVIDIA/CUDA. + // Pure GPU. AMD bit set => HIP, otherwise treat as NVIDIA/CUDA. oroInitialize( ( input.deviceType & hiprtDeviceAMD ) ? ORO_API_HIP : ORO_API_CUDA, 0, diff --git a/hiprt/impl/hiprt_hybrid.cpp b/hiprt/impl/hiprt_hybrid.cpp new file mode 100644 index 0000000..e83f003 --- /dev/null +++ b/hiprt/impl/hiprt_hybrid.cpp @@ -0,0 +1,126 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#include +#include +#include + +#include + +using namespace hiprt; + +namespace +{ +ContextBase* asContext( hiprtContext context ) noexcept +{ + return reinterpret_cast( context ); +} + +uint32_t computeGpuRayCount( uint32_t rayCount, const hiprtHybridTraceConfig& config ) noexcept +{ + if ( rayCount == 0 ) + return 0; + + uint32_t gpuCount = static_cast( static_cast( rayCount ) * config.gpuFraction ); + if ( gpuCount > rayCount ) + gpuCount = rayCount; + + const uint32_t cpuCount = rayCount - gpuCount; + if ( cpuCount > 0 && cpuCount < config.minCpuBatch && rayCount > config.minCpuBatch ) + { + gpuCount = rayCount - config.minCpuBatch; + if ( gpuCount > rayCount ) + gpuCount = rayCount; + } + + return gpuCount; +} + +hiprtError launchSceneTraceBatchKernel( + const hiprtHybridTraceGpuInput& gpuInput, + hiprtScene scene, + uint32_t numRays, + hiprtRay* rays, + hiprtHit* hits, + hiprtApiStream stream ) +{ + if ( gpuInput.traceKernel == nullptr ) + return hiprtErrorInvalidParameter; + + constexpr uint32_t blockSize = 256u; + const uint32_t gridSize = ( numRays + blockSize - 1 ) / blockSize; + + hiprtApiFunction funcApi = gpuInput.traceKernel; + oroFunction func = *reinterpret_cast( &funcApi ); + + void* args[] = { &scene, &numRays, &rays, &hits }; + const oroError err = oroModuleLaunchKernel( + func, gridSize, 1, 1, blockSize, 1, 1, 0, reinterpret_cast( stream ), args, 0 ); + return ( err == oroSuccess ) ? hiprtSuccess : hiprtErrorInternal; +} +} // namespace + +hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + if ( context == nullptr || scene == nullptr || rays == nullptr || hits == nullptr ) + return hiprtErrorInvalidParameter; + + if ( dynamic_cast( asContext( context ) ) == nullptr ) + return hiprtErrorInvalidParameter; + + if ( hiprtGetInternalCpuDataFromScene( scene ) == nullptr ) + return hiprtErrorInvalidParameter; + + const uint32_t gpuCount = computeGpuRayCount( rayCount, config ); + const uint32_t cpuCount = rayCount - gpuCount; + + if ( gpuCount > 0 ) + { + const hiprtError gpuErr = + launchSceneTraceBatchKernel( gpuInput, scene, gpuCount, rays, hits, stream ); + if ( gpuErr != hiprtSuccess ) + return gpuErr; + + const oroStream gpuStream = reinterpret_cast( stream ); + const oroError err = ( gpuStream != nullptr ) ? oroStreamSynchronize( gpuStream ) : oroDeviceSynchronize(); + if ( err != oroSuccess ) + return hiprtErrorInternal; + } + + if ( cpuCount > 0 ) + { + hiprtGeomTraversalClosestCPU::traceBatch( + context, scene, rays + gpuCount, hits + gpuCount, cpuCount ); + } + + return hiprtSuccess; +} diff --git a/hiprt/kernels/HybridSceneTraceBatchKernel.h b/hiprt/kernels/HybridSceneTraceBatchKernel.h new file mode 100644 index 0000000..0e437e6 --- /dev/null +++ b/hiprt/kernels/HybridSceneTraceBatchKernel.h @@ -0,0 +1,47 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include + +#if defined( __KERNELCC__ ) +#include +#endif + +/** \brief Closest-hit scene traversal for a contiguous prefix of a ray batch. + * + * Intended for use with hiprtTraceHybridClosest. Uses the built-in traversal + * stack inside hiprtSceneTraversalClosest (no global stack buffer required). + */ +extern "C" __global__ void HybridSceneTraceBatchKernel( + hiprtScene scene, uint32_t numRays, hiprtRay* rays, hiprtHit* hits ) +{ + const uint32_t index = blockIdx.x * blockDim.x + threadIdx.x; + if ( index >= numRays ) + return; + + hiprtSceneTraversalClosest tr( scene, rays[index] ); + hits[index] = tr.getNextHit(); +} diff --git a/premake5.lua b/premake5.lua index 7231b43..6e08f2a 100644 --- a/premake5.lua +++ b/premake5.lua @@ -132,6 +132,19 @@ workspace "hiprt" files {"contrib/Orochi/contrib/hipew/**.h", "contrib/Orochi/contrib/hipew/**.cpp"} files {"contrib/Orochi/ParallelPrimitives/**.h", "contrib/Orochi/ParallelPrimitives/**.cpp"} + externalincludedirs { "contrib/embree/include/" } + if os.istarget("windows") then + libdirs { "contrib/embree/win/" } + end + if os.istarget("linux") then + local embree_rel_path = "contrib/embree/linux" + local project_root = path.getabsolute(".") + local embree_abs_path = path.join(project_root, embree_rel_path) + libdirs { embree_rel_path } + linkoptions { "-Wl,-rpath," .. embree_abs_path } + end + links { "embree4", "tbb" } + if not _OPTIONS["noUnittest"] then project( "unittest" ) cppdialect "C++17" @@ -149,7 +162,7 @@ workspace "hiprt" if os.ishost("linux") then links { "pthread", "dl" } end - files { "test/hiprtT*.h", "test/hiprtT*.cpp", "test/shared.h", "test/main.cpp", "test/CornellBox.h", "test/kernels/*.h" } + files { "test/hiprtT*.h", "test/hiprtT*.cpp", "test/Cpu*Test.h", "test/Cpu*Test.cpp", "test/Hybrid*Test.h", "test/Hybrid*Test.cpp", "test/shared.h", "test/main.cpp", "test/CornellBox.h", "test/kernels/*.h" } externalincludedirs { "./contrib/Orochi/" } files {"contrib/Orochi/Orochi/**.h", "contrib/Orochi/Orochi/**.cpp"} files {"contrib/Orochi/contrib/cuew/**.h", "contrib/Orochi/contrib/cuew/**.cpp"} diff --git a/test/CpuGeometryTest.cpp b/test/CpuGeometryTest.cpp index 3b94de4..ff819fb 100644 --- a/test/CpuGeometryTest.cpp +++ b/test/CpuGeometryTest.cpp @@ -59,7 +59,8 @@ hiprtGeometry CpuGeometryTest::buildSampleTriangle( namespace { -bool traceHit( hiprtGeometry geom, float ox, float oy, float oz, float dx, float dy, float dz ) +bool traceHit( + hiprtContext ctx, hiprtGeometry geom, float ox, float oy, float oz, float dx, float dy, float dz ) { hiprtRay ray{}; ray.origin = { ox, oy, oz }; @@ -67,8 +68,9 @@ bool traceHit( hiprtGeometry geom, float ox, float oy, float oz, float dx, float ray.direction = { dx, dy, dz }; ray.maxT = 1e30f; - hiprtGeomTraversalClosestCPU tr( geom, ray ); - return tr.getNextHit().hasHit(); + hiprtHit hit{}; + hiprtGeomTraversalClosestCPU::traceBatch( ctx, geom, &ray, &hit, 1 ); + return hit.hasHit(); } } // namespace @@ -111,9 +113,9 @@ TEST_F( CpuGeometryTest, BuildGeometryProducesHittableBvh ) hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); ASSERT_TRUE( geom != nullptr ); - EXPECT_TRUE( traceHit( geom, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceHit( m_context, geom, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Promien ze srodka trojkata powinien trafic w zbudowane BVH."; - EXPECT_FALSE( traceHit( geom, 2.0f, 2.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_FALSE( traceHit( m_context, geom, 2.0f, 2.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Promien poza trojkatem nie powinien trafiac."; ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); @@ -127,7 +129,7 @@ TEST_F( CpuGeometryTest, UpdateGeometryRebuildsBvh ) hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); ASSERT_TRUE( geom != nullptr ); - EXPECT_FALSE( traceHit( geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_FALSE( traceHit( m_context, geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Przed update BVH nie powinno obejmowac y=1.5."; vertices[2] = { 0.0f, 2.0f, 0.0f }; @@ -140,7 +142,7 @@ TEST_F( CpuGeometryTest, UpdateGeometryRebuildsBvh ) hiprtSuccess ) << "hiprtBuildGeometry (Update) failed to rebuild the scene."; - EXPECT_TRUE( traceHit( geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceHit( m_context, geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Po update BVH powinno obejmowac y=1.5 (nowy wierzcholek y=2)."; ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); diff --git a/test/CpuSceneTest.cpp b/test/CpuSceneTest.cpp index df36e67..ef2927d 100644 --- a/test/CpuSceneTest.cpp +++ b/test/CpuSceneTest.cpp @@ -67,7 +67,8 @@ hiprtFrameSRT makeSrt( float tx, float ty, float tz ) return f; } -bool traceScene( hiprtScene scene, float ox, float oy, float oz, float dx, float dy, float dz ) +bool traceScene( + hiprtContext ctx, hiprtScene scene, float ox, float oy, float oz, float dx, float dy, float dz ) { hiprtRay ray{}; ray.origin = { ox, oy, oz }; @@ -75,8 +76,8 @@ bool traceScene( hiprtScene scene, float ox, float oy, float oz, float dx, float ray.direction = { dx, dy, dz }; ray.maxT = 1e30f; - hiprtSceneTraversalClosestCPU tr( scene, ray ); - const hiprtHit hit = tr.getNextHit(); + hiprtHit hit{}; + hiprtGeomTraversalClosestCPU::traceBatch( ctx, scene, &ray, &hit, 1 ); return hit.instanceID != hiprtInvalidValue; } } // namespace @@ -127,9 +128,9 @@ TEST_F( CpuSceneTest, BuildSceneWithSingleInstanceIsHittable ) hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), hiprtSuccess ); - EXPECT_TRUE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Promien w obszar instancji powinien trafic."; - EXPECT_FALSE( traceScene( scene, 5.0f, 5.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_FALSE( traceScene( m_context, scene, 5.0f, 5.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Promien poza instancja nie powinien trafic."; ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); @@ -165,9 +166,9 @@ TEST_F( CpuSceneTest, BuildSceneWithTranslatedInstance ) hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), hiprtSuccess ); - EXPECT_TRUE( traceScene( scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceScene( m_context, scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Po translacji +3X instancja powinna byc trafiona w x=3.25."; - EXPECT_FALSE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_FALSE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Oryginalna pozycja x=0.25 nie powinna byc juz trafiona."; ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); @@ -203,7 +204,7 @@ TEST_F( CpuSceneTest, UpdateSceneRebuildsInstances ) hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, nullptr, nullptr, scene ), hiprtSuccess ); - EXPECT_TRUE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Przed update trafienie w x=0.25."; frames[0] = makeSrt( 3.0f, 0.0f, 0.0f ); @@ -213,9 +214,9 @@ TEST_F( CpuSceneTest, UpdateSceneRebuildsInstances ) hiprtSuccess ) << "hiprtBuildScene (Update) failed."; - EXPECT_FALSE( traceScene( scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_FALSE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Po update instancja przesunieta, stara pozycja pusta."; - EXPECT_TRUE( traceScene( scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + EXPECT_TRUE( traceScene( m_context, scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) << "Po update trafienie w nowej pozycji x=3.25."; ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); diff --git a/test/HybridGeometryTest.cpp b/test/HybridGeometryTest.cpp new file mode 100644 index 0000000..8a44336 --- /dev/null +++ b/test/HybridGeometryTest.cpp @@ -0,0 +1,236 @@ +#include "HybridGeometryTest.h" + +#include + +#include + +extern const char* g_hip_paths[]; +extern const char* g_hiprtc_paths[]; + +namespace +{ +void assertOroSuccess( oroError err ) +{ + if ( err != oroSuccess ) + { + const char* msg = nullptr; + oroGetErrorString( err, &msg ); + FAIL() << "Orochi error: " << ( msg != nullptr ? msg : "unknown" ); + } +} +} // namespace + +void HybridGeometryTest::freeManagedAllocs() +{ + for ( void* ptr : m_managedAllocs ) + assertOroSuccess( oroFree( ptr ) ); + m_managedAllocs.clear(); +} + +void HybridGeometryTest::SetUp() +{ + oroInitialize( (oroApi)( ORO_API_HIP | ORO_API_CUDA ), 0, g_hip_paths, g_hiprtc_paths ); + assertOroSuccess( oroInit( 0 ) ); + assertOroSuccess( oroDeviceGet( &m_oroDevice, 0 ) ); + assertOroSuccess( oroCtxCreate( &m_oroCtx, 0, m_oroDevice ) ); + assertOroSuccess( oroCtxSetCurrent( m_oroCtx ) ); + + oroDeviceProp props{}; + assertOroSuccess( oroGetDeviceProperties( &props, m_oroDevice ) ); + + hiprtContextCreationInput ctxInput{}; + ctxInput.numCpuThreads = 1; + ctxInput.ctxt = oroGetRawCtx( m_oroCtx ); + ctxInput.device = oroGetRawDevice( m_oroDevice ); + ctxInput.deviceType = ( std::string( props.name ).find( "NVIDIA" ) != std::string::npos ) + ? static_cast( hiprtDeviceNVIDIA | hiprtDeviceCPU ) + : static_cast( hiprtDeviceAMD | hiprtDeviceCPU ); + + ASSERT_EQ( hiprtCreateContext( HIPRT_API_VERSION, ctxInput, m_context ), hiprtSuccess ) + << "Failed to create hybrid context."; + ASSERT_TRUE( m_context != nullptr ) << "Context pointer is null after creation."; +} + +void HybridGeometryTest::TearDown() +{ + if ( m_context != nullptr ) + EXPECT_EQ( hiprtDestroyContext( m_context ), hiprtSuccess ) + << "Failed to cleanly destroy hybrid context."; + m_context = nullptr; + + freeManagedAllocs(); + + if ( m_oroCtx != nullptr ) + assertOroSuccess( oroCtxDestroy( m_oroCtx ) ); + m_oroCtx = nullptr; +} + +hiprtDevicePtr HybridGeometryTest::allocGeometryBuildTemp( + const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions& buildOptions ) +{ + size_t tempSize = 0; + hiprtDevicePtr temp = nullptr; + if ( hiprtGetGeometryBuildTemporaryBufferSize( m_context, buildInput, buildOptions, tempSize ) != hiprtSuccess ) + return nullptr; + if ( tempSize == 0 ) + return nullptr; + + assertOroSuccess( oroMalloc( reinterpret_cast( &temp ), tempSize ) ); + m_managedAllocs.push_back( reinterpret_cast( temp ) ); + return temp; +} + +hiprtGeometry HybridGeometryTest::buildSampleTriangle( + std::vector& vertices, + std::vector& indices, + hiprtGeometryBuildInput& geomInputOut ) +{ + hiprtFloat3* hostVertices = nullptr; + uint32_t* hostIndices = nullptr; + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostVertices ), 3 * sizeof( hiprtFloat3 ), oroMemAttachGlobal ) ); + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostIndices ), 3 * sizeof( uint32_t ), oroMemAttachGlobal ) ); + m_managedAllocs.push_back( hostVertices ); + m_managedAllocs.push_back( hostIndices ); + + hostVertices[0] = { 0.0f, 0.0f, 0.0f }; + hostVertices[1] = { 1.0f, 0.0f, 0.0f }; + hostVertices[2] = { 0.0f, 1.0f, 0.0f }; + hostIndices[0] = 0; + hostIndices[1] = 1; + hostIndices[2] = 2; + + vertices = { hostVertices[0], hostVertices[1], hostVertices[2] }; + indices = { hostIndices[0], hostIndices[1], hostIndices[2] }; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = hostVertices; + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = hostIndices; + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + geomInputOut = {}; + geomInputOut.type = hiprtPrimitiveTypeTriangleMesh; + geomInputOut.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + EXPECT_EQ( hiprtCreateGeometry( m_context, geomInputOut, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed."; + EXPECT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + EXPECT_EQ( + hiprtBuildGeometry( + m_context, hiprtBuildOperationBuild, geomInputOut, opts, + allocGeometryBuildTemp( geomInputOut, opts ), nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry (Build) failed."; + return geom; +} + +namespace +{ +bool traceHit( + hiprtContext ctx, hiprtGeometry geom, float ox, float oy, float oz, float dx, float dy, float dz ) +{ + hiprtRay ray{}; + ray.origin = { ox, oy, oz }; + ray.minT = 0.0f; + ray.direction = { dx, dy, dz }; + ray.maxT = 1e30f; + + hiprtHit hit{}; + hiprtGeomTraversalClosestCPU::traceBatch( ctx, geom, &ray, &hit, 1 ); + return hit.hasHit(); +} +} // namespace + +TEST_F( HybridGeometryTest, CreateAndDestroyGeometry ) +{ + hiprtFloat3* vertices = nullptr; + uint32_t* indices = nullptr; + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &vertices ), 3 * sizeof( hiprtFloat3 ), oroMemAttachGlobal ) ); + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &indices ), 3 * sizeof( uint32_t ), oroMemAttachGlobal ) ); + m_managedAllocs.push_back( vertices ); + m_managedAllocs.push_back( indices ); + + vertices[0] = { 0.0f, 0.0f, 0.0f }; + vertices[1] = { 1.0f, 0.0f, 0.0f }; + vertices[2] = { 0.0f, 1.0f, 0.0f }; + indices[0] = 0; + indices[1] = 1; + indices[2] = 2; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = vertices; + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = indices; + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + hiprtGeometryBuildInput in{}; + in.type = hiprtPrimitiveTypeTriangleMesh; + in.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + ASSERT_EQ( hiprtCreateGeometry( m_context, in, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed on hybrid context."; + ASSERT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ) + << "hiprtDestroyGeometry failed to release GPU/CPU mirror."; +} + +TEST_F( HybridGeometryTest, BuildGeometryProducesHittableBvhViaCpuMirror ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometryBuildInput in{}; + hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); + ASSERT_TRUE( geom != nullptr ); + + EXPECT_TRUE( traceHit( m_context, geom, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "CPU traceBatch on GPU handle should hit the Embree mirror."; + EXPECT_FALSE( traceHit( m_context, geom, 2.0f, 2.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Ray outside triangle should miss."; + + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridGeometryTest, UpdateGeometryRebuildsCpuMirror ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometryBuildInput in{}; + hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); + ASSERT_TRUE( geom != nullptr ); + + EXPECT_FALSE( traceHit( m_context, geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Before update CPU mirror should not cover y=1.5."; + + auto* meshVertices = reinterpret_cast( in.primitive.triangleMesh.vertices ); + meshVertices[2] = { 0.0f, 2.0f, 0.0f }; + vertices[2] = meshVertices[2]; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + ASSERT_EQ( + hiprtBuildGeometry( + m_context, hiprtBuildOperationUpdate, in, opts, allocGeometryBuildTemp( in, opts ), nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry (Update) failed to rebuild hybrid mirrors."; + + EXPECT_TRUE( traceHit( m_context, geom, 0.1f, 1.5f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "After update CPU mirror should cover y=1.5."; + + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} diff --git a/test/HybridGeometryTest.h b/test/HybridGeometryTest.h new file mode 100644 index 0000000..1ba16f8 --- /dev/null +++ b/test/HybridGeometryTest.h @@ -0,0 +1,30 @@ +#pragma once + +#include +#include +#include +#include + +class HybridGeometryTest : public ::testing::Test +{ + protected: + void SetUp() override; + void TearDown() override; + + hiprtGeometry buildSampleTriangle( + std::vector& vertices, + std::vector& indices, + hiprtGeometryBuildInput& geomInputOut ); + + hiprtContext m_context = nullptr; + oroCtx m_oroCtx = nullptr; + oroDevice m_oroDevice = 0; + + void freeManagedAllocs(); + std::vector m_managedAllocs; + + hiprtDevicePtr allocGeometryBuildTemp( + const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions& buildOptions ); + hiprtDevicePtr allocSceneBuildTemp( + const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions& buildOptions ); +}; diff --git a/test/HybridMixedTracingTest.cpp b/test/HybridMixedTracingTest.cpp new file mode 100644 index 0000000..119667d --- /dev/null +++ b/test/HybridMixedTracingTest.cpp @@ -0,0 +1,142 @@ +#include "HybridMixedTracingTest.h" + +#include + +#include +#include +#include +#include + +extern const char* g_hip_paths[]; +extern const char* g_hiprtc_paths[]; + +namespace +{ +std::filesystem::path getRootDir() +{ + const char* env = std::getenv( "HIPRT_PATH" ); + return env != nullptr ? std::filesystem::path( env ) : std::filesystem::path( ".." ); +} + +bool readKernelSource( const std::filesystem::path& srcPath, std::string& sourceCode ) +{ + std::ifstream f( srcPath ); + if ( !f.is_open() ) + return false; + sourceCode.assign( std::istreambuf_iterator( f ), std::istreambuf_iterator() ); + return true; +} + +oroFunction buildHybridRowsKernel( hiprtContext context ) +{ + const auto kernelPath = getRootDir() / "test/kernels/HybridRowsKernel.h"; + std::string sourceCode; + const std::string moduleName = kernelPath.string(); + if ( !readKernelSource( kernelPath, sourceCode ) ) + return nullptr; + + const char* funcName = "HybridRowsKernel"; + + std::vector optionStorage; + optionStorage.push_back( "-I" + getRootDir().string() ); + if ( oroGetCurAPI( 0 ) == ORO_API_HIP ) + optionStorage.push_back( "-ffast-math" ); + else + optionStorage.push_back( "--use_fast_math" ); + + std::vector options; + options.reserve( optionStorage.size() ); + for ( const auto& opt : optionStorage ) + options.push_back( opt.c_str() ); + + hiprtApiFunction funcApi = nullptr; + if ( hiprtBuildTraceKernels( + context, 1, &funcName, sourceCode.c_str(), moduleName.c_str(), 0, nullptr, nullptr, + static_cast( options.size() ), options.data(), 0, 0, nullptr, &funcApi, nullptr, + true ) != hiprtSuccess ) + return nullptr; + + return *reinterpret_cast( &funcApi ); +} + +void launchHybridRowsKernel( + oroFunction func, hiprtGeometry geom, uint8_t* image, + uint32_t width, uint32_t height, uint32_t yBegin, uint32_t yEnd, uint8_t hitValue ) +{ + const uint32_t tx = 8; + const uint32_t ty = 8; + const uint32_t nbx = ( width + tx - 1 ) / tx; + const uint32_t nby = ( height + ty - 1 ) / ty; + void* args[] = { + &geom, &image, &width, &height, &yBegin, &yEnd, &hitValue, + }; + const oroError err = oroModuleLaunchKernel( func, nbx, nby, 1, tx, ty, 1, 0, 0, args, 0 ); + if ( err != oroSuccess ) + FAIL() << "oroModuleLaunchKernel failed"; +} + +void renderCpuRows( + hiprtContext context, hiprtGeometry geom, uint8_t* image, + uint32_t width, uint32_t height, uint32_t yBegin, uint32_t yEnd, uint8_t hitValue ) +{ + for ( uint32_t y = yBegin; y < yEnd; ++y ) + { + for ( uint32_t x = 0; x < width; ++x ) + { + hiprtRay ray{}; + ray.origin = { x / static_cast( width ), y / static_cast( height ), -1.0f }; + ray.direction = { 0.0f, 0.0f, 1.0f }; + ray.minT = 0.0f; + ray.maxT = 1e30f; + + hiprtHit hit{}; + hiprtGeomTraversalClosestCPU::traceBatch( context, geom, &ray, &hit, 1 ); + + const uint32_t index = x + y * width; + const uint8_t shade = hit.hasHit() ? hitValue : 0; + image[index * 4 + 0] = shade; + image[index * 4 + 1] = shade; + image[index * 4 + 2] = shade; + image[index * 4 + 3] = 255; + } + } +} +} // namespace + +TEST_F( HybridMixedTracingTest, SplitRowsGpuTopCpuBottom ) +{ + constexpr uint32_t width = 32; + constexpr uint32_t height = 32; + constexpr uint32_t splitY = height / 2; + constexpr uint8_t gpuHit = 200; + constexpr uint8_t cpuHit = 100; + + std::vector vertices; + std::vector indices; + hiprtGeometryBuildInput in{}; + hiprtGeometry geom = buildSampleTriangle( vertices, indices, in ); + ASSERT_TRUE( geom != nullptr ); + + uint8_t* image = nullptr; + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &image ), width * height * 4, oroMemAttachGlobal ), + oroSuccess ); + m_managedAllocs.push_back( image ); + std::memset( image, 0, width * height * 4 ); + + oroFunction kernel = buildHybridRowsKernel( m_context ); + ASSERT_TRUE( kernel != nullptr ) << "Failed to compile HybridRowsKernel."; + + launchHybridRowsKernel( kernel, geom, image, width, height, 0, splitY, gpuHit ); + ASSERT_EQ( oroDeviceSynchronize(), oroSuccess ); + + renderCpuRows( m_context, geom, image, width, height, splitY, height, cpuHit ); + + const auto pixel = [&]( uint32_t x, uint32_t y ) -> uint8_t { return image[( x + y * width ) * 4]; }; + + EXPECT_EQ( pixel( 8, 8 ), gpuHit ) << "GPU top row should hit triangle."; + EXPECT_EQ( pixel( 8, 24 ), cpuHit ) << "CPU bottom row should hit triangle via mirror."; + EXPECT_EQ( pixel( 31, 31 ), 0 ) << "Corner should miss on both backends."; + + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} diff --git a/test/HybridMixedTracingTest.h b/test/HybridMixedTracingTest.h new file mode 100644 index 0000000..de7cfb7 --- /dev/null +++ b/test/HybridMixedTracingTest.h @@ -0,0 +1,7 @@ +#pragma once + +#include "HybridGeometryTest.h" + +class HybridMixedTracingTest : public HybridGeometryTest +{ +}; diff --git a/test/HybridSceneTest.cpp b/test/HybridSceneTest.cpp new file mode 100644 index 0000000..f7e65e3 --- /dev/null +++ b/test/HybridSceneTest.cpp @@ -0,0 +1,339 @@ +#include "HybridSceneTest.h" + +#include + +#include + +extern const char* g_hip_paths[]; +extern const char* g_hiprtc_paths[]; + +namespace +{ +void assertOroSuccess( oroError err ) +{ + if ( err != oroSuccess ) + { + const char* msg = nullptr; + oroGetErrorString( err, &msg ); + FAIL() << "Orochi error: " << ( msg != nullptr ? msg : "unknown" ); + } +} +} // namespace + +void HybridSceneTest::freeManagedAllocs() +{ + for ( void* ptr : m_managedAllocs ) + assertOroSuccess( oroFree( ptr ) ); + m_managedAllocs.clear(); +} + +void HybridSceneTest::SetUp() +{ + oroInitialize( (oroApi)( ORO_API_HIP | ORO_API_CUDA ), 0, g_hip_paths, g_hiprtc_paths ); + assertOroSuccess( oroInit( 0 ) ); + assertOroSuccess( oroDeviceGet( &m_oroDevice, 0 ) ); + assertOroSuccess( oroCtxCreate( &m_oroCtx, 0, m_oroDevice ) ); + assertOroSuccess( oroCtxSetCurrent( m_oroCtx ) ); + + oroDeviceProp props{}; + assertOroSuccess( oroGetDeviceProperties( &props, m_oroDevice ) ); + + hiprtContextCreationInput ctxInput{}; + ctxInput.numCpuThreads = 1; + ctxInput.ctxt = oroGetRawCtx( m_oroCtx ); + ctxInput.device = oroGetRawDevice( m_oroDevice ); + ctxInput.deviceType = ( std::string( props.name ).find( "NVIDIA" ) != std::string::npos ) + ? static_cast( hiprtDeviceNVIDIA | hiprtDeviceCPU ) + : static_cast( hiprtDeviceAMD | hiprtDeviceCPU ); + + ASSERT_EQ( hiprtCreateContext( HIPRT_API_VERSION, ctxInput, m_context ), hiprtSuccess ) + << "Failed to create hybrid context."; + ASSERT_TRUE( m_context != nullptr ) << "Context pointer is null after creation."; +} + +void HybridSceneTest::TearDown() +{ + if ( m_context != nullptr ) + EXPECT_EQ( hiprtDestroyContext( m_context ), hiprtSuccess ) + << "Failed to cleanly destroy hybrid context."; + m_context = nullptr; + + freeManagedAllocs(); + + if ( m_oroCtx != nullptr ) + assertOroSuccess( oroCtxDestroy( m_oroCtx ) ); + m_oroCtx = nullptr; +} + +hiprtGeometry HybridSceneTest::buildTriangleGeometry( + std::vector& vertices, + std::vector& indices ) +{ + hiprtFloat3* hostVertices = nullptr; + uint32_t* hostIndices = nullptr; + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostVertices ), 3 * sizeof( hiprtFloat3 ), oroMemAttachGlobal ) ); + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostIndices ), 3 * sizeof( uint32_t ), oroMemAttachGlobal ) ); + m_managedAllocs.push_back( hostVertices ); + m_managedAllocs.push_back( hostIndices ); + + hostVertices[0] = { 0.0f, 0.0f, 0.0f }; + hostVertices[1] = { 1.0f, 0.0f, 0.0f }; + hostVertices[2] = { 0.0f, 1.0f, 0.0f }; + hostIndices[0] = 0; + hostIndices[1] = 1; + hostIndices[2] = 2; + + vertices = { hostVertices[0], hostVertices[1], hostVertices[2] }; + indices = { hostIndices[0], hostIndices[1], hostIndices[2] }; + + hiprtTriangleMeshPrimitive mesh{}; + mesh.vertices = hostVertices; + mesh.vertexCount = 3; + mesh.vertexStride = sizeof( hiprtFloat3 ); + mesh.triangleIndices = hostIndices; + mesh.triangleCount = 1; + mesh.triangleStride = 3 * sizeof( uint32_t ); + + hiprtGeometryBuildInput in{}; + in.type = hiprtPrimitiveTypeTriangleMesh; + in.primitive.triangleMesh = mesh; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtGeometry geom = nullptr; + EXPECT_EQ( hiprtCreateGeometry( m_context, in, opts, geom ), hiprtSuccess ) + << "hiprtCreateGeometry failed."; + EXPECT_TRUE( geom != nullptr ) << "Returned geometry handle is null."; + EXPECT_EQ( + hiprtBuildGeometry( + m_context, hiprtBuildOperationBuild, in, opts, allocGeometryBuildTemp( in, opts ), nullptr, geom ), + hiprtSuccess ) + << "hiprtBuildGeometry failed."; + return geom; +} + +void HybridSceneTest::stageSceneBuildInput( + hiprtSceneBuildInput& in, + const std::vector& instances, + const std::vector& frames ) +{ + ASSERT_FALSE( instances.empty() ); + ASSERT_EQ( instances.size(), frames.size() ); + + hiprtInstance* hostInstances = nullptr; + hiprtFrameSRT* hostFrames = nullptr; + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostInstances ), + instances.size() * sizeof( hiprtInstance ), + oroMemAttachGlobal ) ); + assertOroSuccess( oroMallocManaged( + reinterpret_cast( &hostFrames ), + frames.size() * sizeof( hiprtFrameSRT ), + oroMemAttachGlobal ) ); + m_managedAllocs.push_back( hostInstances ); + m_managedAllocs.push_back( hostFrames ); + + for ( size_t i = 0; i < instances.size(); ++i ) + hostInstances[i] = instances[i]; + for ( size_t i = 0; i < frames.size(); ++i ) + hostFrames[i] = frames[i]; + + in.instances = hostInstances; + in.instanceCount = static_cast( instances.size() ); + in.instanceFrames = hostFrames; + in.frameCount = static_cast( frames.size() ); + in.frameType = hiprtFrameTypeSRT; +} + +hiprtDevicePtr HybridSceneTest::allocGeometryBuildTemp( + const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions& buildOptions ) +{ + size_t tempSize = 0; + hiprtDevicePtr temp = nullptr; + if ( hiprtGetGeometryBuildTemporaryBufferSize( m_context, buildInput, buildOptions, tempSize ) != hiprtSuccess ) + return nullptr; + if ( tempSize == 0 ) + return nullptr; + + assertOroSuccess( oroMalloc( reinterpret_cast( &temp ), tempSize ) ); + m_managedAllocs.push_back( reinterpret_cast( temp ) ); + return temp; +} + +hiprtDevicePtr HybridSceneTest::allocSceneBuildTemp( + const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions& buildOptions ) +{ + size_t tempSize = 0; + hiprtDevicePtr temp = nullptr; + if ( hiprtGetSceneBuildTemporaryBufferSize( m_context, buildInput, buildOptions, tempSize ) != hiprtSuccess ) + return nullptr; + if ( tempSize == 0 ) + return nullptr; + + assertOroSuccess( oroMalloc( reinterpret_cast( &temp ), tempSize ) ); + m_managedAllocs.push_back( reinterpret_cast( temp ) ); + return temp; +} + +namespace +{ +hiprtFrameSRT makeSrt( float tx, float ty, float tz ) +{ + hiprtFrameSRT f{}; + f.rotation = { 0.0f, 0.0f, 0.0f, 0.0f }; + f.scale = { 1.0f, 1.0f, 1.0f }; + f.translation = { tx, ty, tz }; + f.time = 0.0f; + return f; +} + +bool traceScene( + hiprtContext ctx, hiprtScene scene, float ox, float oy, float oz, float dx, float dy, float dz ) +{ + hiprtRay ray{}; + ray.origin = { ox, oy, oz }; + ray.minT = 0.0f; + ray.direction = { dx, dy, dz }; + ray.maxT = 1e30f; + + hiprtHit hit{}; + hiprtGeomTraversalClosestCPU::traceBatch( ctx, scene, &ray, &hit, 1 ); + return hit.instanceID != hiprtInvalidValue; +} +} // namespace + +TEST_F( HybridSceneTest, CreateAndDestroyScene ) +{ + hiprtSceneBuildInput in{}; + in.instanceCount = 0; + in.frameType = hiprtFrameTypeSRT; + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ) + << "hiprtCreateScene failed on hybrid context."; + ASSERT_TRUE( scene != nullptr ) << "Returned scene handle is null."; + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ) + << "hiprtDestroyScene failed to release GPU/CPU mirror."; +} + +TEST_F( HybridSceneTest, BuildSceneWithSingleInstanceIsHittableViaCpuMirror ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 0.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + stageSceneBuildInput( in, instances, frames ); + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( + m_context, hiprtBuildOperationBuild, in, opts, allocSceneBuildTemp( in, opts ), nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "CPU traceBatch on GPU scene handle should hit the Embree mirror."; + EXPECT_FALSE( traceScene( m_context, scene, 5.0f, 5.0f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Ray outside instance should miss."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridSceneTest, BuildSceneWithTranslatedInstance ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 3.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + stageSceneBuildInput( in, instances, frames ); + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( + m_context, hiprtBuildOperationBuild, in, opts, allocSceneBuildTemp( in, opts ), nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( m_context, scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Translated instance should be hit at x=3.25 via CPU mirror."; + EXPECT_FALSE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Original position should miss after +3X translation."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridSceneTest, UpdateSceneRebuildsCpuMirror ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + ASSERT_TRUE( geom != nullptr ); + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 0.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + stageSceneBuildInput( in, instances, frames ); + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + ASSERT_EQ( hiprtCreateScene( m_context, in, opts, scene ), hiprtSuccess ); + ASSERT_EQ( + hiprtBuildScene( + m_context, hiprtBuildOperationBuild, in, opts, allocSceneBuildTemp( in, opts ), nullptr, scene ), + hiprtSuccess ); + + EXPECT_TRUE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Before update hit at x=0.25."; + + auto* stagedFrames = reinterpret_cast( in.instanceFrames ); + stagedFrames[0] = makeSrt( 3.0f, 0.0f, 0.0f ); + ASSERT_EQ( + hiprtBuildScene( + m_context, hiprtBuildOperationUpdate, in, opts, allocSceneBuildTemp( in, opts ), nullptr, scene ), + hiprtSuccess ) + << "hiprtBuildScene (Update) failed."; + + EXPECT_FALSE( traceScene( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "After update old position should miss."; + EXPECT_TRUE( traceScene( m_context, scene, 3.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "After update new position x=3.25 should hit."; + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} diff --git a/test/HybridSceneTest.h b/test/HybridSceneTest.h new file mode 100644 index 0000000..43b4839 --- /dev/null +++ b/test/HybridSceneTest.h @@ -0,0 +1,34 @@ +#pragma once + +#include +#include +#include +#include + +class HybridSceneTest : public ::testing::Test +{ + protected: + void SetUp() override; + void TearDown() override; + + hiprtGeometry buildTriangleGeometry( + std::vector& vertices, + std::vector& indices ); + + void stageSceneBuildInput( + hiprtSceneBuildInput& in, + const std::vector& instances, + const std::vector& frames ); + + hiprtDevicePtr allocGeometryBuildTemp( + const hiprtGeometryBuildInput& buildInput, const hiprtBuildOptions& buildOptions ); + hiprtDevicePtr allocSceneBuildTemp( + const hiprtSceneBuildInput& buildInput, const hiprtBuildOptions& buildOptions ); + + hiprtContext m_context = nullptr; + oroCtx m_oroCtx = nullptr; + oroDevice m_oroDevice = 0; + + void freeManagedAllocs(); + std::vector m_managedAllocs; +}; diff --git a/test/HybridTraceBatchTest.cpp b/test/HybridTraceBatchTest.cpp new file mode 100644 index 0000000..593e4cb --- /dev/null +++ b/test/HybridTraceBatchTest.cpp @@ -0,0 +1,209 @@ +#include "HybridTraceBatchTest.h" + +#include +#include + +#include +#include +#include +#include +#include + +extern const char* g_hip_paths[]; +extern const char* g_hiprtc_paths[]; + +namespace +{ +std::filesystem::path getRootDir() +{ + const char* env = std::getenv( "HIPRT_PATH" ); + return env != nullptr ? std::filesystem::path( env ) : std::filesystem::path( ".." ); +} + +bool readKernelSource( const std::filesystem::path& srcPath, std::string& sourceCode ) +{ + std::ifstream f( srcPath ); + if ( !f.is_open() ) + return false; + sourceCode.assign( std::istreambuf_iterator( f ), std::istreambuf_iterator() ); + return true; +} + +hiprtApiFunction buildHybridSceneTraceBatchKernel( hiprtContext context ) +{ + const auto kernelPath = getRootDir() / "hiprt/kernels/HybridSceneTraceBatchKernel.h"; + std::string sourceCode; + const std::string moduleName = kernelPath.string(); + if ( !readKernelSource( kernelPath, sourceCode ) ) + return nullptr; + + const char* funcName = "HybridSceneTraceBatchKernel"; + + std::vector optionStorage; + optionStorage.push_back( "-I" + getRootDir().string() ); + if ( oroGetCurAPI( 0 ) == ORO_API_HIP ) + optionStorage.push_back( "-ffast-math" ); + else + optionStorage.push_back( "--use_fast_math" ); + + std::vector options; + options.reserve( optionStorage.size() ); + for ( const auto& opt : optionStorage ) + options.push_back( opt.c_str() ); + + hiprtApiFunction funcApi = nullptr; + if ( hiprtBuildTraceKernels( + context, 1, &funcName, sourceCode.c_str(), moduleName.c_str(), 0, nullptr, nullptr, + static_cast( options.size() ), options.data(), 0, 0, nullptr, &funcApi, nullptr, + true ) != hiprtSuccess ) + return nullptr; + + return funcApi; +} + +hiprtFrameSRT makeSrt( float tx, float ty, float tz ) +{ + hiprtFrameSRT f{}; + f.rotation = { 0.0f, 0.0f, 0.0f, 0.0f }; + f.scale = { 1.0f, 1.0f, 1.0f }; + f.translation = { tx, ty, tz }; + f.time = 0.0f; + return f; +} + +bool hitsEqual( const hiprtHit& a, const hiprtHit& b ) +{ + if ( a.hasHit() != b.hasHit() ) + return false; + if ( !a.hasHit() ) + return true; + if ( a.primID != b.primID || a.instanceID != b.instanceID ) + return false; + return std::fabs( a.t - b.t ) < 1e-4f; +} + +void fillTestRays( hiprtRay* rays, uint32_t count ) +{ + for ( uint32_t i = 0; i < count; ++i ) + { + hiprtRay& ray = rays[i]; + ray.minT = 0.0f; + ray.maxT = 1e30f; + ray.direction = { 0.0f, 0.0f, 1.0f }; + + if ( i % 2 == 0 ) + ray.origin = { 0.25f, 0.25f, -1.0f }; + else + ray.origin = { 5.0f, 5.0f, -1.0f }; + } +} +} // namespace + +hiprtScene HybridTraceBatchTest::buildSingleInstanceScene( hiprtGeometry& geomOut ) +{ + std::vector vertices; + std::vector indices; + hiprtGeometry geom = buildTriangleGeometry( vertices, indices ); + if ( geom == nullptr ) + return nullptr; + + std::vector instances( 1 ); + instances[0].type = hiprtInstanceTypeGeometry; + instances[0].geometry = geom; + + std::vector frames = { makeSrt( 0.0f, 0.0f, 0.0f ) }; + + hiprtSceneBuildInput in{}; + stageSceneBuildInput( in, instances, frames ); + + hiprtBuildOptions opts{}; + opts.buildFlags = hiprtBuildFlagBitPreferFastBuild; + + hiprtScene scene = nullptr; + if ( hiprtCreateScene( m_context, in, opts, scene ) != hiprtSuccess ) + return nullptr; + if ( hiprtBuildScene( m_context, hiprtBuildOperationBuild, in, opts, allocSceneBuildTemp( in, opts ), nullptr, scene ) != + hiprtSuccess ) + return nullptr; + + geomOut = geom; + return scene; +} + +TEST_F( HybridTraceBatchTest, SplitBatchMatchesCpuReference ) +{ + constexpr uint32_t rayCount = 16; + + hiprtGeometry geom = nullptr; + hiprtScene scene = buildSingleInstanceScene( geom ); + ASSERT_TRUE( scene != nullptr ); + ASSERT_TRUE( geom != nullptr ); + + hiprtApiFunction traceKernel = buildHybridSceneTraceBatchKernel( m_context ); + ASSERT_TRUE( traceKernel != nullptr ) << "Failed to compile HybridSceneTraceBatchKernel."; + + hiprtRay* rays = nullptr; + hiprtHit* refHits = nullptr; + hiprtHit* hybridHits = nullptr; + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &rays ), rayCount * sizeof( hiprtRay ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( + reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( + reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + m_managedAllocs.push_back( rays ); + m_managedAllocs.push_back( refHits ); + m_managedAllocs.push_back( hybridHits ); + + fillTestRays( rays, rayCount ); + std::memset( refHits, 0, rayCount * sizeof( hiprtHit ) ); + std::memset( hybridHits, 0, rayCount * sizeof( hiprtHit ) ); + + hiprtGeomTraversalClosestCPU::traceBatch( m_context, scene, rays, refHits, rayCount ); + + hiprtHybridTraceConfig cfg{}; + cfg.gpuFraction = 0.5f; + cfg.minCpuBatch = 1u; + + hiprtHybridTraceGpuInput gpuInput{}; + gpuInput.traceKernel = traceKernel; + + ASSERT_EQ( + hiprtTraceHybridClosest( m_context, scene, cfg, gpuInput, rays, hybridHits, rayCount, nullptr ), + hiprtSuccess ); + + for ( uint32_t i = 0; i < rayCount; ++i ) + { + EXPECT_TRUE( hitsEqual( hybridHits[i], refHits[i] ) ) + << "Ray " << i << " mismatch (even=hit, odd=miss)."; + } + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridTraceBatchTest, RejectsNonHybridContext ) +{ + hiprtContextCreationInput cpuInput{}; + cpuInput.deviceType = hiprtDeviceCPU; + cpuInput.numCpuThreads = 1; + + hiprtContext cpuContext = nullptr; + ASSERT_EQ( hiprtCreateContext( HIPRT_API_VERSION, cpuInput, cpuContext ), hiprtSuccess ); + + hiprtRay ray{}; + hiprtHit hit{}; + hiprtHybridTraceConfig cfg{}; + hiprtHybridTraceGpuInput gpuInput{ nullptr, {} }; + + EXPECT_EQ( + hiprtTraceHybridClosest( cpuContext, nullptr, cfg, gpuInput, &ray, &hit, 1, nullptr ), + hiprtErrorInvalidParameter ); + + EXPECT_EQ( hiprtDestroyContext( cpuContext ), hiprtSuccess ); +} diff --git a/test/HybridTraceBatchTest.h b/test/HybridTraceBatchTest.h new file mode 100644 index 0000000..606faea --- /dev/null +++ b/test/HybridTraceBatchTest.h @@ -0,0 +1,9 @@ +#pragma once + +#include "HybridSceneTest.h" + +class HybridTraceBatchTest : public HybridSceneTest +{ + protected: + hiprtScene buildSingleInstanceScene( hiprtGeometry& geomOut ); +}; diff --git a/test/kernels/HybridRowsKernel.h b/test/kernels/HybridRowsKernel.h new file mode 100644 index 0000000..2b8f7fb --- /dev/null +++ b/test/kernels/HybridRowsKernel.h @@ -0,0 +1,43 @@ +#pragma once + +#include + +#ifndef BLOCK_SIZE +#define BLOCK_SIZE 1 +#endif + +#ifndef SHARED_STACK_SIZE +#define SHARED_STACK_SIZE 1 +#endif + +extern "C" __global__ void HybridRowsKernel( + hiprtGeometry geom, + uint8_t* image, + uint32_t width, + uint32_t height, + uint32_t yBegin, + uint32_t yEnd, + uint8_t hitValue ) +{ + const uint32_t x = blockIdx.x * blockDim.x + threadIdx.x; + const uint32_t y = blockIdx.y * blockDim.y + threadIdx.y; + if ( x >= width || y >= height || y < yBegin || y >= yEnd ) + return; + + const uint32_t index = x + y * width; + + hiprtRay ray{}; + ray.origin = { x / static_cast( width ), y / static_cast( height ), -1.0f }; + ray.direction = { 0.0f, 0.0f, 1.0f }; + ray.minT = 0.0f; + ray.maxT = 1e30f; + + hiprtGeomTraversalClosest tr( geom, ray ); + const hiprtHit hit = tr.getNextHit(); + + const uint8_t shade = hit.hasHit() ? hitValue : 0; + image[index * 4 + 0] = shade; + image[index * 4 + 1] = shade; + image[index * 4 + 2] = shade; + image[index * 4 + 3] = 255; +} From eed91421a92df978b64d767275c6bbf696739af3 Mon Sep 17 00:00:00 2001 From: Krzysztof Date: Thu, 25 Jun 2026 01:28:02 +0200 Subject: [PATCH 7/7] hybrid and CPU API is now similar to GPU API --- hiprt/hiprt_cpu.h | 89 +++++- hiprt/hiprt_hybrid.h | 99 +++++-- hiprt/impl/CpuDataRegistry.h | 11 +- hiprt/impl/HybridContext.cpp | 123 ++++++-- hiprt/impl/HybridContext.h | 42 +-- hiprt/impl/hiprt_hybrid.cpp | 281 ++++++++++++++++--- hiprt/kernels/HybridGeomAnyHitBatchKernel.h | 43 +++ hiprt/kernels/HybridGeomTraceBatchKernel.h | 43 +++ hiprt/kernels/HybridSceneAnyHitBatchKernel.h | 43 +++ hiprt/kernels/HybridSceneTraceBatchKernel.h | 6 +- test/HybridTraceBatchTest.cpp | 179 +++++++++++- 11 files changed, 802 insertions(+), 157 deletions(-) create mode 100644 hiprt/kernels/HybridGeomAnyHitBatchKernel.h create mode 100644 hiprt/kernels/HybridGeomTraceBatchKernel.h create mode 100644 hiprt/kernels/HybridSceneAnyHitBatchKernel.h diff --git a/hiprt/hiprt_cpu.h b/hiprt/hiprt_cpu.h index 5c7678c..dbbb93c 100644 --- a/hiprt/hiprt_cpu.h +++ b/hiprt/hiprt_cpu.h @@ -30,9 +30,7 @@ #include #include -// --------------------------------------------------------------------------- -// Internal helpers – not for external consumption -// --------------------------------------------------------------------------- +// Internal helpers namespace hiprt_cpu_detail { @@ -91,19 +89,51 @@ inline void traceGeometryBatch( } } +inline RTCRay makeRTCRay( const hiprtRay& ray ) noexcept +{ + RTCRay r{}; + r.org_x = ray.origin.x; + r.org_y = ray.origin.y; + r.org_z = ray.origin.z; + r.dir_x = ray.direction.x; + r.dir_y = ray.direction.y; + r.dir_z = ray.direction.z; + r.tnear = ray.minT; + r.tfar = ray.maxT; + r.mask = static_cast( -1 ); + r.flags = 0; + return r; +} + +// Any-hit: rtcOccluded1 result encoded in hit.hasHit(). +inline void traceGeometryAnyHitBatch( + RTCScene scene, const hiprtRay* rays, hiprtHit* hits, uint32_t count ) noexcept +{ + if ( scene == nullptr ) + { + for ( uint32_t i = 0; i < count; ++i ) + hits[i] = hiprtHit{}; + return; + } + + for ( uint32_t i = 0; i < count; ++i ) + { + RTCRay r = makeRTCRay( rays[i] ); + rtcOccluded1( scene, &r ); + + hiprtHit hit{}; + if ( r.tfar < 0.0f ) // occluded + { + hit.primID = 0; + hit.instanceID = hiprtInvalidValue; + } + hits[i] = hit; + } +} + } // namespace hiprt_cpu_detail -// --------------------------------------------------------------------------- -// Public CPU traversal API -// --------------------------------------------------------------------------- - -/** \brief Host-side batch traversal that finds the closest hit. - * - * In CPU-only mode the handle is a direct CpuGeometryData / CpuSceneData pointer. - * In hybrid mode (hiprtDeviceAMD | hiprtDeviceCPU or hiprtDeviceNVIDIA | hiprtDeviceCPU) - * geometry and scene handles are GPU handles; the CPU Embree mirror is resolved - * through the context side table. - */ +// Host-side CPU batch traversal (Embree). class hiprtGeomTraversalClosestCPU { public: @@ -133,3 +163,34 @@ class hiprtGeomTraversalClosestCPU ( data != nullptr ) ? data->rtcScene : nullptr, rays, hits, count ); } }; + +// Any-hit batch traversal (shadow/occlusion rays). +class hiprtGeomTraversalAnyHitCPU +{ + public: + static void traceBatch( + hiprtContext ctx, + hiprtGeometry geom, + const hiprtRay* rays, + hiprtHit* hits, + uint32_t count ) noexcept + { + (void)ctx; + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromGeometry( geom ) ); + hiprt_cpu_detail::traceGeometryAnyHitBatch( + ( data != nullptr ) ? data->rtcScene : nullptr, rays, hits, count ); + } + + static void traceBatch( + hiprtContext ctx, + hiprtScene scene, + const hiprtRay* rays, + hiprtHit* hits, + uint32_t count ) noexcept + { + (void)ctx; + auto* data = reinterpret_cast( hiprtGetInternalCpuDataFromScene( scene ) ); + hiprt_cpu_detail::traceGeometryAnyHitBatch( + ( data != nullptr ) ? data->rtcScene : nullptr, rays, hits, count ); + } +}; diff --git a/hiprt/hiprt_hybrid.h b/hiprt/hiprt_hybrid.h index d2de7f7..5bdecb5 100644 --- a/hiprt/hiprt_hybrid.h +++ b/hiprt/hiprt_hybrid.h @@ -26,45 +26,20 @@ #include -/** \brief Scheduling parameters for hybrid closest-hit tracing. - */ struct hiprtHybridTraceConfig { - /*!< Fraction of rays dispatched to the GPU in [0, 1]. */ - float gpuFraction = 0.8f; - /*!< Minimum number of rays kept for the CPU path when both backends are used. */ + float gpuFraction = 0.8f; uint32_t minCpuBatch = 1024u; }; -/** \brief GPU resources required for hybrid scene tracing. - * - * Compile hiprt/kernels/HybridSceneTraceBatchKernel.h with hiprtBuildTraceKernels - * and pass the resulting function handle here. globalStack is unused by the - * default batch kernel but reserved for advanced kernels such as TraceKernel. - */ +// Optional: pass your own GPU kernel. Default overloads compile one internally. struct hiprtHybridTraceGpuInput { hiprtApiFunction traceKernel = nullptr; hiprtGlobalStackBuffer globalStack = {}; }; -/** \brief Trace a batch of closest-hit scene rays on GPU and CPU concurrently. - * - * Requires a hybrid context (hiprtDeviceAMD | hiprtDeviceCPU or - * hiprtDeviceNVIDIA | hiprtDeviceCPU). The first gpuCount rays are traced on - * the GPU; the remainder on the CPU Embree mirror. rays and hits must be - * host-visible (managed/UVA or pinned mapped) so both backends can access them. - * - * \param context Hybrid HIPRT context. - * \param scene GPU scene handle with a CPU mirror. - * \param config Split policy between GPU and CPU. - * \param gpuInput Precompiled GPU trace kernel (and optional stack buffer). - * \param rays Host-visible ray array of length rayCount. - * \param hits Host-visible hit array of length rayCount. - * \param rayCount Number of rays to trace. - * \param stream GPU stream for the device launch (0 for default stream). - * \return hiprtSuccess or an error code. - */ +// Hybrid closest-hit (scene). Requires hybrid context; rays/hits must be host-visible. HIPRT_API hiprtError hiprtTraceHybridClosest( hiprtContext context, hiprtScene scene, @@ -74,3 +49,71 @@ HIPRT_API hiprtError hiprtTraceHybridClosest( hiprtHit* hits, uint32_t rayCount, hiprtApiStream stream ); + +HIPRT_API hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +// Hybrid any-hit (scene/geometry). Only hit.hasHit() is meaningful on CPU path. +HIPRT_API hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +HIPRT_API hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +// Simple API: kernel compiled and cached inside the hybrid context. +HIPRT_API hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +HIPRT_API hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +HIPRT_API hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); + +HIPRT_API hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ); diff --git a/hiprt/impl/CpuDataRegistry.h b/hiprt/impl/CpuDataRegistry.h index b87aa7e..540830b 100644 --- a/hiprt/impl/CpuDataRegistry.h +++ b/hiprt/impl/CpuDataRegistry.h @@ -7,16 +7,7 @@ namespace hiprt struct CpuGeometryData; struct CpuSceneData; -/** - * Global CPU-data registry. - * - * GPU pointers (oroDeviceptr cast to hiprtGeometry/hiprtScene) are unique - * process-wide, so a single flat map is sufficient. Pure-CPU handles are the - * CpuGeometryData/CpuSceneData pointers themselves — they are also unique. - * Both can therefore share the same table without ambiguity. - * - * All functions are thread-safe via an internal mutex. - */ +// Maps GPU and CPU handles for hybrid traversal. Thread-safe. void registerCpuGeom( hiprtGeometry handle, CpuGeometryData* data ); void unregisterCpuGeom( hiprtGeometry handle ); diff --git a/hiprt/impl/HybridContext.cpp b/hiprt/impl/HybridContext.cpp index 32c536c..1cc6324 100644 --- a/hiprt/impl/HybridContext.cpp +++ b/hiprt/impl/HybridContext.cpp @@ -16,11 +16,32 @@ void syncGpu( oroStream stream ) if ( err != oroSuccess ) throw std::runtime_error( "HybridContext: GPU synchronization failed before CPU mirror build" ); } + +[[noreturn]] void throwHybridNotImplemented( const char* what ) +{ + std::cerr << "[HybridContext] not implemented: " << what << std::endl; + throw std::runtime_error( what ); +} + +#ifndef NDEBUG +// Debug: warn if Embree input buffer is device-only memory. +void warnIfDeviceOnly( const void* ptr, const char* what ) +{ + if ( ptr == nullptr ) return; + oroPointerAttribute_t attr{}; + if ( oroPointerGetAttributes( &attr, const_cast( ptr ) ) == oroSuccess && attr.type == oroMemoryTypeDevice ) + { + std::cerr << "[HybridContext] WARNING: " << what + << " points to device-only memory; Embree (CPU) cannot read it. " + << "Use oroMallocManaged / host-visible memory for CPU/hybrid mode.\n"; + } +} +#else +inline void warnIfDeviceOnly( const void*, const char* ) {} +#endif } // namespace -// --------------------------------------------------------------------------- -// Construction / destruction -// --------------------------------------------------------------------------- +// Construction HybridContext::HybridContext( const hiprtContextCreationInput& input ) : m_gpu( input ), m_cpu( input ) @@ -28,15 +49,48 @@ HybridContext::HybridContext( const hiprtContextCreationInput& input ) std::cerr << "[HybridContext] constructed - GPU BVH + Embree CPU mirror active.\n"; } -HybridContext::~HybridContext() +HybridContext::~HybridContext() {} + +oroFunction HybridContext::getBatchKernel( HybridBatchKernel kind ) { - // Both sub-contexts clean up their own resources. - // Maps are empty at this point (destroy calls already handled them). + const int idx = static_cast( kind ); + + std::lock_guard lock( m_batchKernelMutex ); + if ( m_batchKernels[idx] != nullptr ) + return m_batchKernels[idx]; + + struct KernelDef + { + const char* include; + const char* func; + }; + static const KernelDef defs[static_cast( HybridBatchKernel::Count )] = { + { "#include \n", "HybridSceneTraceBatchKernel" }, + { "#include \n", "HybridGeomTraceBatchKernel" }, + { "#include \n", "HybridSceneAnyHitBatchKernel" }, + { "#include \n", "HybridGeomAnyHitBatchKernel" }, + }; + + std::vector funcNames = { defs[idx].func }; + std::string src = defs[idx].include; + std::vector headers; + std::vector includeNames; + std::vector options; + std::vector funcNameSets; + std::vector functions; + oroModule module = nullptr; + + m_gpu.buildKernels( + funcNames, src, defs[idx].func, headers, includeNames, options, 0, 0, funcNameSets, functions, module, true ); + + if ( functions.empty() || functions[0] == nullptr ) + throw std::runtime_error( "HybridContext: failed to compile internal batch kernel" ); + + m_batchKernels[idx] = functions[0]; + return m_batchKernels[idx]; } -// --------------------------------------------------------------------------- // Geometry -// --------------------------------------------------------------------------- std::vector HybridContext::createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) @@ -51,7 +105,7 @@ std::vector HybridContext::createGeometries( registerCpuGeom( gpuHandles[i], cpuData ); } - return gpuHandles; // return GPU handles; CPU side is hidden in the maps. + return gpuHandles; } void HybridContext::destroyGeometries( const std::vector& gpuGeometries ) @@ -80,7 +134,7 @@ void HybridContext::dispatchCpuGeomBuild( const std::vector& gpuBuffers, bool isUpdate ) { - // Translate GPU buffers -> CPU buffers for CpuContext + // Map GPU buffers to CPU buffers. std::vector cpuBuffers; cpuBuffers.reserve( gpuBuffers.size() ); for ( auto gp : gpuBuffers ) @@ -91,6 +145,20 @@ void HybridContext::dispatchCpuGeomBuild( cpuBuffers.push_back( reinterpret_cast( it->second ) ); } + // Debug: check host visibility of primitive data. + for ( const auto& in : buildInputs ) + { + if ( in.type == hiprtPrimitiveTypeTriangleMesh ) + { + warnIfDeviceOnly( in.primitive.triangleMesh.vertices, "geometry vertices" ); + warnIfDeviceOnly( in.primitive.triangleMesh.triangleIndices, "geometry triangleIndices" ); + } + else if ( in.type == hiprtPrimitiveTypeAABBList ) + { + warnIfDeviceOnly( in.primitive.aabbList.aabbs, "geometry aabbs" ); + } + } + if ( isUpdate ) m_cpu.updateGeometries( buildInputs, {}, nullptr, nullptr, cpuBuffers ); else @@ -128,15 +196,12 @@ size_t HybridContext::getGeometriesBuildTempBufferSize( } std::vector HybridContext::compactGeometries( - const std::vector& geometries, oroStream stream ) + const std::vector& /*geometries*/, oroStream /*stream*/ ) { - // GPU compaction only in this pass; CPU mirror is unaffected. - return m_gpu.compactGeometries( geometries, stream ); + throwHybridNotImplemented( "compactGeometries is not supported in hybrid mode (CPU mirror rebuild required)" ); } -// --------------------------------------------------------------------------- // Scene -// --------------------------------------------------------------------------- std::vector HybridContext::createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) @@ -180,11 +245,9 @@ void HybridContext::dispatchCpuSceneBuild( const std::vector& gpuBuffers, bool isUpdate ) { - // Build CPU-side buffers and remap instance handles (GPU->CPU). - // We deep-copy each hiprtSceneBuildInput and patch its instances array. + // Remap GPU instance handles to CPU handles. std::vector cpuBuffers; std::vector cpuInputs; - // Storage for the remapped instances arrays (keep alive during the call). std::vector> remappedInstances( buildInputs.size() ); cpuBuffers.reserve( gpuBuffers.size() ); @@ -192,15 +255,17 @@ void HybridContext::dispatchCpuSceneBuild( for ( size_t i = 0; i < buildInputs.size(); ++i ) { - // Translate GPU buffer to CPU buffer. auto it = m_gpuToCpuScene.find( reinterpret_cast( gpuBuffers[i] ) ); if ( it == m_gpuToCpuScene.end() ) throw std::runtime_error( "HybridContext: GPU scene buffer not found in map" ); cpuBuffers.push_back( reinterpret_cast( it->second ) ); - // Copy the build input and remap instances. hiprtSceneBuildInput cpuIn = buildInputs[i]; + warnIfDeviceOnly( buildInputs[i].instances, "scene instances" ); + warnIfDeviceOnly( buildInputs[i].instanceTransformHeaders, "scene instanceTransformHeaders" ); + warnIfDeviceOnly( buildInputs[i].instanceFrames, "scene instanceFrames" ); + if ( buildInputs[i].instanceCount > 0 && buildInputs[i].instances != nullptr ) { const auto* srcInst = reinterpret_cast( buildInputs[i].instances ); @@ -266,14 +331,12 @@ size_t HybridContext::getScenesBuildTempBufferSize( } std::vector HybridContext::compactScenes( - const std::vector& scenes, oroStream stream ) + const std::vector& /*scenes*/, oroStream /*stream*/ ) { - return m_gpu.compactScenes( scenes, stream ); + throwHybridNotImplemented( "compactScenes is not supported in hybrid mode (CPU mirror rebuild required)" ); } -// --------------------------------------------------------------------------- -// GPU-authoritative delegators -// --------------------------------------------------------------------------- +// GPU delegators hiprtFuncTable HybridContext::createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) { @@ -302,9 +365,9 @@ void HybridContext::saveGeometry( hiprtGeometry inGeometry, const std::string& f m_gpu.saveGeometry( inGeometry, filename ); } -hiprtGeometry HybridContext::loadGeometry( const std::string& filename ) +hiprtGeometry HybridContext::loadGeometry( const std::string& /*filename*/ ) { - return m_gpu.loadGeometry( filename ); + throwHybridNotImplemented( "loadGeometry is not supported in hybrid mode (CPU mirror rebuild required)" ); } void HybridContext::saveScene( hiprtScene inScene, const std::string& filename ) @@ -312,9 +375,9 @@ void HybridContext::saveScene( hiprtScene inScene, const std::string& filename ) m_gpu.saveScene( inScene, filename ); } -hiprtScene HybridContext::loadScene( const std::string& filename ) +hiprtScene HybridContext::loadScene( const std::string& /*filename*/ ) { - return m_gpu.loadScene( filename ); + throwHybridNotImplemented( "loadScene is not supported in hybrid mode (CPU mirror rebuild required)" ); } void HybridContext::exportGeometryAabb( hiprtGeometry inGeometry, float3& outAabbMin, float3& outAabbMax ) @@ -366,9 +429,7 @@ void HybridContext::setLogLevel( hiprtLogLevel level ) m_gpu.setLogLevel( level ); } -// --------------------------------------------------------------------------- // CPU data accessors -// --------------------------------------------------------------------------- CpuGeometryData* HybridContext::getCpuGeom( hiprtGeometry geometry ) { diff --git a/hiprt/impl/HybridContext.h b/hiprt/impl/HybridContext.h index 5206b33..2148b0f 100644 --- a/hiprt/impl/HybridContext.h +++ b/hiprt/impl/HybridContext.h @@ -5,30 +5,31 @@ #include #include +#include #include namespace hiprt { -/** - * HybridContext — composes GpuContext + CpuContext. - * - * Every build operation is performed twice: - * 1. On the GPU (via GpuContext) — the GPU handle is returned to the caller. - * 2. On the CPU (via CpuContext / Embree) — the CpuGeometryData / CpuSceneData ptr - * is stored in the instance maps and registered in the global CpuDataRegistry - * so that hiprt_cpu.h traversal classes can resolve it by the GPU handle. - * - * This preserves 100 % binary compatibility with existing GPU-only code: the - * opaque handles visible to the application remain GPU device pointers. - */ +// GPU + CPU (Embree) context. Build ops update both sides; callers keep GPU handles. class HybridContext : public ContextBase { public: + enum class HybridBatchKernel : int + { + SceneClosest = 0, + GeomClosest, + SceneAnyHit, + GeomAnyHit, + Count + }; + explicit HybridContext( const hiprtContextCreationInput& input ); ~HybridContext() override; - // ----- Geometry -------------------------------------------------------- + oroFunction getBatchKernel( HybridBatchKernel kind ); + + // Geometry std::vector createGeometries( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; @@ -53,7 +54,7 @@ class HybridContext : public ContextBase std::vector compactGeometries( const std::vector& geometries, oroStream stream ) override; - // ----- Scene ----------------------------------------------------------- + // Scene std::vector createScenes( const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; @@ -78,7 +79,7 @@ class HybridContext : public ContextBase std::vector compactScenes( const std::vector& scenes, oroStream stream ) override; - // ----- GPU-authoritative operations (delegate to GpuContext) ----------- + // GPU ops (delegate to GpuContext) hiprtFuncTable createFuncTable( uint32_t numGeomTypes, uint32_t numRayTypes ) override; void setFuncTable( hiprtFuncTable funcTable, uint32_t geomType, uint32_t rayType, hiprtFuncDataSet set ) override; void destroyFuncTable( hiprtFuncTable funcTable ) override; @@ -121,18 +122,16 @@ class HybridContext : public ContextBase void setCacheDir( const std::filesystem::path& path ) override; void setLogLevel( hiprtLogLevel level ) override; - // ----- CPU-data accessors (used by hiprt_cpu.h via registry) ----------- + // CPU data accessors CpuGeometryData* getCpuGeom( hiprtGeometry geometry ) override; CpuSceneData* getCpuScene( hiprtScene scene ) override; private: - // Translate GPU geometry buffers to their CPU counterparts and call fn. void dispatchCpuGeomBuild( const std::vector& buildInputs, const std::vector& gpuBuffers, bool isUpdate ); - // Translate GPU scene buffers (remapping instance handles) and call fn. void dispatchCpuSceneBuild( const std::vector& buildInputs, const std::vector& gpuBuffers, @@ -141,8 +140,11 @@ class HybridContext : public ContextBase GpuContext m_gpu; CpuContext m_cpu; - std::unordered_map m_gpuToCpuGeom; // GPU handle -> CPU handle - std::unordered_map m_gpuToCpuScene; // GPU handle -> CPU handle + std::unordered_map m_gpuToCpuGeom; + std::unordered_map m_gpuToCpuScene; + + std::mutex m_batchKernelMutex; + oroFunction m_batchKernels[static_cast( HybridBatchKernel::Count )] = {}; }; } // namespace hiprt diff --git a/hiprt/impl/hiprt_hybrid.cpp b/hiprt/impl/hiprt_hybrid.cpp index e83f003..78ab893 100644 --- a/hiprt/impl/hiprt_hybrid.cpp +++ b/hiprt/impl/hiprt_hybrid.cpp @@ -28,6 +28,9 @@ #include +#include +#include + using namespace hiprt; namespace @@ -57,47 +60,78 @@ uint32_t computeGpuRayCount( uint32_t rayCount, const hiprtHybridTraceConfig& co return gpuCount; } -hiprtError launchSceneTraceBatchKernel( - const hiprtHybridTraceGpuInput& gpuInput, - hiprtScene scene, - uint32_t numRays, - hiprtRay* rays, - hiprtHit* hits, - hiprtApiStream stream ) +// Batch kernel: (handle, numRays, rays, hits). +hiprtError launchBatchKernel( + oroFunction func, + void* handle, + uint32_t numRays, + hiprtRay* rays, + hiprtHit* hits, + hiprtApiStream stream ) { - if ( gpuInput.traceKernel == nullptr ) + if ( func == nullptr ) return hiprtErrorInvalidParameter; + if ( numRays == 0 ) + return hiprtSuccess; constexpr uint32_t blockSize = 256u; const uint32_t gridSize = ( numRays + blockSize - 1 ) / blockSize; - hiprtApiFunction funcApi = gpuInput.traceKernel; - oroFunction func = *reinterpret_cast( &funcApi ); - - void* args[] = { &scene, &numRays, &rays, &hits }; + void* args[] = { &handle, &numRays, &rays, &hits }; const oroError err = oroModuleLaunchKernel( func, gridSize, 1, 1, blockSize, 1, 1, 0, reinterpret_cast( stream ), args, 0 ); return ( err == oroSuccess ) ? hiprtSuccess : hiprtErrorInternal; } -} // namespace -hiprtError hiprtTraceHybridClosest( +hiprtError syncStream( hiprtApiStream stream ) +{ + const oroStream gpuStream = reinterpret_cast( stream ); + const oroError err = ( gpuStream != nullptr ) ? oroStreamSynchronize( gpuStream ) : oroDeviceSynchronize(); + return ( err == oroSuccess ) ? hiprtSuccess : hiprtErrorInternal; +} + +// Caller kernel, or compile/cache internal kernel for kernelKind. +oroFunction resolveBatchKernel( + HybridContext* hybrid, + const hiprtHybridTraceGpuInput* gpuInput, + HybridContext::HybridBatchKernel kernelKind ) +{ + if ( gpuInput != nullptr && gpuInput->traceKernel != nullptr ) + { + hiprtApiFunction funcApi = gpuInput->traceKernel; + return *reinterpret_cast( &funcApi ); + } + + try + { + return hybrid->getBatchKernel( kernelKind ); + } + catch ( const std::exception& e ) + { + std::cerr << e.what() << std::endl; + return nullptr; + } +} + +// GPU batch first (sync), then CPU batch on the remainder. +template +hiprtError dispatchHybrid( hiprtContext context, - hiprtScene scene, + HandleT handle, const hiprtHybridTraceConfig& config, - const hiprtHybridTraceGpuInput& gpuInput, + const hiprtHybridTraceGpuInput* gpuInput, + HybridContext::HybridBatchKernel kernelKind, hiprtRay* rays, hiprtHit* hits, uint32_t rayCount, - hiprtApiStream stream ) + hiprtApiStream stream, + CpuFn cpuFn ) { - if ( context == nullptr || scene == nullptr || rays == nullptr || hits == nullptr ) - return hiprtErrorInvalidParameter; - - if ( dynamic_cast( asContext( context ) ) == nullptr ) + if ( context == nullptr || handle == nullptr || rays == nullptr || hits == nullptr ) return hiprtErrorInvalidParameter; - if ( hiprtGetInternalCpuDataFromScene( scene ) == nullptr ) + HybridContext* hybrid = dynamic_cast( asContext( context ) ); + if ( hybrid == nullptr ) return hiprtErrorInvalidParameter; const uint32_t gpuCount = computeGpuRayCount( rayCount, config ); @@ -105,22 +139,201 @@ hiprtError hiprtTraceHybridClosest( if ( gpuCount > 0 ) { - const hiprtError gpuErr = - launchSceneTraceBatchKernel( gpuInput, scene, gpuCount, rays, hits, stream ); - if ( gpuErr != hiprtSuccess ) - return gpuErr; - - const oroStream gpuStream = reinterpret_cast( stream ); - const oroError err = ( gpuStream != nullptr ) ? oroStreamSynchronize( gpuStream ) : oroDeviceSynchronize(); - if ( err != oroSuccess ) + oroFunction func = resolveBatchKernel( hybrid, gpuInput, kernelKind ); + if ( func == nullptr ) return hiprtErrorInternal; + + const hiprtError e = launchBatchKernel( func, handle, gpuCount, rays, hits, stream ); + if ( e != hiprtSuccess ) + return e; + const hiprtError s = syncStream( stream ); + if ( s != hiprtSuccess ) + return s; } if ( cpuCount > 0 ) - { - hiprtGeomTraversalClosestCPU::traceBatch( - context, scene, rays + gpuCount, hits + gpuCount, cpuCount ); - } + cpuFn( rays + gpuCount, hits + gpuCount, cpuCount ); return hiprtSuccess; } +} // namespace + +static hiprtError traceHybridClosestScene( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput* gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + if ( scene != nullptr && hiprtGetInternalCpuDataFromScene( scene ) == nullptr ) + return hiprtErrorInvalidParameter; + + return dispatchHybrid( + context, scene, config, gpuInput, HybridContext::HybridBatchKernel::SceneClosest, rays, hits, rayCount, stream, + [context, scene]( hiprtRay* r, hiprtHit* h, uint32_t c ) { + hiprtGeomTraversalClosestCPU::traceBatch( context, scene, r, h, c ); + } ); +} + +hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridClosestScene( context, scene, config, &gpuInput, rays, hits, rayCount, stream ); +} + +hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridClosestScene( context, scene, config, nullptr, rays, hits, rayCount, stream ); +} + +static hiprtError traceHybridClosestGeom( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput* gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + if ( geometry != nullptr && hiprtGetInternalCpuDataFromGeometry( geometry ) == nullptr ) + return hiprtErrorInvalidParameter; + + return dispatchHybrid( + context, geometry, config, gpuInput, HybridContext::HybridBatchKernel::GeomClosest, rays, hits, rayCount, stream, + [context, geometry]( hiprtRay* r, hiprtHit* h, uint32_t c ) { + hiprtGeomTraversalClosestCPU::traceBatch( context, geometry, r, h, c ); + } ); +} + +hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridClosestGeom( context, geometry, config, &gpuInput, rays, hits, rayCount, stream ); +} + +hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridClosestGeom( context, geometry, config, nullptr, rays, hits, rayCount, stream ); +} + +static hiprtError traceHybridAnyHitScene( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput* gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + if ( scene != nullptr && hiprtGetInternalCpuDataFromScene( scene ) == nullptr ) + return hiprtErrorInvalidParameter; + + return dispatchHybrid( + context, scene, config, gpuInput, HybridContext::HybridBatchKernel::SceneAnyHit, rays, hits, rayCount, stream, + [context, scene]( hiprtRay* r, hiprtHit* h, uint32_t c ) { + hiprtGeomTraversalAnyHitCPU::traceBatch( context, scene, r, h, c ); + } ); +} + +hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridAnyHitScene( context, scene, config, &gpuInput, rays, hits, rayCount, stream ); +} + +hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridAnyHitScene( context, scene, config, nullptr, rays, hits, rayCount, stream ); +} + +static hiprtError traceHybridAnyHitGeom( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput* gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + if ( geometry != nullptr && hiprtGetInternalCpuDataFromGeometry( geometry ) == nullptr ) + return hiprtErrorInvalidParameter; + + return dispatchHybrid( + context, geometry, config, gpuInput, HybridContext::HybridBatchKernel::GeomAnyHit, rays, hits, rayCount, stream, + [context, geometry]( hiprtRay* r, hiprtHit* h, uint32_t c ) { + hiprtGeomTraversalAnyHitCPU::traceBatch( context, geometry, r, h, c ); + } ); +} + +hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridAnyHitGeom( context, geometry, config, &gpuInput, rays, hits, rayCount, stream ); +} + +hiprtError hiprtTraceHybridAnyHit( + hiprtContext context, + hiprtGeometry geometry, + const hiprtHybridTraceConfig& config, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream ) +{ + return traceHybridAnyHitGeom( context, geometry, config, nullptr, rays, hits, rayCount, stream ); +} diff --git a/hiprt/kernels/HybridGeomAnyHitBatchKernel.h b/hiprt/kernels/HybridGeomAnyHitBatchKernel.h new file mode 100644 index 0000000..0b09636 --- /dev/null +++ b/hiprt/kernels/HybridGeomAnyHitBatchKernel.h @@ -0,0 +1,43 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include + +#if defined( __KERNELCC__ ) +#include +#endif + +// Any-hit geometry batch kernel for hiprtTraceHybridAnyHit. +extern "C" __global__ void HybridGeomAnyHitBatchKernel( + hiprtGeometry geom, uint32_t numRays, hiprtRay* rays, hiprtHit* hits ) +{ + const uint32_t index = blockIdx.x * blockDim.x + threadIdx.x; + if ( index >= numRays ) + return; + + hiprtGeomTraversalAnyHit tr( geom, rays[index] ); + hits[index] = tr.getNextHit(); +} diff --git a/hiprt/kernels/HybridGeomTraceBatchKernel.h b/hiprt/kernels/HybridGeomTraceBatchKernel.h new file mode 100644 index 0000000..1a6cb80 --- /dev/null +++ b/hiprt/kernels/HybridGeomTraceBatchKernel.h @@ -0,0 +1,43 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include + +#if defined( __KERNELCC__ ) +#include +#endif + +// Closest-hit geometry batch kernel for hiprtTraceHybridClosest. +extern "C" __global__ void HybridGeomTraceBatchKernel( + hiprtGeometry geom, uint32_t numRays, hiprtRay* rays, hiprtHit* hits ) +{ + const uint32_t index = blockIdx.x * blockDim.x + threadIdx.x; + if ( index >= numRays ) + return; + + hiprtGeomTraversalClosest tr( geom, rays[index] ); + hits[index] = tr.getNextHit(); +} diff --git a/hiprt/kernels/HybridSceneAnyHitBatchKernel.h b/hiprt/kernels/HybridSceneAnyHitBatchKernel.h new file mode 100644 index 0000000..c11061f --- /dev/null +++ b/hiprt/kernels/HybridSceneAnyHitBatchKernel.h @@ -0,0 +1,43 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// Copyright (c) 2024 Advanced Micro Devices, Inc. All Rights Reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. +// +////////////////////////////////////////////////////////////////////////////////////////// + +#pragma once + +#include + +#if defined( __KERNELCC__ ) +#include +#endif + +// Any-hit scene batch kernel for hiprtTraceHybridAnyHit. +extern "C" __global__ void HybridSceneAnyHitBatchKernel( + hiprtScene scene, uint32_t numRays, hiprtRay* rays, hiprtHit* hits ) +{ + const uint32_t index = blockIdx.x * blockDim.x + threadIdx.x; + if ( index >= numRays ) + return; + + hiprtSceneTraversalAnyHit tr( scene, rays[index] ); + hits[index] = tr.getNextHit(); +} diff --git a/hiprt/kernels/HybridSceneTraceBatchKernel.h b/hiprt/kernels/HybridSceneTraceBatchKernel.h index 0e437e6..b7c3fe7 100644 --- a/hiprt/kernels/HybridSceneTraceBatchKernel.h +++ b/hiprt/kernels/HybridSceneTraceBatchKernel.h @@ -30,11 +30,7 @@ #include #endif -/** \brief Closest-hit scene traversal for a contiguous prefix of a ray batch. - * - * Intended for use with hiprtTraceHybridClosest. Uses the built-in traversal - * stack inside hiprtSceneTraversalClosest (no global stack buffer required). - */ +// Closest-hit scene batch kernel for hiprtTraceHybridClosest. extern "C" __global__ void HybridSceneTraceBatchKernel( hiprtScene scene, uint32_t numRays, hiprtRay* rays, hiprtHit* hits ) { diff --git a/test/HybridTraceBatchTest.cpp b/test/HybridTraceBatchTest.cpp index 593e4cb..06c0dfa 100644 --- a/test/HybridTraceBatchTest.cpp +++ b/test/HybridTraceBatchTest.cpp @@ -29,16 +29,14 @@ bool readKernelSource( const std::filesystem::path& srcPath, std::string& source return true; } -hiprtApiFunction buildHybridSceneTraceBatchKernel( hiprtContext context ) +hiprtApiFunction buildHybridKernel( hiprtContext context, const char* relPath, const char* funcName ) { - const auto kernelPath = getRootDir() / "hiprt/kernels/HybridSceneTraceBatchKernel.h"; - std::string sourceCode; + const auto kernelPath = getRootDir() / relPath; + std::string sourceCode; const std::string moduleName = kernelPath.string(); if ( !readKernelSource( kernelPath, sourceCode ) ) return nullptr; - const char* funcName = "HybridSceneTraceBatchKernel"; - std::vector optionStorage; optionStorage.push_back( "-I" + getRootDir().string() ); if ( oroGetCurAPI( 0 ) == ORO_API_HIP ) @@ -61,6 +59,11 @@ hiprtApiFunction buildHybridSceneTraceBatchKernel( hiprtContext context ) return funcApi; } +hiprtApiFunction buildHybridSceneTraceBatchKernel( hiprtContext context ) +{ + return buildHybridKernel( context, "hiprt/kernels/HybridSceneTraceBatchKernel.h", "HybridSceneTraceBatchKernel" ); +} + hiprtFrameSRT makeSrt( float tx, float ty, float tz ) { hiprtFrameSRT f{}; @@ -82,6 +85,10 @@ bool hitsEqual( const hiprtHit& a, const hiprtHit& b ) return std::fabs( a.t - b.t ) < 1e-4f; } +// Any-hit / occlusion only reports presence of an intersection; primID/uv are +// not provided by the CPU occlusion path, so compare hasHit() alone. +bool anyHitEqual( const hiprtHit& a, const hiprtHit& b ) { return a.hasHit() == b.hasHit(); } + void fillTestRays( hiprtRay* rays, uint32_t count ) { for ( uint32_t i = 0; i < count; ++i ) @@ -132,9 +139,60 @@ hiprtScene HybridTraceBatchTest::buildSingleInstanceScene( hiprtGeometry& geomOu TEST_F( HybridTraceBatchTest, SplitBatchMatchesCpuReference ) { + // Uses the simple API: no kernel to compile or pass - the hybrid context + // compiles and caches the batch kernel internally. + constexpr uint32_t rayCount = 16; + + hiprtGeometry geom = nullptr; + hiprtScene scene = buildSingleInstanceScene( geom ); + ASSERT_TRUE( scene != nullptr ); + ASSERT_TRUE( geom != nullptr ); + + hiprtRay* rays = nullptr; + hiprtHit* refHits = nullptr; + hiprtHit* hybridHits = nullptr; + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &rays ), rayCount * sizeof( hiprtRay ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + m_managedAllocs.push_back( rays ); + m_managedAllocs.push_back( refHits ); + m_managedAllocs.push_back( hybridHits ); + + fillTestRays( rays, rayCount ); + std::memset( refHits, 0, rayCount * sizeof( hiprtHit ) ); + std::memset( hybridHits, 0, rayCount * sizeof( hiprtHit ) ); + + hiprtGeomTraversalClosestCPU::traceBatch( m_context, scene, rays, refHits, rayCount ); + + hiprtHybridTraceConfig cfg{}; + cfg.gpuFraction = 0.5f; + cfg.minCpuBatch = 1u; + + ASSERT_EQ( + hiprtTraceHybridClosest( m_context, scene, cfg, rays, hybridHits, rayCount, nullptr ), hiprtSuccess ); + + for ( uint32_t i = 0; i < rayCount; ++i ) + { + EXPECT_TRUE( hitsEqual( hybridHits[i], refHits[i] ) ) + << "Ray " << i << " mismatch (even=hit, odd=miss)."; + } + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridTraceBatchTest, ExplicitGpuInputStillWorks ) +{ + // Advanced path: caller compiles the kernel and passes it via gpuInput. constexpr uint32_t rayCount = 16; - hiprtGeometry geom = nullptr; + hiprtGeometry geom = nullptr; hiprtScene scene = buildSingleInstanceScene( geom ); ASSERT_TRUE( scene != nullptr ); ASSERT_TRUE( geom != nullptr ); @@ -142,19 +200,17 @@ TEST_F( HybridTraceBatchTest, SplitBatchMatchesCpuReference ) hiprtApiFunction traceKernel = buildHybridSceneTraceBatchKernel( m_context ); ASSERT_TRUE( traceKernel != nullptr ) << "Failed to compile HybridSceneTraceBatchKernel."; - hiprtRay* rays = nullptr; - hiprtHit* refHits = nullptr; + hiprtRay* rays = nullptr; + hiprtHit* refHits = nullptr; hiprtHit* hybridHits = nullptr; ASSERT_EQ( oroMallocManaged( reinterpret_cast( &rays ), rayCount * sizeof( hiprtRay ), oroMemAttachGlobal ), oroSuccess ); ASSERT_EQ( - oroMallocManaged( - reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroMallocManaged( reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), oroSuccess ); ASSERT_EQ( - oroMallocManaged( - reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroMallocManaged( reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), oroSuccess ); m_managedAllocs.push_back( rays ); m_managedAllocs.push_back( refHits ); @@ -174,8 +230,7 @@ TEST_F( HybridTraceBatchTest, SplitBatchMatchesCpuReference ) gpuInput.traceKernel = traceKernel; ASSERT_EQ( - hiprtTraceHybridClosest( m_context, scene, cfg, gpuInput, rays, hybridHits, rayCount, nullptr ), - hiprtSuccess ); + hiprtTraceHybridClosest( m_context, scene, cfg, gpuInput, rays, hybridHits, rayCount, nullptr ), hiprtSuccess ); for ( uint32_t i = 0; i < rayCount; ++i ) { @@ -187,6 +242,100 @@ TEST_F( HybridTraceBatchTest, SplitBatchMatchesCpuReference ) ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); } +TEST_F( HybridTraceBatchTest, AnyHitSceneMatchesCpuReference ) +{ + constexpr uint32_t rayCount = 16; + + hiprtGeometry geom = nullptr; + hiprtScene scene = buildSingleInstanceScene( geom ); + ASSERT_TRUE( scene != nullptr ); + ASSERT_TRUE( geom != nullptr ); + + hiprtRay* rays = nullptr; + hiprtHit* refHits = nullptr; + hiprtHit* hybridHits = nullptr; + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &rays ), rayCount * sizeof( hiprtRay ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + m_managedAllocs.push_back( rays ); + m_managedAllocs.push_back( refHits ); + m_managedAllocs.push_back( hybridHits ); + + fillTestRays( rays, rayCount ); + std::memset( refHits, 0, rayCount * sizeof( hiprtHit ) ); + std::memset( hybridHits, 0, rayCount * sizeof( hiprtHit ) ); + + hiprtGeomTraversalAnyHitCPU::traceBatch( m_context, scene, rays, refHits, rayCount ); + + hiprtHybridTraceConfig cfg{}; + cfg.gpuFraction = 0.5f; + cfg.minCpuBatch = 1u; + + ASSERT_EQ( + hiprtTraceHybridAnyHit( m_context, scene, cfg, rays, hybridHits, rayCount, nullptr ), hiprtSuccess ); + + for ( uint32_t i = 0; i < rayCount; ++i ) + { + EXPECT_TRUE( anyHitEqual( hybridHits[i], refHits[i] ) ) << "Ray " << i << " any-hit mismatch."; + } + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + +TEST_F( HybridTraceBatchTest, GeometryClosestMatchesCpuReference ) +{ + constexpr uint32_t rayCount = 16; + + hiprtGeometry geom = nullptr; + hiprtScene scene = buildSingleInstanceScene( geom ); + ASSERT_TRUE( scene != nullptr ); + ASSERT_TRUE( geom != nullptr ); + + hiprtRay* rays = nullptr; + hiprtHit* refHits = nullptr; + hiprtHit* hybridHits = nullptr; + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &rays ), rayCount * sizeof( hiprtRay ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &refHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + ASSERT_EQ( + oroMallocManaged( reinterpret_cast( &hybridHits ), rayCount * sizeof( hiprtHit ), oroMemAttachGlobal ), + oroSuccess ); + m_managedAllocs.push_back( rays ); + m_managedAllocs.push_back( refHits ); + m_managedAllocs.push_back( hybridHits ); + + fillTestRays( rays, rayCount ); + std::memset( refHits, 0, rayCount * sizeof( hiprtHit ) ); + std::memset( hybridHits, 0, rayCount * sizeof( hiprtHit ) ); + + hiprtGeomTraversalClosestCPU::traceBatch( m_context, geom, rays, refHits, rayCount ); + + hiprtHybridTraceConfig cfg{}; + cfg.gpuFraction = 0.5f; + cfg.minCpuBatch = 1u; + + ASSERT_EQ( + hiprtTraceHybridClosest( m_context, geom, cfg, rays, hybridHits, rayCount, nullptr ), hiprtSuccess ); + + for ( uint32_t i = 0; i < rayCount; ++i ) + { + EXPECT_TRUE( hitsEqual( hybridHits[i], refHits[i] ) ) << "Ray " << i << " geometry-closest mismatch."; + } + + ASSERT_EQ( hiprtDestroyScene( m_context, scene ), hiprtSuccess ); + ASSERT_EQ( hiprtDestroyGeometry( m_context, geom ), hiprtSuccess ); +} + TEST_F( HybridTraceBatchTest, RejectsNonHybridContext ) { hiprtContextCreationInput cpuInput{}; @@ -202,7 +351,7 @@ TEST_F( HybridTraceBatchTest, RejectsNonHybridContext ) hiprtHybridTraceGpuInput gpuInput{ nullptr, {} }; EXPECT_EQ( - hiprtTraceHybridClosest( cpuContext, nullptr, cfg, gpuInput, &ray, &hit, 1, nullptr ), + hiprtTraceHybridClosest( cpuContext, static_cast( nullptr ), cfg, gpuInput, &ray, &hit, 1, nullptr ), hiprtErrorInvalidParameter ); EXPECT_EQ( hiprtDestroyContext( cpuContext ), hiprtSuccess );