Add boundary hit termination, small fixes #81

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged

tobre1 merged 21 commits into main from fix-warnings

Nov 10, 2025

.github/workflows/test.yml

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -43,14 +43,20 @@ jobs:
  
            with:

              os: ${{ matrix.os }}

          - name: 📋 Install MacOS Dependencies

            shell: bash

            if: ${{ matrix.os == 'macos-latest' }}

            run: |

              brew install embree tbb

          - name: 🦥 Cache Dependencies

            uses: actions/cache@v4

            with:

              key: test-${{ matrix.os }}-${{ matrix.config }}

              path: build

          - name: 🏗️ Compile

            run: cmake -B build -DCPM_SOURCE_CACHE=deps-cache -DVIENNARAY_BUILD_TESTS=ON && cmake --build build --config ${{ matrix.config }}

            run: cmake -DVIENNARAY_BUILD_TESTS=ON -B build && cmake --build build --config ${{ matrix.config }}

          - name: 🧪 Test

            run: ctest -C ${{ matrix.config }} --test-dir build

            run: ctest -E "Benchmark|Performance" -C ${{ matrix.config }} --test-dir build

CMakeLists.txt

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -2,7 +2,7 @@ cmake_minimum_required(VERSION 3.20 FATAL_ERROR)
  
    project(

      ViennaRay

      LANGUAGES CXX

      VERSION 3.7.0)

      VERSION 3.7.1)

    # --------------------------------------------------------------------------------------------------------

    # Library switches

    @@ -92,7 +92,7 @@ include("cmake/cpm.cmake")
  
    CPMAddPackage(

      NAME ViennaCore

      VERSION 1.6.2

      VERSION 1.6.3

      GIT_REPOSITORY "https://github.com/ViennaTools/ViennaCore"

      OPTIONS "VIENNACORE_USE_GPU ${VIENNARAY_USE_GPU}")

README.md

-Original file line number
+Diff line change
@@ Expand Up @@
     * Installation with CPM
       ```cmake
-      CPMAddPackage("gh:viennatools/viennaray@3.7.0")
+      CPMAddPackage("gh:viennatools/viennaray@3.7.1") # Use the latest release version
       ```
     * With a local installation
@@ Expand Down @@

examples/disk2D/disk2D.cpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -49,7 +49,12 @@ int main() { @@
       rayTracer.setNumberOfRaysPerPoint(2000);
       // Run the ray tracer
+      Timer timer;
+      timer.start();
       rayTracer.apply();
+      timer.finish();
+      std::cout << "Tracing time: " << timer.currentDuration / 1e9 << " s\n";
       // Extract the normalized hit counts for each geometry point
       auto &flux = rayTracer.getLocalData().getVectorData("flux");
@@ Expand Down @@

examples/triangle3D/CMakeLists.txt

-Original file line number
+Diff line change
@@ Expand Up / @@ -2,7 +2,6 @@ project(triangle3D LANGUAGES CXX) @@
     add_executable(${PROJECT_NAME} "${PROJECT_NAME}.cpp")
     target_link_libraries(${PROJECT_NAME} PRIVATE ViennaRay)
-    configure_file(${CMAKE_SOURCE_DIR}/gpu/examples/Resources/trenchMesh.dat
-                   ${CMAKE_CURRENT_BINARY_DIR}/trenchMesh.dat COPYONLY)
+    configure_file(trenchMesh.dat ${CMAKE_CURRENT_BINARY_DIR}/trenchMesh.dat COPYONLY)
     add_dependencies(ViennaRay_Examples ${PROJECT_NAME})

gpu/examples/Resources/trenchMesh.dat → examples/triangle3D/trenchMesh.dat

-Original file line number
+Diff line change
@@ -1,6 +1,6 @@
     grid_delta 1
     min_extent -25 -25 -40
-    max_extent 25 25 2.22507e-308
+    max_extent 25 25 0
     n_nodes 6579
     n_triangles 12800
     n -7 -24 -40
@@ Expand Down @@

gpu/examples/CMakeLists.txt

-Original file line number
+Diff line change
@@ Expand Up / @@ -3,7 +3,8 @@ project(ViennaRay-GPU_Examples) @@
     add_custom_target(ViennaRay-GPU_Examples ALL)
     add_gpu_executable(GPU_trenchTriangles target_name trenchTriangles.cpp)
-    configure_file(Resources/trenchMesh.dat ${CMAKE_CURRENT_BINARY_DIR}/trenchMesh.dat COPYONLY)
+    configure_file(${CMAKE_SOURCE_DIR}/examples/triangle3D/trenchMesh.dat
+                   ${CMAKE_CURRENT_BINARY_DIR}/trenchMesh.dat COPYONLY)
     add_dependencies(ViennaRay-GPU_Examples ${target_name})
     add_gpu_executable(GPU_trenchDisks target_name trenchDisks.cpp)
@@ Expand Down @@

gpu/examples/Resources/CMakeLists.txt

This file was deleted.

gpu/include/raygBoundary.hpp

-Original file line number
+Diff line change
@@ Expand Up @@
     // this can only get compiled if included in a cuda kernel
     #ifdef __CUDACC__
+    namespace viennaray::gpu {
+    using namespace viennacore;
     template <typename SBTData>
-    __device__ __inline__ void reflectFromBoundary(viennaray::gpu::PerRayData *prd,
-                                                   const SBTData *hsd,
-                                                   const int D) {
-      using namespace viennacore;
+    __device__ __inline__ void
+    reflectFromBoundary(PerRayData *prd, const SBTData *hsd, const int D) {
       const unsigned int primID = optixGetPrimitiveIndex();
       prd->numBoundaryHits++;
-      if constexpr (std::is_same<SBTData, viennaray::gpu::HitSBTDataDisk>::value) {
+      if constexpr (std::is_same<SBTData, HitSBTDataDisk>::value) {
         prd->pos =
             prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold);
         if (primID == 0 || primID == 1) {
           prd->dir[0] -= 2 * prd->dir[0]; // x boundary
         } else if ((primID == 2 || primID == 3) && D == 3) {
           prd->dir[1] -= 2 * prd->dir[1]; // y boundary
         }
-      } else if constexpr (std::is_same<
-                               SBTData,
-                               viennaray::gpu::HitSBTDataTriangle>::value) {
+      } else if constexpr (std::is_same<SBTData, HitSBTDataTriangle>::value) {
         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<SBTData,
-                                        viennaray::gpu::HitSBTDataLine>::value) {
+      } else if constexpr (std::is_same<SBTData, HitSBTDataLine>::value) {
         prd->pos = prd->pos + prd->dir * (optixGetRayTmax());
         if (primID == 0 || primID == 1) // x boundary
           prd->dir[0] -= 2 * prd->dir[0];
@@ Expand All @@
     template <typename SBTData>
     __device__ __inline__ void
-    applyPeriodicBoundary(viennaray::gpu::PerRayData *prd, const SBTData *hsd,
-                          const int D) {
+    applyPeriodicBoundary(PerRayData *prd, const SBTData *hsd, const int D) {
       using namespace viennacore;
       const unsigned int primID = optixGetPrimitiveIndex();
       prd->numBoundaryHits++;
-      if constexpr (std::is_same<SBTData, viennaray::gpu::HitSBTDataDisk>::value) {
+      if constexpr (std::is_same<SBTData, HitSBTDataDisk>::value) {
         prd->pos =
             prd->pos + prd->dir * (optixGetRayTmax() - launchParams.tThreshold);
         if (primID == 0) { // xmin
@@ Expand All @@
         } else if (D == 3 && primID == 3) { // ymax
           prd->pos[1] = hsd->point[2][1];
         }
-      } else if constexpr (std::is_same<
-                               SBTData,
-                               viennaray::gpu::HitSBTDataTriangle>::value) {
+      } else if constexpr (std::is_same<SBTData, HitSBTDataTriangle>::value) {
         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<SBTData,
-                                        viennaray::gpu::HitSBTDataLine>::value) {
+      } else if constexpr (std::is_same<SBTData, HitSBTDataLine>::value) {
         prd->pos = prd->pos + prd->dir * (optixGetRayTmax());
         if (primID == 0) { // xmin
           prd->pos[0] = hsd->nodes[1][0];
@@ Expand All @@
         }
       }
     }
+    } // namespace viennaray::gpu
     #endif

gpu/include/raygPerRayData.hpp

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -12,12 +12,14 @@
  
    namespace viennaray::gpu {

    using namespace viennacore;

    // Per-ray data structure associated with each ray. Should be kept small to

    // optimize memory usage and performance.

    struct PerRayData {

      // Position and direction

      viennacore::Vec3Df pos;

      viennacore::Vec3Df dir;

      Vec3Df pos;

      Vec3Df dir;

      // Simulation specific data

      float rayWeight = 1.f;

    @@ -40,8 +42,6 @@ struct PerRayData {
  
      bool hitFromBack = false;

    };

    } // namespace viennaray::gpu

    // this can only get compiled if included in a cuda kernel

    #ifdef __CUDACC__

    static __forceinline__ __device__ void *unpackPointer(uint32_t i0,

    @@ -64,10 +64,12 @@ template <typename T> static __forceinline__ __device__ T *getPRD() {
  
      return reinterpret_cast<T *>(unpackPointer(u0, u1));

    }

    static __device__ void initializeRNGState(viennaray::gpu::PerRayData *prd,

    static __device__ void initializeRNGState(PerRayData *prd,

                                              unsigned int linearLaunchIndex,

                                              unsigned int seed) {

      auto rngSeed = tea<4>(linearLaunchIndex, seed);

      curand_init(rngSeed, 0, 0, &prd->RNGstate);

    }

    #endif

    } // namespace viennaray::gpu

gpu/include/raygRNG.hpp

-Original file line number
+Diff line change
@@ Expand Up / @@ -7,8 +7,6 @@ namespace viennaray::gpu { @@
     typedef curandStatePhilox4_32_10_t RNGState;
-    }
     // Other possible RNGState types:
     // typedef curandStateXORWOW_t curtRNGState; // bad
     // typedef curandStateMRG32k3a_t curtRNGState // not tested
@@ Expand All / @@ -17,10 +15,8 @@ typedef curandStatePhilox4_32_10_t RNGState; @@
     #ifdef __CUDACC__
     template <unsigned int N>
-    static __device__ __inline__ unsigned int tea(unsigned int val0,
-                                                  unsigned int val1) {
-      unsigned int v0 = val0;
-      unsigned int v1 = val1;
+    static __device__ __inline__ unsigned int tea(unsigned int v0,
+                                                  unsigned int v1) {
       unsigned int s0 = 0;
       for (unsigned int n = 0; n < N; n++) {
@@ Expand All @@
       return v0;
     }
-    __device__ __inline__ float getNextRand(viennaray::gpu::RNGState *state) {
+    __device__ __inline__ float getNextRand(RNGState *state) {
       return (float)(curand_uniform(state));
     }
-    __device__ __inline__ float getNormalDistRand(viennaray::gpu::RNGState *state) {
+    __device__ __inline__ float getNormalDistRand(RNGState *state) {
       float4 u0 = curand_uniform4(state);
       float r = sqrtf(-2.f * logf(u0.x));
       float theta = 2.f * M_PIf * u0.y;
       return r * sinf(theta);
     }
     #endif
+    } // namespace viennaray::gpu

gpu/include/raygReflection.hpp

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -1,5 +1,8 @@
  
    #pragma once

    // Device code for handling reflections on various geometry types

    #ifdef __CUDACC__

    #include <curand.h>

    #include "raygPerRayData.hpp"

    @@ -8,12 +11,12 @@
  
    #include <vcVectorType.hpp>

    using namespace viennaray::gpu;

    namespace viennaray::gpu {

    #ifdef __CUDACC__

    __device__ __inline__ viennacore::Vec3Df

    computeNormal(const void *sbtData, const unsigned int primID) {

      using namespace viennacore;

    using namespace viennacore;

    __device__ __inline__ Vec3Df computeNormal(const void *sbtData,

                                               const unsigned int primID) {

      const HitSBTDataBase *baseData =

          reinterpret_cast<const HitSBTDataBase *>(sbtData);

      switch (baseData->geometryType) {

    @@ -48,37 +51,32 @@ computeNormal(const void *sbtData, const unsigned int primID) {
  
      }

    }

    __device__ __forceinline__ viennacore::Vec3Df

    getNormal(const void *sbtData, const unsigned int primID) {

    __device__ __forceinline__ Vec3Df getNormal(const void *sbtData,

                                                const unsigned int primID) {

      return reinterpret_cast<const HitSBTDataBase *>(sbtData)->normal[primID];

    }

    static __device__ __forceinline__ void

    specularReflection(viennaray::gpu::PerRayData *prd,

                       const viennacore::Vec3Df &geoNormal) {

      using namespace viennacore;

    specularReflection(PerRayData *prd, const Vec3Df &geoNormal) {

    #ifndef VIENNARAY_TEST

      prd->pos = prd->pos + prd->tMin * prd->dir;

    #endif

      prd->dir = prd->dir - (2 * DotProduct(prd->dir, geoNormal)) * geoNormal;

    }

    static __device__ viennacore::Vec3Df

    PickRandomPointOnUnitSphere(viennaray::gpu::RNGState *state) {

    static __device__ Vec3Df PickRandomPointOnUnitSphere(RNGState *state) {

      const float4 u = curand_uniform4(state); // (0,1]

      const float z = 1.0f - 2.0f * u.x;       // uniform in [-1,1]

      const float r2 = fmaxf(0.0f, 1.0f - z * z);

      const float r = sqrtf(r2);

      const float phi = 2.0f * M_PIf * u.y;

      float s, c;

      __sincosf(phi, &s, &c); // branch-free sin+cos

      return viennacore::Vec3Df{r * c, r * s, z};

      return Vec3Df{r * c, r * s, z};

    }

    static __device__ void diffuseReflection(viennaray::gpu::PerRayData *prd,

                                             const viennacore::Vec3Df &geoNormal,

                                             const uint8_t D) {

      using namespace viennacore;

    static __device__ void

    diffuseReflection(PerRayData *prd, const Vec3Df &geoNormal, const uint8_t D) {

    #ifndef VIENNARAY_TEST

      prd->pos = prd->pos + prd->tMin * prd->dir;

    #endif

    @@ -91,22 +89,17 @@ static __device__ void diffuseReflection(viennaray::gpu::PerRayData *prd,
  
      Normalize(prd->dir);

    }

    static __device__ void diffuseReflection(viennaray::gpu::PerRayData *prd,

                                             const uint8_t D) {

      using namespace viennacore;

    static __device__ void diffuseReflection(PerRayData *prd, const uint8_t D) {

      const viennaray::gpu::HitSBTDataDisk *sbtData =

          (const viennaray::gpu::HitSBTDataDisk *)optixGetSbtDataPointer();

      const HitSBTDataDisk *sbtData =

          (const HitSBTDataDisk *)optixGetSbtDataPointer();

      const Vec3Df geoNormal = computeNormal(sbtData, optixGetPrimitiveIndex());

      diffuseReflection(prd, geoNormal, D);

    }

    static __device__ __forceinline__ void

    conedCosineReflection(viennaray::gpu::PerRayData *prd,

                          const viennacore::Vec3Df &geomNormal,

    conedCosineReflection(PerRayData *prd, const Vec3Df &geomNormal,

                          const float maxConeAngle, const uint8_t D) {

      using namespace viennacore;

      // Calculate specular direction

      specularReflection(prd, geomNormal);

    @@ -165,4 +158,5 @@ conedCosineReflection(viennaray::gpu::PerRayData *prd,
  
      Normalize(dir);

      prd->dir = dir;

    }

    } // namespace viennaray::gpu

    #endif

gpu/include/raygSBTRecords.hpp

            
                      Original file line number
                      Diff line number
                      Diff line change
                  
    @@ -5,29 +5,31 @@
  
    namespace viennaray::gpu {

    using namespace viennacore;

    struct HitSBTDataBase {

      void *cellData;

      bool isBoundary;

      int geometryType;

      viennacore::Vec3Df *normal; // optional normal buffer

      Vec3Df *normal; // optional normal buffer

    };

    struct HitSBTDataDisk {

      HitSBTDataBase base;

      viennacore::Vec3Df *point;

      Vec3Df *point;

      float radius;

    };

    struct HitSBTDataTriangle {

      HitSBTDataBase base;

      viennacore::Vec3Df *vertex;

      viennacore::Vec3D<unsigned> *index;

      Vec3Df *vertex;

      Vec3D<unsigned> *index;

    };

    struct HitSBTDataLine {

      HitSBTDataBase base;

      viennacore::Vec3Df *nodes;

      viennacore::Vec2D<unsigned> *lines;

      Vec3Df *nodes;

      Vec2D<unsigned> *lines;

    };

    // SBT record for a raygen program

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add boundary hit termination, small fixes #81

Uh oh!

Diff view

Diff view

There are no files selected for viewing

Uh oh!

Uh oh!