diff --git a/CMakeLists.txt b/CMakeLists.txt index 625ea77..651e607 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,18 @@ 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/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" + "${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}) @@ -720,3 +742,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/hiprt.h.in b/hiprt/hiprt.h.in index b83ae53..4dc1d46 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 (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. + * + * \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 (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. + * + * \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 new file mode 100644 index 0000000..dbbb93c --- /dev/null +++ b/hiprt/hiprt_cpu.h @@ -0,0 +1,196 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// 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 +#include +#include +#include +#include + +// Internal helpers +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; + + 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; + return hit; +} + +inline void traceGeometryBatch( + 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 ) + { + RTCRayHit rh = makeRTCRayHit( rays[i] ); + rtcIntersect1( scene, &rh ); + hits[i] = makeHiprtHit( rh ); + } +} + +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 + +// Host-side CPU batch traversal (Embree). +class hiprtGeomTraversalClosestCPU +{ + 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::traceGeometryBatch( + ( 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::traceGeometryBatch( + ( 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 new file mode 100644 index 0000000..5bdecb5 --- /dev/null +++ b/hiprt/hiprt_hybrid.h @@ -0,0 +1,119 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// 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 + +struct hiprtHybridTraceConfig +{ + float gpuFraction = 0.8f; + uint32_t minCpuBatch = 1024u; +}; + +// Optional: pass your own GPU kernel. Default overloads compile one internally. +struct hiprtHybridTraceGpuInput +{ + hiprtApiFunction traceKernel = nullptr; + hiprtGlobalStackBuffer globalStack = {}; +}; + +// Hybrid closest-hit (scene). Requires hybrid context; rays/hits must be host-visible. +HIPRT_API hiprtError hiprtTraceHybridClosest( + hiprtContext context, + hiprtScene scene, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput& gpuInput, + hiprtRay* rays, + 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/hiprt_types.h b/hiprt/hiprt_types.h index 2e631c3..555b3d1 100644 --- a/hiprt/hiprt_types.h +++ b/hiprt/hiprt_types.h @@ -309,16 +309,21 @@ 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. * + * 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 { @@ -328,6 +333,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..841aa15 --- /dev/null +++ b/hiprt/impl/ContextBase.h @@ -0,0 +1,102 @@ + +#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 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; + 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..d2f71cd --- /dev/null +++ b/hiprt/impl/CpuContext.cpp @@ -0,0 +1,483 @@ +#include "CpuContext.h" +#include + +#include + +#include +#include + +namespace hiprt +{ +namespace +{ +[[noreturn]] void throwNotImplemented( const char* what ) +{ + 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 ) +{ + 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& 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" ); + } + + hiprtGeometry handle = reinterpret_cast( data ); + registerCpuGeom( handle, data ); + geometries.push_back( handle ); + } + + return geometries; +} + +void CpuContext::destroyGeometries( const std::vector& geometries ) +{ + for ( hiprtGeometry geometry : geometries ) + { + CpuGeometryData* data = reinterpret_cast( geometry ); + if ( data == nullptr ) continue; + + unregisterCpuGeom( geometry ); + + if ( data->rtcScene != nullptr ) + { + rtcReleaseScene( data->rtcScene ); + data->rtcScene = nullptr; + } + + delete data; + } +} + +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& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + (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& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + (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 ) +{ + 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& buildInputs, const hiprtBuildOptions ) +{ + 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" ); + } + + hiprtScene handle = reinterpret_cast( data ); + registerCpuScene( handle, data ); + scenes.push_back( handle ); + } + return scenes; +} + +void CpuContext::destroyScenes( const std::vector& scenes ) +{ + for ( hiprtScene scene : scenes ) + { + CpuSceneData* data = getCpuScene( scene ); + if ( data == nullptr ) continue; + + unregisterCpuScene( scene ); + + 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& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + (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& buildInputs, + const hiprtBuildOptions buildOptions, + hiprtDevicePtr temporaryBuffer, + oroStream stream, + std::vector& buffers ) +{ + (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 ) +{ + 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..9a86e34 --- /dev/null +++ b/hiprt/impl/CpuContext.h @@ -0,0 +1,129 @@ +#pragma once +#include "ContextBase.h" +#include "CpuTypes.h" + +#include +#include +#include +#include +#include +#include + +#include + + +namespace hiprt{ +class CpuContext : public ContextBase +{ + public: + CpuContext( const hiprtContextCreationInput& input ); + ~CpuContext() override; + + 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 ) 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; } + + 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 geometry ) override + { + return reinterpret_cast( geometry ); + }; + struct CpuSceneData* getCpuScene( hiprtScene scene ) override + { + 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; +}; +}//namespace hiprt \ No newline at end of file diff --git a/hiprt/impl/CpuDataRegistry.h b/hiprt/impl/CpuDataRegistry.h new file mode 100644 index 0000000..540830b --- /dev/null +++ b/hiprt/impl/CpuDataRegistry.h @@ -0,0 +1,20 @@ +#pragma once + +#include + +namespace hiprt +{ +struct CpuGeometryData; +struct CpuSceneData; + +// Maps GPU and CPU handles for hybrid traversal. Thread-safe. + +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/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/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 75% rename from hiprt/impl/Context.h rename to hiprt/impl/GpuContext.h index c078dbd..7549193 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 ); + size_t getScenesBuildTempBufferSize( + const std::vector& buildInputs, const hiprtBuildOptions buildOptions ) override; - 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 ); + 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 ); - 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 setCacheDir( const std::filesystem::path& path ) override; - 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/HybridContext.cpp b/hiprt/impl/HybridContext.cpp new file mode 100644 index 0000000..1cc6324 --- /dev/null +++ b/hiprt/impl/HybridContext.cpp @@ -0,0 +1,448 @@ +#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" ); +} + +[[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 + +HybridContext::HybridContext( const hiprtContextCreationInput& input ) + : m_gpu( input ), m_cpu( input ) +{ + std::cerr << "[HybridContext] constructed - GPU BVH + Embree CPU mirror active.\n"; +} + +HybridContext::~HybridContext() {} + +oroFunction HybridContext::getBatchKernel( HybridBatchKernel kind ) +{ + 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 ) +{ + 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; +} + +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 ) +{ + // Map GPU buffers to CPU buffers. + 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 ) ); + } + + // 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 + 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 ); + syncGpu( stream ); + 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 ); + syncGpu( stream ); + 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*/ ) +{ + throwHybridNotImplemented( "compactGeometries is not supported in hybrid mode (CPU mirror rebuild required)" ); +} + +// 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 ) +{ + // Remap GPU instance handles to CPU handles. + std::vector cpuBuffers; + std::vector cpuInputs; + std::vector> remappedInstances( buildInputs.size() ); + + cpuBuffers.reserve( gpuBuffers.size() ); + cpuInputs.reserve( buildInputs.size() ); + + for ( size_t i = 0; i < buildInputs.size(); ++i ) + { + 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 ) ); + + 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 ); + 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 ); + syncGpu( stream ); + 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 ); + syncGpu( stream ); + 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*/ ) +{ + throwHybridNotImplemented( "compactScenes is not supported in hybrid mode (CPU mirror rebuild required)" ); +} + +// GPU 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*/ ) +{ + throwHybridNotImplemented( "loadGeometry is not supported in hybrid mode (CPU mirror rebuild required)" ); +} + +void HybridContext::saveScene( hiprtScene inScene, const std::string& filename ) +{ + m_gpu.saveScene( inScene, filename ); +} + +hiprtScene HybridContext::loadScene( const std::string& /*filename*/ ) +{ + throwHybridNotImplemented( "loadScene is not supported in hybrid mode (CPU mirror rebuild required)" ); +} + +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..2148b0f --- /dev/null +++ b/hiprt/impl/HybridContext.h @@ -0,0 +1,150 @@ +#pragma once + +#include +#include +#include +#include + +#include +#include + +namespace hiprt +{ + +// 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; + + oroFunction getBatchKernel( HybridBatchKernel kind ); + + // 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 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; + + 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 + CpuGeometryData* getCpuGeom( hiprtGeometry geometry ) override; + CpuSceneData* getCpuScene( hiprtScene scene ) override; + + private: + void dispatchCpuGeomBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ); + + void dispatchCpuSceneBuild( + const std::vector& buildInputs, + const std::vector& gpuBuffers, + bool isUpdate ); + + GpuContext m_gpu; + CpuContext m_cpu; + + 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/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..cfa70df 100644 --- a/hiprt/impl/hiprt.cpp +++ b/hiprt/impl/hiprt.cpp @@ -24,22 +24,119 @@ #include #include +#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 +// 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 { - Context* ctxt = new Context( 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 if ( input.deviceType & hiprtDeviceCPU ) + { + // 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. 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 +150,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 +185,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 +216,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 +266,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 +279,7 @@ hiprtError hiprtBuildGeometries( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -211,11 +308,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 +347,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 +385,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 +413,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 +462,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 +475,7 @@ hiprtError hiprtBuildScenes( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -407,11 +504,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 +537,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 +555,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 +571,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 +586,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 +602,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 +617,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 +632,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 +647,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 +662,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 +677,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 +692,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 +707,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 +767,7 @@ hiprtError hiprtBuildTraceKernels( std::vector functions; oroModule module = nullptr; - reinterpret_cast( context )->buildKernels( + asContext( context )->buildKernels( funcNames, src, moduleName, @@ -691,7 +788,7 @@ hiprtError hiprtBuildTraceKernels( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -730,7 +827,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 +835,7 @@ hiprtError hiprtBuildTraceKernelsFromBitcode( } catch ( std::exception& e ) { - reinterpret_cast( context )->logError( e.what() ); + std::cerr << e.what() << std::endl; return hiprtErrorInternal; } @@ -750,11 +847,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,12 +862,22 @@ 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; } + +void* hiprtGetInternalCpuDataFromGeometry( hiprtGeometry geometry ) +{ + return static_cast( hiprt::lookupCpuGeom( geometry ) ); +} + +void* hiprtGetInternalCpuDataFromScene( hiprtScene scene ) +{ + return static_cast( hiprt::lookupCpuScene( scene ) ); +} diff --git a/hiprt/impl/hiprt_hybrid.cpp b/hiprt/impl/hiprt_hybrid.cpp new file mode 100644 index 0000000..78ab893 --- /dev/null +++ b/hiprt/impl/hiprt_hybrid.cpp @@ -0,0 +1,339 @@ +////////////////////////////////////////////////////////////////////////////////////////// +// +// 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 + +#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; +} + +// Batch kernel: (handle, numRays, rays, hits). +hiprtError launchBatchKernel( + oroFunction func, + void* handle, + uint32_t numRays, + hiprtRay* rays, + hiprtHit* hits, + hiprtApiStream stream ) +{ + if ( func == nullptr ) + return hiprtErrorInvalidParameter; + if ( numRays == 0 ) + return hiprtSuccess; + + constexpr uint32_t blockSize = 256u; + const uint32_t gridSize = ( numRays + blockSize - 1 ) / blockSize; + + 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; +} + +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, + HandleT handle, + const hiprtHybridTraceConfig& config, + const hiprtHybridTraceGpuInput* gpuInput, + HybridContext::HybridBatchKernel kernelKind, + hiprtRay* rays, + hiprtHit* hits, + uint32_t rayCount, + hiprtApiStream stream, + CpuFn cpuFn ) +{ + if ( context == nullptr || handle == nullptr || rays == nullptr || hits == nullptr ) + return hiprtErrorInvalidParameter; + + HybridContext* hybrid = dynamic_cast( asContext( context ) ); + if ( hybrid == nullptr ) + return hiprtErrorInvalidParameter; + + const uint32_t gpuCount = computeGpuRayCount( rayCount, config ); + const uint32_t cpuCount = rayCount - gpuCount; + + if ( gpuCount > 0 ) + { + 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 ) + 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 new file mode 100644 index 0000000..b7c3fe7 --- /dev/null +++ b/hiprt/kernels/HybridSceneTraceBatchKernel.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 scene batch kernel for hiprtTraceHybridClosest. +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 new file mode 100644 index 0000000..ff819fb --- /dev/null +++ b/test/CpuGeometryTest.cpp @@ -0,0 +1,149 @@ +#include "CpuGeometryTest.h" + +#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( + 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( 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( 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( 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 ); +} + +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( 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 }; + 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( 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/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; +}; diff --git a/test/CpuSceneTest.cpp b/test/CpuSceneTest.cpp new file mode 100644 index 0000000..ef2927d --- /dev/null +++ b/test/CpuSceneTest.cpp @@ -0,0 +1,224 @@ +#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( + 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( 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( m_context, scene, 0.25f, 0.25f, -1.0f, 0.0f, 0.0f, 1.0f ) ) + << "Promien w obszar instancji powinien trafic."; + 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 ); + 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( 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( 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 ); + 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( 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 ); + in.instanceFrames = frames.data(); + ASSERT_EQ( + hiprtBuildScene( m_context, hiprtBuildOperationUpdate, in, opts, nullptr, nullptr, scene ), + hiprtSuccess ) + << "hiprtBuildScene (Update) failed."; + + 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( 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 ); + 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; +}; 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..06c0dfa --- /dev/null +++ b/test/HybridTraceBatchTest.cpp @@ -0,0 +1,358 @@ +#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 buildHybridKernel( hiprtContext context, const char* relPath, const char* funcName ) +{ + const auto kernelPath = getRootDir() / relPath; + std::string sourceCode; + const std::string moduleName = kernelPath.string(); + if ( !readKernelSource( kernelPath, sourceCode ) ) + return nullptr; + + 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; +} + +hiprtApiFunction buildHybridSceneTraceBatchKernel( hiprtContext context ) +{ + return buildHybridKernel( context, "hiprt/kernels/HybridSceneTraceBatchKernel.h", "HybridSceneTraceBatchKernel" ); +} + +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; +} + +// 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 ) + { + 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 ) +{ + // 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; + 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, 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{}; + 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, static_cast( 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; +}