From 7d1acee59fc20e9a20644fda91f580ac03e9d71a Mon Sep 17 00:00:00 2001 From: reiter Date: Thu, 27 Nov 2025 17:01:48 +0100 Subject: [PATCH 1/4] Add test for GPU triangle boundaries --- cmake/generate_ptx.cmake | 8 +- gpu/include/raygBoundary.hpp | 2 +- gpu/include/raygTrace.hpp | 40 ++++--- gpu/include/raygTriangleGeometry.hpp | 5 + gpu/tests/boundaries/CMakeLists.txt | 8 ++ gpu/tests/boundaries/TestPipelineTriangle.cu | 109 +++++++++++++++++ gpu/tests/boundaries/boundaries.cpp | 118 +++++++++++++++++++ include/viennaray/rayUtil.hpp | 7 ++ 8 files changed, 278 insertions(+), 19 deletions(-) create mode 100644 gpu/tests/boundaries/CMakeLists.txt create mode 100644 gpu/tests/boundaries/TestPipelineTriangle.cu create mode 100644 gpu/tests/boundaries/boundaries.cpp diff --git a/cmake/generate_ptx.cmake b/cmake/generate_ptx.cmake index e57af19..0f05f62 100644 --- a/cmake/generate_ptx.cmake +++ b/cmake/generate_ptx.cmake @@ -5,11 +5,11 @@ function(generate_pipeline target_name generated_files) cuda_get_sources_and_options(cu_optix_source_files cmake_options options ${ARGN}) # Add the path to the OptiX headers to our include paths. - include_directories(${OptiX_INCLUDE_DIR}) + cuda_include_directories(${OptiX_INCLUDE_DIR}) # Include ViennaRay headers which are used in pipelines - include_directories(${VIENNARAY_GPU_INCLUDE_DIR}) - include_directories(${ViennaCore_SOURCE_DIR}/include/viennacore) # needed for Context + cuda_include_directories(${VIENNARAY_GPU_INCLUDE}) + cuda_include_directories(${ViennaCore_SOURCE_DIR}/include/viennacore) add_compile_definitions(VIENNACORE_COMPILE_GPU) # Generate OptiX IR files if enabled @@ -38,7 +38,7 @@ function(generate_pipeline target_name generated_files) list(APPEND generated_files_local ${generated_ptx_files}) endif() - list(APPEND ${generated_files} ${generated_files_local}) + set(${generated_files} ${generated_files_local} PARENT_SCOPE) endfunction() function(generate_kernel generated_files) diff --git a/gpu/include/raygBoundary.hpp b/gpu/include/raygBoundary.hpp index ef477e2..8b4719c 100644 --- a/gpu/include/raygBoundary.hpp +++ b/gpu/include/raygBoundary.hpp @@ -35,7 +35,7 @@ reflectFromBoundary(PerRayData *prd, const SBTData *hsd, const int D) { prd->dir[dim] -= 2 * prd->dir[dim]; prd->pos[dim] = hsd->vertex[hsd->index[primID][0]][dim]; } else if constexpr (std::is_same::value) { - prd->pos = prd->pos + prd->dir * (optixGetRayTmax()); + prd->pos = prd->pos + prd->dir * optixGetRayTmax(); if (primID == 0 || primID == 1) // x boundary prd->dir[0] -= 2 * prd->dir[0]; } diff --git a/gpu/include/raygTrace.hpp b/gpu/include/raygTrace.hpp index e6078a4..569a55e 100644 --- a/gpu/include/raygTrace.hpp +++ b/gpu/include/raygTrace.hpp @@ -53,6 +53,10 @@ template class Trace { destroyMembers(); } + void setPipelineFileName(const std::string &fileName) { + pipelineFileName = fileName; + } + void setCallables(std::string fileName, const std::filesystem::path &path) { // check if filename ends in .optixir if (fileName.find(".optixir") == std::string::npos) { @@ -494,7 +498,7 @@ template class Trace { char log[2048]; size_t sizeof_log = sizeof(log); - std::string pipelineFile = "GeneralPipeline" + geometryType_ + ".optixir"; + std::string pipelineFile = pipelineFileName + geometryType_ + ".optixir"; std::filesystem::path pipelinePath = context_->modulePath / pipelineFile; if (!std::filesystem::exists(pipelinePath)) { Logger::getInstance() @@ -518,6 +522,10 @@ template class Trace { char logCallable[2048]; size_t sizeof_log_callable = sizeof(logCallable); + if (callableFile_.empty()) { + Logger::getInstance().addWarning("No callable file set.").print(); + return; + } auto callableInput = getInputData(callableFile_.c_str(), inputSize); if (!callableInput) { Logger::getInstance() @@ -597,8 +605,9 @@ template class Trace { void createDirectCallablePrograms() { if (callableMap_.empty()) { Logger::getInstance() - .addError("No particleType->callable mapping provided.") + .addWarning("No particleType->callable mapping provided.") .print(); + return; } unsigned maxParticleTypeId = 0; for (const auto &p : particleMap_) { @@ -701,18 +710,20 @@ template class Trace { buildHitGroups(); // callable programs - std::vector callableRecords(directCallablePGs.size()); - for (size_t j = 0; j < directCallablePGs.size(); ++j) { - CallableRecord callableRecord = {}; - optixSbtRecordPackHeader(directCallablePGs[j], &callableRecord); - callableRecords[j] = callableRecord; - } - directCallableRecordBuffer.allocUpload(callableRecords); - - sbt.callablesRecordBase = directCallableRecordBuffer.dPointer(); - sbt.callablesRecordStrideInBytes = sizeof(CallableRecord); - sbt.callablesRecordCount = - static_cast(directCallablePGs.size()); + if (!directCallablePGs.empty()) { + std::vector callableRecords(directCallablePGs.size()); + for (size_t j = 0; j < directCallablePGs.size(); ++j) { + CallableRecord callableRecord = {}; + optixSbtRecordPackHeader(directCallablePGs[j], &callableRecord); + callableRecords[j] = callableRecord; + } + directCallableRecordBuffer.allocUpload(callableRecords); + + sbt.callablesRecordBase = directCallableRecordBuffer.dPointer(); + sbt.callablesRecordStrideInBytes = sizeof(CallableRecord); + sbt.callablesRecordCount = + static_cast(directCallablePGs.size()); + } } protected: @@ -775,6 +786,7 @@ template class Trace { const std::string normModuleName = "normKernels.ptx"; std::string normKernelName = "normalize_surface_"; + std::string pipelineFileName = "GeneralPipeline"; }; } // namespace viennaray::gpu diff --git a/gpu/include/raygTriangleGeometry.hpp b/gpu/include/raygTriangleGeometry.hpp index 3c0b3a4..3b1ac45 100644 --- a/gpu/include/raygTriangleGeometry.hpp +++ b/gpu/include/raygTriangleGeometry.hpp @@ -263,6 +263,11 @@ struct TriangleGeometry { boundaryMesh.maximumExtent = bbMax; } +#ifndef NDEBUG + rayInternal::writeVTP(boundaryMesh, "triangleMesh_boundary.vtp", + std::vector()); +#endif + return boundaryMesh; } diff --git a/gpu/tests/boundaries/CMakeLists.txt b/gpu/tests/boundaries/CMakeLists.txt new file mode 100644 index 0000000..e388be4 --- /dev/null +++ b/gpu/tests/boundaries/CMakeLists.txt @@ -0,0 +1,8 @@ +project(boundaries LANGUAGES CXX) + +generate_pipeline(${PROJECT_NAME} generated_files TestPipelineTriangle.cu) +add_executable(${PROJECT_NAME} ${PROJECT_NAME}.cpp ${generated_files}) +target_link_libraries(${PROJECT_NAME} PRIVATE ViennaRay) + +add_dependencies(ViennaRay-GPU_Tests ${PROJECT_NAME}) +add_test(NAME ${PROJECT_NAME} COMMAND $) diff --git a/gpu/tests/boundaries/TestPipelineTriangle.cu b/gpu/tests/boundaries/TestPipelineTriangle.cu new file mode 100644 index 0000000..488e65d --- /dev/null +++ b/gpu/tests/boundaries/TestPipelineTriangle.cu @@ -0,0 +1,109 @@ +#include + +#ifndef __CUDACC__ +#define __CUDACC__ +#endif + +#include +#include +#include +#include +#include +#include +#include + +#include + +// #define COUNT_RAYS + +using namespace viennaray::gpu; + +extern "C" __constant__ viennaray::gpu::LaunchParams launchParams; + +extern "C" __global__ void __closesthit__() { + const HitSBTDataTriangle *sbtData = + (const HitSBTDataTriangle *)optixGetSbtDataPointer(); + PerRayData *prd = getPRD(); + + const unsigned int primID = optixGetPrimitiveIndex(); + prd->tMin = optixGetRayTmax(); + prd->primID = primID; + + if (sbtData->base.isBoundary) { + if (launchParams.periodicBoundary) { + applyPeriodicBoundary(prd, sbtData, launchParams.D); + } else { + reflectFromBoundary(prd, sbtData, launchParams.D); + } + } else { + atomicAdd(&launchParams.resultBuffer[prd->primID], prd->rayWeight); + prd->rayWeight = 0.f; + } +} + +extern "C" __global__ void __miss__() { getPRD()->rayWeight = 0.f; } + +extern "C" __global__ void __raygen__() { + const uint3 idx = optixGetLaunchIndex(); + const uint3 dims = optixGetLaunchDimensions(); + const int linearLaunchIndex = + idx.x + idx.y * dims.x + idx.z * dims.x * dims.y; + + // per-ray data + PerRayData prd; + // each ray has its own RNG state + initializeRNGState(&prd, linearLaunchIndex, launchParams.seed); + + // initialize ray position and direction + if (launchParams.D == 2) { + if (linearLaunchIndex == 0) { + prd.pos[0] = 0.5f; + prd.pos[1] = 1.1f; + prd.pos[2] = 0.f; + + prd.dir[0] = -1.f; + prd.dir[1] = -.5f; + prd.dir[2] = 0.f; + Normalize(prd.dir); + } else if (linearLaunchIndex == 1) { + prd.pos[0] = 0.5f; + prd.pos[1] = 1.5f; + prd.pos[2] = 0.f; + + prd.dir[0] = 0.6f; + prd.dir[1] = -.5f; + prd.dir[2] = 0.f; + Normalize(prd.dir); + } else { + return; + } + } else { + return; + } + + // the values we store the PRD pointer in: + uint32_t u0, u1; + packPointer((void *)&prd, u0, u1); + + while (continueRay(launchParams, prd)) { + // printf("Tracing ray %u from pos (%f, %f, %f) in dir (%f, %f, %f)\n", + // linearLaunchIndex, prd.pos[0], prd.pos[1], prd.pos[2], prd.dir[0], + // prd.dir[1], prd.dir[2]); + optixTrace(launchParams.traversable, // traversable GAS + make_float3(prd.pos[0], prd.pos[1], prd.pos[2]), // origin + make_float3(prd.dir[0], prd.dir[1], prd.dir[2]), // direction + 1e-4f, // tmin + 1e20f, // tmax + 0.0f, // rayTime + OptixVisibilityMask(255), + OPTIX_RAY_FLAG_DISABLE_ANYHIT, // OPTIX_RAY_FLAG_NONE, + 0, // SBT offset + 1, // SBT stride + 0, // missSBTIndex + u0, u1); // Payload + // printf("Ray %u hit primID %u with tMin %f\n", linearLaunchIndex, + // prd.primID, + // prd.tMin); + prd.numReflections++; + } +} diff --git a/gpu/tests/boundaries/boundaries.cpp b/gpu/tests/boundaries/boundaries.cpp new file mode 100644 index 0000000..9ca3654 --- /dev/null +++ b/gpu/tests/boundaries/boundaries.cpp @@ -0,0 +1,118 @@ +#include +#include + +#include + +using namespace viennaray; + +template TriangleMesh createGeometry() { + TriangleMesh mesh; + if constexpr (D == 3) { + mesh.nodes = {{1.f, 0.f, 0.f}, {0.f, 0.f, 0.f}, {1.f, 0.5f, 0.f}, + {0.f, 0.5f, 0.f}, {1.f, 0.5f, 1.f}, {0.f, 0.5f, 1.f}, + {1.f, 1.f, 1.f}, {0.f, 1.f, 1.f}}; + mesh.triangles = { + {0, 1, 2}, {1, 3, 2}, // bottom + {2, 4, 3}, {3, 4, 5}, // side + {5, 4, 6}, {5, 6, 7} // top + }; + mesh.minimumExtent = Vec3Df{0.f, 0.f, 0.f}; + mesh.maximumExtent = Vec3Df{1.f, 1.f, 1.f}; + mesh.gridDelta = 0.1f; + } else if constexpr (D == 2) { + LineMesh lineMesh; + lineMesh.nodes = { + {0.f, 0.f, 0.f}, {0.5f, 0.f, 0.f}, {0.5f, 1.f, 0.f}, {1.f, 1.f, 0.f}}; + lineMesh.lines = {{0, 1}, {1, 2}, {2, 3}}; + lineMesh.gridDelta = 0.5f; + lineMesh.calculateNormals(); + computeBoundingBox(lineMesh); + + mesh = convertLinesToTriangles(lineMesh); + + mesh.nodes.push_back(Vec3Df{0.6f, 0.f, -0.25f}); + mesh.nodes.push_back(Vec3Df{0.6f, 0.f, 0.25f}); + mesh.nodes.push_back(Vec3Df{0.6f, 1.f, 0.25f}); + mesh.nodes.push_back(Vec3Df{0.6f, 1.f, -0.25f}); + mesh.triangles.push_back(Vec3D{8, 9, 10}); + mesh.triangles.push_back(Vec3D{8, 11, 10}); + } + return mesh; +} + +int main() { + + auto context = DeviceContext::createContext("../../../lib/ptx", 0); + + gpu::Particle particle; + particle.name = "Particle"; + particle.dataLabels = {"flux"}; + std::unordered_map particleMap; + particleMap["Particle"] = 0; + + // // Run 3D test + // { + // constexpr int D = 3; + // TriangleMesh mesh = createGeometry(); + // std::vector flux(mesh.triangles.size(), 0.f); + // rayInternal::writeVTP(mesh, "boundary_3d.vtp", flux); + // } + + // Run 2D test + { + // reflective boundary test + constexpr int D = 2; + TriangleMesh mesh = createGeometry(); + + gpu::TraceTriangle tracer(context); + tracer.setPipelineFileName("TestPipeline"); + tracer.insertNextParticle(particle); + tracer.setParticleCallableMap({particleMap, {}}); + tracer.prepareParticlePrograms(); + + tracer.setNumberOfRaysFixed(2); + tracer.setGeometry(mesh); + + tracer.apply(); + + // auto flux = tracer.getFlux(0, 0, 0); + // std::cout << "Flux values at each triangle:" << std::endl; + // for (size_t i = 0; i < flux.size(); ++i) { + // std::cout << "Triangle " << i << ": " << flux[i] << std::endl; + // } + // rayInternal::writeVTP(mesh, "boundary_2d_reflective.vtp", flux); + + VC_TEST_ASSERT(flux[3] > 0.f); + VC_TEST_ASSERT(flux[5] > 0.f); + } + + { + // periodic boundary test + constexpr int D = 2; + TriangleMesh mesh = createGeometry(); + + gpu::TraceTriangle tracer(context); + tracer.setPeriodicBoundary(true); + tracer.setPipelineFileName("TestPipeline"); + tracer.insertNextParticle(particle); + tracer.setParticleCallableMap({particleMap, {}}); + tracer.prepareParticlePrograms(); + + tracer.setNumberOfRaysFixed(2); + tracer.setGeometry(mesh); + + tracer.apply(); + + // auto flux = tracer.getFlux(0, 0, 0); + // std::cout << "Flux values at each triangle:" << std::endl; + // for (size_t i = 0; i < flux.size(); ++i) { + // std::cout << "Triangle " << i << ": " << flux[i] << std::endl; + // } + // rayInternal::writeVTP(mesh, "boundary_2d_periodic.vtp", flux); + + VC_TEST_ASSERT(flux[3] > 0.f); + VC_TEST_ASSERT(flux[7] > 0.f); + } + + return 0; +} \ No newline at end of file diff --git a/include/viennaray/rayUtil.hpp b/include/viennaray/rayUtil.hpp index 0184156..6f3e584 100644 --- a/include/viennaray/rayUtil.hpp +++ b/include/viennaray/rayUtil.hpp @@ -4,6 +4,8 @@ #include #include +#include + #if VIENNARAY_EMBREE_VERSION < 4 #include #else @@ -558,6 +560,11 @@ void writeVTP(const std::string &filename, f.close(); } +void writeVTP(TriangleMesh const &mesh, const std::string &filename, + const std::vector &flux) { + writeVTP(filename, mesh.nodes, mesh.triangles, flux); +} + /* -------------------------------------------------------------- */ template From f468b1aa7e454c2b961fd7beab2eaa078c222063 Mon Sep 17 00:00:00 2001 From: reiter Date: Thu, 27 Nov 2025 17:23:02 +0100 Subject: [PATCH 2/4] Add test in 3D --- gpu/tests/boundaries/TestPipelineTriangle.cu | 22 +++++++++++- gpu/tests/boundaries/boundaries.cpp | 38 +++++++++++++++----- 2 files changed, 50 insertions(+), 10 deletions(-) diff --git a/gpu/tests/boundaries/TestPipelineTriangle.cu b/gpu/tests/boundaries/TestPipelineTriangle.cu index 488e65d..2eb5f3d 100644 --- a/gpu/tests/boundaries/TestPipelineTriangle.cu +++ b/gpu/tests/boundaries/TestPipelineTriangle.cu @@ -78,7 +78,27 @@ extern "C" __global__ void __raygen__() { return; } } else { - return; + if (linearLaunchIndex == 0) { + prd.pos[0] = 0.5f; + prd.pos[1] = 0.5f; + prd.pos[2] = 1.1f; + + prd.dir[0] = 0.f; + prd.dir[1] = -1.f; + prd.dir[2] = -.5f; + Normalize(prd.dir); + } else if (linearLaunchIndex == 1) { + prd.pos[0] = 0.5f; + prd.pos[1] = 0.5f; + prd.pos[2] = 1.5f; + + prd.dir[0] = 0.f; + prd.dir[1] = 0.6f; + prd.dir[2] = -.5f; + Normalize(prd.dir); + } else { + return; + } } // the values we store the PRD pointer in: diff --git a/gpu/tests/boundaries/boundaries.cpp b/gpu/tests/boundaries/boundaries.cpp index 9ca3654..2ccdf38 100644 --- a/gpu/tests/boundaries/boundaries.cpp +++ b/gpu/tests/boundaries/boundaries.cpp @@ -18,7 +18,8 @@ template TriangleMesh createGeometry() { }; mesh.minimumExtent = Vec3Df{0.f, 0.f, 0.f}; mesh.maximumExtent = Vec3Df{1.f, 1.f, 1.f}; - mesh.gridDelta = 0.1f; + mesh.gridDelta = 0.5f; + mesh.calculateNormals(); } else if constexpr (D == 2) { LineMesh lineMesh; lineMesh.nodes = { @@ -51,12 +52,31 @@ int main() { particleMap["Particle"] = 0; // // Run 3D test - // { - // constexpr int D = 3; - // TriangleMesh mesh = createGeometry(); - // std::vector flux(mesh.triangles.size(), 0.f); - // rayInternal::writeVTP(mesh, "boundary_3d.vtp", flux); - // } + { + constexpr int D = 3; + TriangleMesh mesh = createGeometry(); + + gpu::TraceTriangle tracer(context); + tracer.setPipelineFileName("TestPipeline"); + tracer.insertNextParticle(particle); + tracer.setParticleCallableMap({particleMap, {}}); + tracer.prepareParticlePrograms(); + + tracer.setNumberOfRaysFixed(2); + tracer.setGeometry(mesh); + + tracer.apply(); + + auto flux = tracer.getFlux(0, 0, 0); + // std::cout << "Flux values at each triangle:" << std::endl; + // for (size_t i = 0; i < flux.size(); ++i) { + // std::cout << "Triangle " << i << ": " << flux[i] << std::endl; + // } + // rayInternal::writeVTP(mesh, "boundary_3d_reflective.vtp", flux); + + VC_TEST_ASSERT(flux[3] > 0.f); + VC_TEST_ASSERT(flux[5] > 0.f); + } // Run 2D test { @@ -75,7 +95,7 @@ int main() { tracer.apply(); - // auto flux = tracer.getFlux(0, 0, 0); + auto flux = tracer.getFlux(0, 0, 0); // std::cout << "Flux values at each triangle:" << std::endl; // for (size_t i = 0; i < flux.size(); ++i) { // std::cout << "Triangle " << i << ": " << flux[i] << std::endl; @@ -103,7 +123,7 @@ int main() { tracer.apply(); - // auto flux = tracer.getFlux(0, 0, 0); + auto flux = tracer.getFlux(0, 0, 0); // std::cout << "Flux values at each triangle:" << std::endl; // for (size_t i = 0; i < flux.size(); ++i) { // std::cout << "Triangle " << i << ": " << flux[i] << std::endl; From a2afb058cbcb31fb55e3aa1712d2b59f9a2c8297 Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Fri, 28 Nov 2025 12:00:40 +0100 Subject: [PATCH 3/4] Clang-tidy fixes --- cmake/generate_ptx.cmake | 4 +- gpu/include/raygBoundary.hpp | 14 +-- gpu/include/raygCallableConfig.hpp | 3 +- gpu/include/raygDiskGeometry.hpp | 10 +- gpu/include/raygLineGeometry.hpp | 8 +- gpu/include/raygPerRayData.hpp | 18 ++-- gpu/include/raygSource.hpp | 23 ++--- gpu/include/raygTrace.hpp | 16 ++-- gpu/include/raygTraceDisk.hpp | 109 +++++++++++----------- gpu/include/raygTraceLine.hpp | 10 +- gpu/include/raygTraceTriangle.hpp | 14 +-- gpu/include/raygTriangleGeometry.hpp | 11 +-- include/viennaray/rayGeometry.hpp | 1 + include/viennaray/rayGeometryDisk.hpp | 6 +- include/viennaray/rayGeometryTriangle.hpp | 2 +- include/viennaray/rayMesh.hpp | 39 ++++---- include/viennaray/rayTrace.hpp | 2 +- include/viennaray/rayTraceDisk.hpp | 8 +- include/viennaray/rayTraceKernel.hpp | 3 +- include/viennaray/rayTraceTriangle.hpp | 13 ++- include/viennaray/rayTracingData.hpp | 1 - include/viennaray/rayUtil.hpp | 7 +- 22 files changed, 152 insertions(+), 170 deletions(-) diff --git a/cmake/generate_ptx.cmake b/cmake/generate_ptx.cmake index 0f05f62..c28aa17 100644 --- a/cmake/generate_ptx.cmake +++ b/cmake/generate_ptx.cmake @@ -38,7 +38,9 @@ function(generate_pipeline target_name generated_files) list(APPEND generated_files_local ${generated_ptx_files}) endif() - set(${generated_files} ${generated_files_local} PARENT_SCOPE) + set(${generated_files} + ${generated_files_local} + PARENT_SCOPE) endfunction() function(generate_kernel generated_files) diff --git a/gpu/include/raygBoundary.hpp b/gpu/include/raygBoundary.hpp index 8b4719c..305e41b 100644 --- a/gpu/include/raygBoundary.hpp +++ b/gpu/include/raygBoundary.hpp @@ -21,7 +21,7 @@ reflectFromBoundary(PerRayData *prd, const SBTData *hsd, const int D) { const unsigned int primID = optixGetPrimitiveIndex(); prd->numBoundaryHits++; - if constexpr (std::is_same::value) { + if constexpr (std::is_same_v) { prd->pos = prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold); if (primID == 0 || primID == 1) { @@ -29,12 +29,12 @@ reflectFromBoundary(PerRayData *prd, const SBTData *hsd, const int D) { } else if ((primID == 2 || primID == 3) && D == 3) { prd->dir[1] -= 2 * prd->dir[1]; // y boundary } - } else if constexpr (std::is_same::value) { + } else if constexpr (std::is_same_v) { prd->pos = prd->pos + prd->dir * optixGetRayTmax(); unsigned dim = primID / 4; prd->dir[dim] -= 2 * prd->dir[dim]; prd->pos[dim] = hsd->vertex[hsd->index[primID][0]][dim]; - } else if constexpr (std::is_same::value) { + } else if constexpr (std::is_same_v) { prd->pos = prd->pos + prd->dir * optixGetRayTmax(); if (primID == 0 || primID == 1) // x boundary prd->dir[0] -= 2 * prd->dir[0]; @@ -48,7 +48,7 @@ applyPeriodicBoundary(PerRayData *prd, const SBTData *hsd, const int D) { const unsigned int primID = optixGetPrimitiveIndex(); prd->numBoundaryHits++; - if constexpr (std::is_same::value) { + if constexpr (std::is_same_v) { prd->pos = prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold); if (primID == 0) { // xmin @@ -60,12 +60,12 @@ applyPeriodicBoundary(PerRayData *prd, const SBTData *hsd, const int D) { } else if (D == 3 && primID == 3) { // ymax prd->pos[1] = hsd->point[2][1]; } - } else if constexpr (std::is_same::value) { + } else if constexpr (std::is_same_v) { prd->pos = prd->pos + prd->dir * optixGetRayTmax(); unsigned dim = primID / 4; prd->pos[dim] = hsd->vertex[hsd->index[primID ^ 2][0]][dim]; - } else if constexpr (std::is_same::value) { - prd->pos = prd->pos + prd->dir * (optixGetRayTmax()); + } else if constexpr (std::is_same_v) { + prd->pos = prd->pos + prd->dir * optixGetRayTmax(); if (primID == 0) { // xmin prd->pos[0] = hsd->nodes[1][0]; } else if (primID == 1) { // xmax diff --git a/gpu/include/raygCallableConfig.hpp b/gpu/include/raygCallableConfig.hpp index 2b1c6d6..e09dc29 100644 --- a/gpu/include/raygCallableConfig.hpp +++ b/gpu/include/raygCallableConfig.hpp @@ -1,7 +1,6 @@ #pragma once -#include -#include +#include namespace viennaray::gpu { diff --git a/gpu/include/raygDiskGeometry.hpp b/gpu/include/raygDiskGeometry.hpp index 363bb8b..ee9e9ce 100644 --- a/gpu/include/raygDiskGeometry.hpp +++ b/gpu/include/raygDiskGeometry.hpp @@ -49,7 +49,6 @@ template struct DiskGeometry { // 2 inputs: one for the geometry, one for the boundary std::array diskInput{}; - std::array diskInputFlags{}; // ------------------- geometry input ------------------- // upload the model to the device: the builder @@ -105,11 +104,6 @@ template struct DiskGeometry { boundaryPointBuffer.allocUpload(boundaryMesh.nodes); boundaryNormalBuffer.allocUpload(boundaryMesh.normals); - // create local variables, because we need a *pointer* to the - // device pointers - CUdeviceptr d_boundPoints = boundaryPointBuffer.dPointer(); - CUdeviceptr d_boundNormals = boundaryNormalBuffer.dPointer(); - // AABB build input for boundary disks std::vector aabbBoundary(boundaryMesh.nodes.size()); for (size_t i = 0; i < boundaryMesh.nodes.size(); ++i) { @@ -217,12 +211,12 @@ template struct DiskGeometry { } // Find maximum extent in each dimension - Vec3Df extent = bbMax - bbMin; + const Vec3Df extent = bbMax - bbMin; float maxExtent = std::max(std::max(extent[0], extent[1]), extent[2]); // has to be the same as in raygTrace.hpp (hitGroupRecords) if constexpr (D == 2) { - boundaryMesh.radius = 0.5 * maxExtent; + boundaryMesh.radius = 0.5f * maxExtent; } else { boundaryMesh.radius = maxExtent * rayInternal::DiskFactor; } diff --git a/gpu/include/raygLineGeometry.hpp b/gpu/include/raygLineGeometry.hpp index d9055e1..ff396cd 100644 --- a/gpu/include/raygLineGeometry.hpp +++ b/gpu/include/raygLineGeometry.hpp @@ -25,7 +25,7 @@ struct LineGeometry { CudaBuffer asBuffer; /// build acceleration structure from triangle mesh - void buildAccel(DeviceContext &context, const LineMesh &mesh, + void buildAccel(const DeviceContext &context, const LineMesh &mesh, LaunchParams &launchParams, const bool ignoreBoundary, const float sourceOffset = 0.f) { assert(context.deviceID != -1 && "Context not initialized."); @@ -39,7 +39,6 @@ struct LineGeometry { // 2 inputs: one for the geometry, one for the boundary std::array lineInput{}; - std::array lineInputFlags{}; // ------------------- geometry input ------------------- // upload the model to the device: the builder @@ -47,11 +46,6 @@ struct LineGeometry { geometryLinesBuffer.allocUpload(mesh.lines); geometryNormalsBuffer.allocUpload(mesh.normals); - // create local variables, because we need a *pointer* to the - // device pointers - CUdeviceptr d_geoNodes = geometryNodesBuffer.dPointer(); - CUdeviceptr d_geoLines = geometryLinesBuffer.dPointer(); - // AABB build input std::vector aabb(mesh.lines.size()); diff --git a/gpu/include/raygPerRayData.hpp b/gpu/include/raygPerRayData.hpp index 9b25828..45e73da 100644 --- a/gpu/include/raygPerRayData.hpp +++ b/gpu/include/raygPerRayData.hpp @@ -1,11 +1,9 @@ #pragma once -#include - #include #include -#include +#include #define MAX_NEIGHBORS 8 @@ -44,29 +42,31 @@ struct PerRayData { // this can only get compiled if included in a cuda kernel #ifdef __CUDACC__ -static __forceinline__ __device__ void *unpackPointer(uint32_t i0, +#include + +static __device__ __forceinline__ void *unpackPointer(uint32_t i0, uint32_t i1) { const uint64_t uptr = static_cast(i0) << 32 | i1; void *ptr = reinterpret_cast(uptr); return ptr; } -static __forceinline__ __device__ void packPointer(void *ptr, uint32_t &i0, +static __device__ __forceinline__ void packPointer(void *ptr, uint32_t &i0, uint32_t &i1) { const uint64_t uptr = reinterpret_cast(ptr); i0 = uptr >> 32; i1 = uptr & 0x00000000ffffffff; } -static __forceinline__ __device__ PerRayData *getPRD() { +static __device__ __forceinline__ PerRayData *getPRD() { const uint32_t u0 = optixGetPayload_0(); const uint32_t u1 = optixGetPayload_1(); return reinterpret_cast(unpackPointer(u0, u1)); } -static __device__ void initializeRNGState(PerRayData *prd, - unsigned int linearLaunchIndex, - unsigned int seed) { +static __device__ __forceinline__ void +initializeRNGState(PerRayData *prd, unsigned int linearLaunchIndex, + unsigned int seed) { auto rngSeed = tea<3>(linearLaunchIndex, seed); curand_init(rngSeed, 0, 0, &prd->RNGstate); } diff --git a/gpu/include/raygSource.hpp b/gpu/include/raygSource.hpp index f5bbcf2..b780496 100644 --- a/gpu/include/raygSource.hpp +++ b/gpu/include/raygSource.hpp @@ -25,8 +25,8 @@ getOrthonormalBasis(const Vec3Df &n) { return {n, t, b2}; } -__device__ void initializeRayDirection(PerRayData *prd, const float power, - const uint16_t D) { +__device__ __forceinline__ void +initializeRayDirection(PerRayData *prd, const float power, const uint16_t D) { // source direction const float4 u = curand_uniform4(&prd->RNGstate); // (0,1] const float tt = powf(u.w, 2.f / (power + 1.f)); @@ -44,9 +44,9 @@ __device__ void initializeRayDirection(PerRayData *prd, const float power, Normalize(prd->dir); } -__device__ void initializeRayDirection(PerRayData *prd, const float power, - const std::array &basis, - const uint16_t D) { +__device__ __forceinline__ void +initializeRayDirection(PerRayData *prd, const float power, + const std::array &basis, const uint16_t D) { // source direction do { const float4 u = curand_uniform4(&prd->RNGstate); // (0,1] @@ -72,9 +72,9 @@ __device__ void initializeRayDirection(PerRayData *prd, const float power, Normalize(prd->dir); } -__device__ void initializeRayPosition(PerRayData *prd, - const LaunchParams::SourcePlane &source, - const uint16_t D) { +__device__ __forceinline__ void +initializeRayPosition(PerRayData *prd, const LaunchParams::SourcePlane &source, + const uint16_t D) { const float4 u = curand_uniform4(&prd->RNGstate); // (0,1] prd->pos[0] = source.minPoint[0] + u.x * (source.maxPoint[0] - source.minPoint[0]); @@ -90,8 +90,9 @@ __device__ void initializeRayPosition(PerRayData *prd, } // This is slightly faster because there is only one call to curand_uniform4 -__device__ void initializeRayPositionAndDirection(PerRayData *prd, - LaunchParams *launchParams) { +__device__ __forceinline__ void +initializeRayPositionAndDirection(PerRayData *prd, + const LaunchParams *launchParams) { const float4 u = curand_uniform4(&prd->RNGstate); // (0,1] prd->pos[0] = launchParams->source.minPoint[0] + u.x * (launchParams->source.maxPoint[0] - @@ -104,7 +105,7 @@ __device__ void initializeRayPositionAndDirection(PerRayData *prd, const float tt = powf(u.w, 2.f / (launchParams->cosineExponent + 1.f)); float s, c; __sincosf(2.f * M_PIf * u.z, &s, &c); - float sqrt1mtt = sqrtf(1 - tt); + const float sqrt1mtt = sqrtf(1 - tt); prd->dir[0] = c * sqrt1mtt; prd->dir[1] = s * sqrt1mtt; prd->dir[2] = -1.f * sqrtf(tt); diff --git a/gpu/include/raygTrace.hpp b/gpu/include/raygTrace.hpp index 569a55e..05d0fa2 100644 --- a/gpu/include/raygTrace.hpp +++ b/gpu/include/raygTrace.hpp @@ -29,13 +29,13 @@ using namespace viennacore; template class Trace { public: - Trace(std::shared_ptr &passedContext, + Trace(std::shared_ptr const &passedContext, std::string &&geometryType) : context_(passedContext), geometryType_(std::move(geometryType)) { initRayTracer(); } - Trace(std::string &&geometryType, unsigned deviceID = 0) + Trace(std::string &&geometryType, const int deviceID = 0) : geometryType_(std::move(geometryType)) { context_ = DeviceContext::getContextFromRegistry(deviceID); if (!context_) { @@ -48,7 +48,7 @@ template class Trace { initRayTracer(); } - ~Trace() { + virtual ~Trace() { freeBuffers(); destroyMembers(); } @@ -240,7 +240,8 @@ template class Trace { resultsDownloaded = false; } - void setElementData(CudaBuffer &passedCellDataBuffer, unsigned numData) { + void setElementData(const CudaBuffer &passedCellDataBuffer, + const unsigned numData) { if (passedCellDataBuffer.sizeInBytes / sizeof(float) / numData != launchParams.numElements) { Logger::getInstance() @@ -463,6 +464,7 @@ template class Trace { private: void initRayTracer() { + launchParams.D = D; context_->addModule(normModuleName); normKernelName.append(geometryType_ + "_f"); // launchParamsBuffer.alloc(sizeof(launchParams)); @@ -647,8 +649,8 @@ template class Trace { programGroups.push_back(missPG); programGroups.push_back(hitgroupPG); - for (size_t j = 0; j < directCallablePGs.size(); j++) { - programGroups.push_back(directCallablePGs[j]); + for (auto const &directCallablePG : directCallablePGs) { + programGroups.push_back(directCallablePG); } char log[2048]; @@ -780,7 +782,7 @@ template class Trace { bool ignoreBoundary = false; bool resultsDownloaded = false; - size_t numRays; + size_t numRays = 0; unsigned numCellData = 0; const std::string globalParamsName = "launchParams"; diff --git a/gpu/include/raygTraceDisk.hpp b/gpu/include/raygTraceDisk.hpp index 619b31a..4afa26a 100644 --- a/gpu/include/raygTraceDisk.hpp +++ b/gpu/include/raygTraceDisk.hpp @@ -3,21 +3,20 @@ #include "raygDiskGeometry.hpp" #include "raygTrace.hpp" -#include #include namespace viennaray::gpu { using namespace viennacore; -template class TraceDisk : public Trace { +template class TraceDisk final : public Trace { public: - TraceDisk(std::shared_ptr &passedContext) + explicit TraceDisk(std::shared_ptr &passedContext) : Trace(passedContext, "Disk") {} - TraceDisk(unsigned deviceID = 0) : Trace("Disk", deviceID) {} + explicit TraceDisk(unsigned deviceID = 0) : Trace("Disk", deviceID) {} - ~TraceDisk() { diskGeometry.freeBuffers(); } + ~TraceDisk() override { diskGeometry.freeBuffers(); } void setGeometry(const DiskMesh &passedMesh, float sourceOffset = 0.f) { assert(context_ && "Context not initialized."); @@ -38,7 +37,6 @@ template class TraceDisk : public Trace { maxBox[2] = diskMesh.gridDelta; } this->gridDelta_ = static_cast(diskMesh.gridDelta); - launchParams.D = D; pointNeighborhood_.template init<3>(diskMesh.nodes, 2 * diskMesh.radius, diskMesh.minimumExtent, diskMesh.maximumExtent); @@ -48,6 +46,7 @@ template class TraceDisk : public Trace { void smoothFlux(std::vector &flux, int smoothingNeighbors) override { auto oldFlux = flux; + const T requiredDistance = smoothingNeighbors * 2.0 * diskMesh.radius; PointNeighborhood *pointNeighborhood; // use pointer to avoid copies if (smoothingNeighbors == 1) { @@ -56,16 +55,17 @@ template class TraceDisk : public Trace { } else if (pointNeighborhoodCache_.getNumPoints() == launchParams.numElements && std::abs(pointNeighborhoodCache_.getDistance() - - smoothingNeighbors * 2 * diskMesh.radius) < 1e-6f) { + requiredDistance) < 1e-6) { // re-use cached neighborhood pointNeighborhood = &pointNeighborhoodCache_; } else { // create a new neighborhood with a larger radius and cache it - pointNeighborhoodCache_.template init<3>( - diskMesh.nodes, smoothingNeighbors * 2 * diskMesh.radius, - diskMesh.minimumExtent, diskMesh.maximumExtent); + pointNeighborhoodCache_.template init<3>(diskMesh.nodes, requiredDistance, + diskMesh.minimumExtent, + diskMesh.maximumExtent); pointNeighborhood = &pointNeighborhoodCache_; } + #pragma omp parallel for for (int idx = 0; idx < launchParams.numElements; idx++) { float vv = oldFlux[idx]; @@ -85,6 +85,8 @@ template class TraceDisk : public Trace { } void normalizeResults() override { + assert(resultBuffer.sizeInBytes != 0 && + "Normalization: Result buffer not initialized."); float sourceArea = 0.f; if constexpr (D == 2) { sourceArea = @@ -94,74 +96,68 @@ template class TraceDisk : public Trace { (launchParams.source.maxPoint[0] - launchParams.source.minPoint[0]) * (launchParams.source.maxPoint[1] - launchParams.source.minPoint[1]); } - assert(resultBuffer.sizeInBytes != 0 && - "Normalization: Result buffer not initialized."); - CUdeviceptr d_data = resultBuffer.dPointer(); - CUdeviceptr d_points = diskGeometry.geometryPointBuffer.dPointer(); - CUdeviceptr d_normals = diskGeometry.geometryNormalBuffer.dPointer(); // calculate areas on host and send to device for now - Vec2D bdBox = {minBox, maxBox}; + const Vec2D bdBox = {minBox, maxBox}; std::vector areas(launchParams.numElements); - DiskBoundingBoxXYIntersector bdDiskIntersector(bdBox); - - // 0 = REFLECTIVE, 1 = PERIODIC, 2 = IGNORE - std::array boundaryConds = { - BoundaryCondition::REFLECTIVE, BoundaryCondition::REFLECTIVE}; - const std::array boundaryDirs = {0, 1}; - constexpr float eps = 1e-4f; -#pragma omp for + DiskBoundingBoxXYIntersector xy_intersector(bdBox); + + const auto radius = diskMesh.radius; + constexpr std::array boundaryDirs = {0, 1}; +#pragma omp parallel for for (long idx = 0; idx < launchParams.numElements; ++idx) { - std::array disk{0.f, 0.f, 0.f, diskMesh.radius}; - Vec3Df coord = diskMesh.nodes[idx]; - Vec3Df normal = diskMesh.normals[idx]; - disk[0] = coord[0]; - disk[1] = coord[1]; - disk[2] = coord[2]; + const Vec3Df &coord = diskMesh.nodes[idx]; + const Vec3Df &normal = diskMesh.normals[idx]; if constexpr (D == 3) { - areas[idx] = disk[3] * disk[3] * M_PIf; // full disk area - if (boundaryConds[boundaryDirs[0]] == BoundaryCondition::IGNORE && - boundaryConds[boundaryDirs[1]] == BoundaryCondition::IGNORE) { + areas[idx] = radius * radius * M_PIf; // full disk area + if (this->ignoreBoundary) { // no boundaries continue; } - - if (boundaryDirs[0] != 2 && boundaryDirs[1] != 2) { - // Disk-BBox intersection only works with boundaries in x and y - // direction - areas[idx] = bdDiskIntersector.areaInside(disk, normal); + std::array disk{0.f, 0.f, 0.f, radius}; + disk[0] = coord[0]; + disk[1] = coord[1]; + disk[2] = coord[2]; + + // Disk-BBox intersection only works with boundaries in x and y + // direction + areas[idx] = xy_intersector.areaInside(disk, normal); + } else { + constexpr float eps = 1e-4f; + // 2D + areas[idx] = 2.f * radius; // full disk area + if (this->ignoreBoundary) { + // no boundaries continue; } - } else { // 2D - areas[idx] = 2 * disk[3]; // test min boundary - if ((boundaryConds[boundaryDirs[0]] != BoundaryCondition::IGNORE) && - (std::abs(disk[boundaryDirs[0]] - bdBox[0][boundaryDirs[0]]) < - disk[3])) { - T insideTest = 1 - normal[boundaryDirs[0]] * normal[boundaryDirs[0]]; + if (std::abs(coord[boundaryDirs[0]] - bdBox[0][boundaryDirs[0]]) < + radius) { + float insideTest = + 1.f - normal[boundaryDirs[0]] * normal[boundaryDirs[0]]; if (insideTest > eps) { insideTest = - std::abs(disk[boundaryDirs[0]] - bdBox[0][boundaryDirs[0]]) / + std::abs(coord[boundaryDirs[0]] - bdBox[0][boundaryDirs[0]]) / std::sqrt(insideTest); - if (insideTest < disk[3]) { - areas[idx] -= disk[3] - insideTest; + if (insideTest < radius) { + areas[idx] -= radius - insideTest; } } } // test max boundary - if ((boundaryConds[boundaryDirs[0]] != BoundaryCondition::IGNORE) && - (std::abs(disk[boundaryDirs[0]] - bdBox[1][boundaryDirs[0]]) < - disk[3])) { - T insideTest = 1 - normal[boundaryDirs[0]] * normal[boundaryDirs[0]]; + if (std::abs(coord[boundaryDirs[0]] - bdBox[1][boundaryDirs[0]]) < + radius) { + float insideTest = + 1.f - normal[boundaryDirs[0]] * normal[boundaryDirs[0]]; if (insideTest > eps) { insideTest = - std::abs(disk[boundaryDirs[0]] - bdBox[1][boundaryDirs[0]]) / + std::abs(coord[boundaryDirs[0]] - bdBox[1][boundaryDirs[0]]) / std::sqrt(insideTest); - if (insideTest < disk[3]) { - areas[idx] -= disk[3] - insideTest; + if (insideTest < radius) { + areas[idx] -= radius - insideTest; } } } @@ -171,6 +167,7 @@ template class TraceDisk : public Trace { CudaBuffer areaBuffer; areaBuffer.allocUpload(areas); CUdeviceptr d_areas = areaBuffer.dPointer(); + CUdeviceptr d_data = resultBuffer.dPointer(); void *kernel_args[] = { &d_data, &d_areas, &launchParams.numElements, @@ -222,8 +219,8 @@ template class TraceDisk : public Trace { PointNeighborhood pointNeighborhood_; PointNeighborhood pointNeighborhoodCache_; - Vec3Df minBox; - Vec3Df maxBox; + Vec3Df minBox{}; + Vec3Df maxBox{}; using Trace::context_; using Trace::geometryType_; diff --git a/gpu/include/raygTraceLine.hpp b/gpu/include/raygTraceLine.hpp index b33be0e..aa58be6 100644 --- a/gpu/include/raygTraceLine.hpp +++ b/gpu/include/raygTraceLine.hpp @@ -2,25 +2,23 @@ #include "raygLineGeometry.hpp" #include "raygTrace.hpp" -#include namespace viennaray::gpu { using namespace viennacore; -template class TraceLine : public Trace { +template class TraceLine final : public Trace { public: - TraceLine(std::shared_ptr passedContext) + explicit TraceLine(std::shared_ptr passedContext) : Trace(passedContext, "Line") {} - TraceLine(unsigned deviceID = 0) : Trace("Line", deviceID) {} + explicit TraceLine(int deviceID = 0) : Trace("Line", deviceID) {} - ~TraceLine() { lineGeometry.freeBuffers(); } + ~TraceLine() override { lineGeometry.freeBuffers(); } void setGeometry(const LineMesh &passedMesh, const float sourceOffset = 0.f) { this->gridDelta_ = static_cast(passedMesh.gridDelta); lineMesh = passedMesh; - launchParams.D = D; lineGeometry.buildAccel(*context_, lineMesh, launchParams, this->ignoreBoundary, sourceOffset); } diff --git a/gpu/include/raygTraceTriangle.hpp b/gpu/include/raygTraceTriangle.hpp index 277c6f2..b07a9ef 100644 --- a/gpu/include/raygTraceTriangle.hpp +++ b/gpu/include/raygTraceTriangle.hpp @@ -7,29 +7,29 @@ namespace viennaray::gpu { using namespace viennacore; -template class TraceTriangle : public Trace { +template class TraceTriangle final : public Trace { public: - TraceTriangle(std::shared_ptr &passedContext) + explicit TraceTriangle(std::shared_ptr &passedContext) : Trace(passedContext, "Triangle") { if constexpr (D == 2) { this->normKernelName.append("_2D"); } } - TraceTriangle(unsigned deviceID = 0) : Trace("Triangle", deviceID) { + explicit TraceTriangle(unsigned deviceID = 0) + : Trace("Triangle", deviceID) { if constexpr (D == 2) { this->normKernelName.append("_2D"); } } - ~TraceTriangle() { triangleGeometry.freeBuffers(); } + ~TraceTriangle() override { triangleGeometry.freeBuffers(); } void setGeometry(const TriangleMesh &passedMesh, const float sourceOffset = 0.f) { assert(context_); - assert(passedMesh.triangles.size() > 0 && - "Triangle mesh has no triangles."); - assert(passedMesh.nodes.size() > 0 && "Triangle mesh has no vertices."); + assert(!passedMesh.triangles.empty() && "Triangle mesh has no triangles."); + assert(!passedMesh.nodes.empty() && "Triangle mesh has no vertices."); this->gridDelta_ = static_cast(passedMesh.gridDelta); triangleGeometry.buildAccel(*context_, passedMesh, launchParams, diff --git a/gpu/include/raygTriangleGeometry.hpp b/gpu/include/raygTriangleGeometry.hpp index 3b1ac45..0a5bec5 100644 --- a/gpu/include/raygTriangleGeometry.hpp +++ b/gpu/include/raygTriangleGeometry.hpp @@ -46,7 +46,6 @@ struct TriangleGeometry { mesh.maximumExtent[2] + mesh.gridDelta + sourceOffset; } launchParams.numElements = mesh.triangles.size(); - launchParams.D = D; // 2 inputs: one for the geometry, one for the boundary std::array triangleInput{}; @@ -66,19 +65,18 @@ struct TriangleGeometry { // device pointers CUdeviceptr d_geoVertices = geometryVertexBuffer.dPointer(); CUdeviceptr d_geoIndices = geometryIndexBuffer.dPointer(); - CUdeviceptr d_geoNormals = geometryNormalBuffer.dPointer(); triangleInput[0].triangleArray.vertexFormat = OPTIX_VERTEX_FORMAT_FLOAT3; triangleInput[0].triangleArray.vertexStrideInBytes = sizeof(Vec3Df); triangleInput[0].triangleArray.numVertices = - (unsigned int)mesh.nodes.size(); + static_cast(mesh.nodes.size()); triangleInput[0].triangleArray.vertexBuffers = &d_geoVertices; triangleInput[0].triangleArray.indexFormat = OPTIX_INDICES_FORMAT_UNSIGNED_INT3; triangleInput[0].triangleArray.indexStrideInBytes = sizeof(Vec3D); triangleInput[0].triangleArray.numIndexTriplets = - (unsigned int)mesh.triangles.size(); + static_cast(mesh.triangles.size()); triangleInput[0].triangleArray.indexBuffer = d_geoIndices; // one SBT entry, and no per-primitive materials: @@ -108,14 +106,15 @@ struct TriangleGeometry { triangleInput[1].triangleArray.vertexFormat = OPTIX_VERTEX_FORMAT_FLOAT3; triangleInput[1].triangleArray.vertexStrideInBytes = sizeof(Vec3Df); - triangleInput[1].triangleArray.numVertices = (int)boundaryMesh.nodes.size(); + triangleInput[1].triangleArray.numVertices = + static_cast(boundaryMesh.nodes.size()); triangleInput[1].triangleArray.vertexBuffers = &d_boundVertices; triangleInput[1].triangleArray.indexFormat = OPTIX_INDICES_FORMAT_UNSIGNED_INT3; triangleInput[1].triangleArray.indexStrideInBytes = sizeof(Vec3D); triangleInput[1].triangleArray.numIndexTriplets = - (int)boundaryMesh.triangles.size(); + static_cast(boundaryMesh.triangles.size()); triangleInput[1].triangleArray.indexBuffer = d_boundIndices; // one SBT entry, and no per-primitive materials: diff --git a/include/viennaray/rayGeometry.hpp b/include/viennaray/rayGeometry.hpp index 86152d4..dd3f1f7 100644 --- a/include/viennaray/rayGeometry.hpp +++ b/include/viennaray/rayGeometry.hpp @@ -12,6 +12,7 @@ using namespace viennacore; template class Geometry { public: + virtual ~Geometry() = default; Geometry(GeometryType type) : geoType_(type) {} template diff --git a/include/viennaray/rayGeometryDisk.hpp b/include/viennaray/rayGeometryDisk.hpp index 39eb71c..87706eb 100644 --- a/include/viennaray/rayGeometryDisk.hpp +++ b/include/viennaray/rayGeometryDisk.hpp @@ -134,9 +134,9 @@ class GeometryDisk : public Geometry { } for (size_t i = 0; i < this->numPrimitives_; ++i) { - pPointBuffer_[i].xx = (float)points[i][0]; - pPointBuffer_[i].yy = (float)points[i][1]; - pPointBuffer_[i].radius = (float)discRadii_; + pPointBuffer_[i].xx = static_cast(points[i][0]); + pPointBuffer_[i].yy = static_cast(points[i][1]); + pPointBuffer_[i].radius = static_cast(discRadii_); if (points[i][0] < this->minCoords_[0]) this->minCoords_[0] = points[i][0]; if (points[i][1] < this->minCoords_[1]) diff --git a/include/viennaray/rayGeometryTriangle.hpp b/include/viennaray/rayGeometryTriangle.hpp index 911cc5d..d18b2e7 100644 --- a/include/viennaray/rayGeometryTriangle.hpp +++ b/include/viennaray/rayGeometryTriangle.hpp @@ -6,7 +6,7 @@ namespace viennaray { using namespace viennacore; -template +template class GeometryTriangle : public Geometry { public: GeometryTriangle() : Geometry(GeometryType::TRIANGLE) {} diff --git a/include/viennaray/rayMesh.hpp b/include/viennaray/rayMesh.hpp index 1427450..3730e4c 100644 --- a/include/viennaray/rayMesh.hpp +++ b/include/viennaray/rayMesh.hpp @@ -3,7 +3,6 @@ #include #include -#include #include namespace viennaray { @@ -26,7 +25,7 @@ template void computeBoundingBox(MeshType &mesh) { } struct LineMesh { - LineMesh() {} + LineMesh() = default; LineMesh(const std::vector &pts, const std::vector> &lns, float delta) : nodes(pts), lines(lns), gridDelta(delta) { @@ -39,12 +38,12 @@ struct LineMesh { std::vector> lines; std::vector normals; - Vec3Df minimumExtent; - Vec3Df maximumExtent; + Vec3Df minimumExtent{}; + Vec3Df maximumExtent{}; float gridDelta = 0.f; void calculateNormals() { - assert(lines.size() > 0 && "No lines in mesh."); + assert(!lines.empty() && "No lines in mesh."); normals.clear(); normals.resize(lines.size()); #pragma omp parallel for @@ -52,7 +51,7 @@ struct LineMesh { Vec3Df const &p0 = nodes[lines[i][0]]; Vec3Df const &p1 = nodes[lines[i][1]]; Vec3Df lineDir = p1 - p0; - Vec3Df normal = Vec3Df{-lineDir[1], lineDir[0], 0.0f}; + auto normal = Vec3Df{-lineDir[1], lineDir[0], 0.0f}; Normalize(normal); normals[i] = normal; } @@ -81,7 +80,7 @@ struct LineMesh { }; struct TriangleMesh { - TriangleMesh() {} + TriangleMesh() = default; TriangleMesh(std::vector const &pts, std::vector> const &tris, float delta) : nodes(pts), triangles(tris), gridDelta(delta) { @@ -93,12 +92,12 @@ struct TriangleMesh { std::vector> triangles; std::vector normals; - Vec3Df minimumExtent; - Vec3Df maximumExtent; + Vec3Df minimumExtent{}; + Vec3Df maximumExtent{}; float gridDelta = 0.f; void calculateNormals() { - assert(triangles.size() > 0 && "No triangles in mesh."); + assert(!triangles.empty() && "No triangles in mesh."); normals.clear(); normals.resize(triangles.size()); #pragma omp parallel for @@ -114,7 +113,7 @@ struct TriangleMesh { }; struct DiskMesh { - DiskMesh() {} + DiskMesh() = default; DiskMesh(const std::vector &pts, const std::vector &nms, float delta) : nodes(pts), normals(nms), gridDelta(delta) { @@ -125,13 +124,13 @@ struct DiskMesh { std::vector normals; std::vector radii; - Vec3Df minimumExtent; - Vec3Df maximumExtent; + Vec3Df minimumExtent{}; + Vec3Df maximumExtent{}; float radius = 0.f; float gridDelta = 0.f; }; -TriangleMesh convertLinesToTriangles(const LineMesh &lineMesh) { +inline TriangleMesh convertLinesToTriangles(const LineMesh &lineMesh) { TriangleMesh mesh; mesh.gridDelta = lineMesh.gridDelta; mesh.minimumExtent = lineMesh.minimumExtent; @@ -142,17 +141,17 @@ TriangleMesh convertLinesToTriangles(const LineMesh &lineMesh) { auto const &points = lineMesh.nodes; mesh.nodes.reserve(points.size() * 2); - for (size_t i = 0; i < points.size(); ++i) { - mesh.nodes.push_back(Vec3Df{points[i][0], points[i][1], lineWidth2}); - mesh.nodes.push_back(Vec3Df{points[i][0], points[i][1], -lineWidth2}); + for (auto const &point : points) { + mesh.nodes.push_back(Vec3Df{point[0], point[1], lineWidth2}); + mesh.nodes.push_back(Vec3Df{point[0], point[1], -lineWidth2}); } auto const &lines = lineMesh.lines; mesh.triangles.reserve(lines.size() * 2); mesh.normals.reserve(lines.size() * 2); - for (size_t i = 0; i < lines.size(); ++i) { - const unsigned p0 = lines[i][0] * 2; - const unsigned p1 = lines[i][1] * 2; + for (auto const &line : lines) { + const unsigned p0 = line[0] * 2; + const unsigned p1 = line[1] * 2; // first triangle Vec3D tri1{p0, p1, static_cast(p0 + 1)}; diff --git a/include/viennaray/rayTrace.hpp b/include/viennaray/rayTrace.hpp index 7bf795a..1c75745 100644 --- a/include/viennaray/rayTrace.hpp +++ b/include/viennaray/rayTrace.hpp @@ -26,7 +26,7 @@ template class Trace { Trace(Trace &&) = delete; Trace &operator=(Trace &&) = delete; - ~Trace() { rtcReleaseDevice(device_); } + virtual ~Trace() { rtcReleaseDevice(device_); } /// Run the ray tracer virtual void apply() {} diff --git a/include/viennaray/rayTraceDisk.hpp b/include/viennaray/rayTraceDisk.hpp index ead52d6..95cd414 100644 --- a/include/viennaray/rayTraceDisk.hpp +++ b/include/viennaray/rayTraceDisk.hpp @@ -10,10 +10,10 @@ namespace viennaray { using namespace viennacore; template -class TraceDisk : public Trace { +class TraceDisk final : public Trace { public: - TraceDisk() {} - ~TraceDisk() { geometry_.releaseGeometry(); } + TraceDisk() = default; + ~TraceDisk() override { geometry_.releaseGeometry(); } /// Run the ray tracer void apply() override { @@ -51,7 +51,7 @@ class TraceDisk : public Trace { this->config_, this->dataLog_, this->RTInfo_); tracer.setTracingData(&this->localData_, this->pGlobalData_); tracer.apply(); - this->config_.runNumber++; + ++this->config_.runNumber; boundary.releaseGeometry(); } diff --git a/include/viennaray/rayTraceKernel.hpp b/include/viennaray/rayTraceKernel.hpp index eee350b..738cef8 100644 --- a/include/viennaray/rayTraceKernel.hpp +++ b/include/viennaray/rayTraceKernel.hpp @@ -1,7 +1,6 @@ #pragma once #include -#include #include #include #include @@ -421,7 +420,7 @@ template class TraceKernel { traceInfo_.reflections = totalReflections; traceInfo_.time = static_cast(timer.currentDuration) * 1e-9; - if (raysTerminated > 1e3) { + if (raysTerminated > 1000) { Logger::getInstance() .addDebug( "A total of " + std::to_string(raysTerminated) + diff --git a/include/viennaray/rayTraceTriangle.hpp b/include/viennaray/rayTraceTriangle.hpp index 14aa470..4306154 100644 --- a/include/viennaray/rayTraceTriangle.hpp +++ b/include/viennaray/rayTraceTriangle.hpp @@ -10,10 +10,10 @@ namespace viennaray { using namespace viennacore; template -class TraceTriangle : public Trace { +class TraceTriangle final : public Trace { public: - TraceTriangle() {} - ~TraceTriangle() { geometry_.releaseGeometry(); } + TraceTriangle() = default; + ~TraceTriangle() override { geometry_.releaseGeometry(); } /// Run the ray tracer void apply() override { @@ -57,7 +57,7 @@ class TraceTriangle : public Trace { this->config_, this->dataLog_, this->RTInfo_); tracer.setTracingData(&this->localData_, this->pGlobalData_); tracer.apply(); - this->config_.runNumber++; + ++this->config_.runNumber; boundary.releaseGeometry(); } @@ -135,9 +135,8 @@ class TraceTriangle : public Trace { /// Helper function to smooth the recorded flux by averaging over the /// neighborhood in a post-processing step. - void smoothFlux(std::vector &flux, - int numNeighbors = 1) override { - // no smoothing + void smoothFlux(std::vector &flux, int numNeighbors) override { + // no smoothing on elements } private: diff --git a/include/viennaray/rayTracingData.hpp b/include/viennaray/rayTracingData.hpp index 5b5f4f1..5ed0524 100644 --- a/include/viennaray/rayTracingData.hpp +++ b/include/viennaray/rayTracingData.hpp @@ -14,7 +14,6 @@ enum class TracingDataMergeEnum : unsigned { }; template class TracingData { -private: using scalarDataType = NumericType; using vectorDataType = std::vector; using mergeType = std::vector; diff --git a/include/viennaray/rayUtil.hpp b/include/viennaray/rayUtil.hpp index 6f3e584..53b934d 100644 --- a/include/viennaray/rayUtil.hpp +++ b/include/viennaray/rayUtil.hpp @@ -17,7 +17,6 @@ #include #endif -#include #include #include #include @@ -560,8 +559,8 @@ void writeVTP(const std::string &filename, f.close(); } -void writeVTP(TriangleMesh const &mesh, const std::string &filename, - const std::vector &flux) { +inline void writeVTP(TriangleMesh const &mesh, const std::string &filename, + const std::vector &flux) { writeVTP(filename, mesh.nodes, mesh.triangles, flux); } @@ -581,7 +580,7 @@ createSourceGrid(const std::array, 2> &pBdBox, auto firstDir = pTraceSettings[1]; auto secondDir = pTraceSettings[2]; auto minMax = pTraceSettings[3]; - assert((!(D == 2) || rayDir != 2) && "Source direction z in 2D geometry"); + assert((D != 2 || rayDir != 2) && "Source direction z in 2D geometry"); auto len1 = pBdBox[1][firstDir] - pBdBox[0][firstDir]; auto len2 = pBdBox[1][secondDir] - pBdBox[0][secondDir]; From aca2e71719f5e6faadafc529c3a88a96710a7f28 Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Fri, 28 Nov 2025 19:02:48 +0100 Subject: [PATCH 4/4] Bump version --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 2c8bfd4..7b8a6a8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.20 FATAL_ERROR) project( ViennaRay LANGUAGES CXX - VERSION 3.8.2) + VERSION 3.8.3) # -------------------------------------------------------------------------------------------------------- # Library switches