From 249c7368ca2c018d9d5a54e2ebffc4b63a194d16 Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Tue, 21 Oct 2025 16:08:00 +0200 Subject: [PATCH 1/3] Small fixes --- gpu/include/raygTrace.hpp | 6 ++++-- gpu/include/raygTraceDisk.hpp | 2 +- 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/gpu/include/raygTrace.hpp b/gpu/include/raygTrace.hpp index fc062f6..5b66a8b 100644 --- a/gpu/include/raygTrace.hpp +++ b/gpu/include/raygTrace.hpp @@ -196,8 +196,8 @@ template class Trace { generateSBT(); #ifndef NDEBUG // Launch on single stream in debug mode - for (size_t i = 0; i < particles.size(); i++) { - OPTIX_CHECK(optixLaunch(pipeline, streams[0], + for (size_t i = 0; i < particles_.size(); i++) { + OPTIX_CHECK(optixLaunch(pipeline_, streams[0], /*! parameters and SBT */ launchParamsBuffers[i].dPointer(), launchParamsBuffers[i].sizeInBytes, &sbt, @@ -300,6 +300,8 @@ template class Trace { callableMap_ = std::get<1>(maps); } + size_t getNumberOfRays() const { return numRays; } + void getFlux(float *flux, int particleIdx, int dataIdx, int smoothingNeighbors = 0) { unsigned int offset = 0; diff --git a/gpu/include/raygTraceDisk.hpp b/gpu/include/raygTraceDisk.hpp index 7cd3bdd..4f2213f 100644 --- a/gpu/include/raygTraceDisk.hpp +++ b/gpu/include/raygTraceDisk.hpp @@ -20,7 +20,7 @@ template class TraceDisk : public Trace { ~TraceDisk() { diskGeometry.freeBuffers(); } void setGeometry(const DiskMesh &passedMesh) { - assert(context); + assert(context_ && "Context not initialized."); diskMesh = passedMesh; if (diskMesh.gridDelta <= 0.f) { Logger::getInstance() From 3b3ee6f59fbaf3701f71a7a83d5cb2341f79671d Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Tue, 21 Oct 2025 19:16:47 +0200 Subject: [PATCH 2/3] Small fixes --- gpu/include/raygDiskGeometry.hpp | 1 - gpu/include/raygLaunchParams.hpp | 1 - gpu/include/raygLineGeometry.hpp | 8 +++----- gpu/include/raygSource.hpp | 4 ---- gpu/include/raygTraceTriangle.hpp | 6 ++++++ gpu/include/raygTriangleGeometry.hpp | 1 - gpu/pipelines/GeneralPipelineDisk.cu | 5 ++--- gpu/pipelines/GeneralPipelineLine.cu | 5 ++--- gpu/pipelines/GeneralPipelineTriangle.cu | 1 - 9 files changed, 13 insertions(+), 19 deletions(-) diff --git a/gpu/include/raygDiskGeometry.hpp b/gpu/include/raygDiskGeometry.hpp index edff9e5..5c66bd2 100644 --- a/gpu/include/raygDiskGeometry.hpp +++ b/gpu/include/raygDiskGeometry.hpp @@ -30,7 +30,6 @@ template struct DiskGeometry { LaunchParams &launchParams) { assert(context.deviceID != -1 && "Context not initialized."); - launchParams.source.gridDelta = mesh.gridDelta; if constexpr (D == 2) { launchParams.source.minPoint[0] = mesh.minimumExtent[0]; launchParams.source.maxPoint[0] = mesh.maximumExtent[0]; diff --git a/gpu/include/raygLaunchParams.hpp b/gpu/include/raygLaunchParams.hpp index 6803e64..0d454b4 100644 --- a/gpu/include/raygLaunchParams.hpp +++ b/gpu/include/raygLaunchParams.hpp @@ -39,7 +39,6 @@ struct LaunchParams { struct { viennacore::Vec2Df minPoint; viennacore::Vec2Df maxPoint; - float gridDelta; float planeHeight; std::array directionBasis; } source; diff --git a/gpu/include/raygLineGeometry.hpp b/gpu/include/raygLineGeometry.hpp index f51b766..19b28e3 100644 --- a/gpu/include/raygLineGeometry.hpp +++ b/gpu/include/raygLineGeometry.hpp @@ -27,22 +27,20 @@ template struct LineGeometry { void buildAccel(DeviceContext &context, const LineMesh &mesh, LaunchParams &launchParams) { assert(context.deviceID != -1 && "Context not initialized."); + assert(mesh.gridDelta > 0.f && "Grid delta must be positive."); - launchParams.source.gridDelta = mesh.gridDelta; if constexpr (D == 2) { launchParams.source.minPoint[0] = mesh.minimumExtent[0]; launchParams.source.maxPoint[0] = mesh.maximumExtent[0]; launchParams.source.planeHeight = - mesh.maximumExtent[1] + - 2 * mesh.gridDelta * rayInternal::DiskFactor; + mesh.maximumExtent[1] + 2 * mesh.gridDelta; } else { launchParams.source.minPoint[0] = mesh.minimumExtent[0]; launchParams.source.minPoint[1] = mesh.minimumExtent[1]; launchParams.source.maxPoint[0] = mesh.maximumExtent[0]; launchParams.source.maxPoint[1] = mesh.maximumExtent[1]; launchParams.source.planeHeight = - mesh.maximumExtent[2] + - 2 * mesh.gridDelta * rayInternal::DiskFactor; + mesh.maximumExtent[2] + 2 * mesh.gridDelta; } launchParams.numElements = mesh.lines.size(); diff --git a/gpu/include/raygSource.hpp b/gpu/include/raygSource.hpp index 00a4eac..70c4dc4 100644 --- a/gpu/include/raygSource.hpp +++ b/gpu/include/raygSource.hpp @@ -76,10 +76,6 @@ initializeRayPosition(viennaray::gpu::PerRayData *prd, prd->pos[0] = launchParams->source.minPoint[0] + u.x * (launchParams->source.maxPoint[0] - launchParams->source.minPoint[0]); - // prd->pos[1] = launchParams->source.minPoint[1] + - // u.y * (launchParams->source.maxPoint[1] - - // launchParams->source.minPoint[1]); - // prd->pos[2] = launchParams->source.planeHeight; if (D == 2) { prd->pos[1] = launchParams->source.planeHeight; diff --git a/gpu/include/raygTraceTriangle.hpp b/gpu/include/raygTraceTriangle.hpp index 425be43..80b40a7 100644 --- a/gpu/include/raygTraceTriangle.hpp +++ b/gpu/include/raygTraceTriangle.hpp @@ -50,6 +50,7 @@ template class TraceTriangle : public Trace { void buildHitGroups() override { // geometry hitgroup std::vector hitgroupRecords; + HitgroupRecordTriangle geometryHitgroupRecord = {}; optixSbtRecordPackHeader(hitgroupPG, &geometryHitgroupRecord); geometryHitgroupRecord.data.vertex = @@ -60,6 +61,8 @@ template class TraceTriangle : public Trace { geometryHitgroupRecord.data.base.isBoundary = false; geometryHitgroupRecord.data.base.cellData = (void *)this->cellDataBuffer_.dPointer(); + + // add geometry hitgroup record hitgroupRecords.push_back(geometryHitgroupRecord); // boundary hitgroup @@ -71,8 +74,11 @@ template class TraceTriangle : public Trace { (Vec3D *)triangleGeometry.boundaryIndexBuffer.dPointer(); boundaryHitgroupRecord.data.base.geometryType = 0; boundaryHitgroupRecord.data.base.isBoundary = true; + + // add boundary hitgroup record hitgroupRecords.push_back(boundaryHitgroupRecord); + // upload hitgroup records hitgroupRecordBuffer.allocUpload(hitgroupRecords); sbt.hitgroupRecordBase = hitgroupRecordBuffer.dPointer(); sbt.hitgroupRecordStrideInBytes = sizeof(HitgroupRecordTriangle); diff --git a/gpu/include/raygTriangleGeometry.hpp b/gpu/include/raygTriangleGeometry.hpp index 2a061f5..7d9a15b 100644 --- a/gpu/include/raygTriangleGeometry.hpp +++ b/gpu/include/raygTriangleGeometry.hpp @@ -27,7 +27,6 @@ struct TriangleGeometry { LaunchParams &launchParams) { assert(context.deviceID != -1 && "Context not initialized."); - launchParams.source.gridDelta = mesh.gridDelta; launchParams.source.minPoint[0] = mesh.minimumExtent[0]; launchParams.source.minPoint[1] = mesh.minimumExtent[1]; launchParams.source.maxPoint[0] = mesh.maximumExtent[0]; diff --git a/gpu/pipelines/GeneralPipelineDisk.cu b/gpu/pipelines/GeneralPipelineDisk.cu index 910de43..09e5ed3 100644 --- a/gpu/pipelines/GeneralPipelineDisk.cu +++ b/gpu/pipelines/GeneralPipelineDisk.cu @@ -134,9 +134,8 @@ extern "C" __global__ void __closesthit__() { // } // ------------- SURFACE COLLISION --------------- // - unsigned callIdx; - - callIdx = callableIndex(launchParams.particleType, CallableSlot::COLLISION); + unsigned callIdx = + callableIndex(launchParams.particleType, CallableSlot::COLLISION); optixDirectCall(callIdx, sbtData, prd); diff --git a/gpu/pipelines/GeneralPipelineLine.cu b/gpu/pipelines/GeneralPipelineLine.cu index 0eec7f0..ce9995b 100644 --- a/gpu/pipelines/GeneralPipelineLine.cu +++ b/gpu/pipelines/GeneralPipelineLine.cu @@ -116,9 +116,8 @@ extern "C" __global__ void __closesthit__() { // } // ------------- SURFACE COLLISION --------------- // - unsigned callIdx; - - callIdx = callableIndex(launchParams.particleType, CallableSlot::COLLISION); + unsigned callIdx = + callableIndex(launchParams.particleType, CallableSlot::COLLISION); optixDirectCall(callIdx, sbtData, prd); diff --git a/gpu/pipelines/GeneralPipelineTriangle.cu b/gpu/pipelines/GeneralPipelineTriangle.cu index e66710e..073053a 100644 --- a/gpu/pipelines/GeneralPipelineTriangle.cu +++ b/gpu/pipelines/GeneralPipelineTriangle.cu @@ -54,7 +54,6 @@ extern "C" __global__ void __closesthit__() { PerRayData *>(callIdx, sbtData, prd); // ------------- REFLECTION --------------- // - callIdx = callableIndex(launchParams.particleType, CallableSlot::REFLECTION); optixDirectCall( From 5226fcb44bd24b0f5a0439405d474602951f14af Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Tue, 21 Oct 2025 19:20:07 +0200 Subject: [PATCH 3/3] Bump ViennaCore --- CMakeLists.txt | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index f233f8d..c458dec 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -92,8 +92,7 @@ include("cmake/cpm.cmake") CPMAddPackage( NAME ViennaCore - VERSION 1.6.1 - GIT_TAG main + VERSION 1.6.2 GIT_REPOSITORY "https://github.com/ViennaTools/ViennaCore" OPTIONS "VIENNACORE_USE_GPU ${VIENNARAY_USE_GPU}")