fix self shadowing in gaussian GI mode

Author: Ka1serM

assets/tests/hotel_old_mill_realistic.nrscene

@@ -2,17 +2,17 @@
   "objects": [
     {
       "type": "camera",
-      "name": "Canon Automotive Fisheye",
+      "name": "Samy & Gao 220deg",
       "position": [0.0, 2.0, 5.0],
       "rotation_euler": [0.0, 0.0, 0.0],
       "scale": [1.0, 1.0, 1.0],
       "camera": {
         "projection": "realistic",
         "active": true,
         "focus_distance_cm": 500.0,
-        "exposure": 3.5,
+        "exposure": 4.0,
         "resolution": [1928, 1088],
-        "lens": "/home/marcel/GitRepositories/NoorRay/external/ROSS/resources/lenses/canon_automotive_fisheye/canon_automotive_fisheye.zmx",
+        "lens": "/home/marcel/GitRepositories/ROSS/resources/lenses/samy_gao/Samy&Gao-1(220deg).zmx",
         "sensor": "/home/marcel/GitRepositories/NoorRay/external/ROSS/resources/sensors/onsemi_AR0237.json",
         "glass_catalogs": "/home/marcel/GitRepositories/NoorRay/external/ROSS/resources/glasscatalogs/ohara.AGF;/home/marcel/GitRepositories/NoorRay/external/ROSS/resources/glasscatalogs/schott.AGF"
       }

src/libnoorray/Kernels/Connect.cu

@@ -31,7 +31,8 @@ NR_GPU inline bool shadowOccluded(
         const uint32_t gaussianSampleIndex = gaussianEnabled ? randomUint(rng) : 0;
         const RayHit hit = intersectRay(params.scene.tlasHandle, shadow.origin,
             shadow.direction, rayMin, shadow.tMax, gaussianSampleIndex,
-            gaussianEnabled, meshVisibilityBoundEnabled);
+            gaussianEnabled, meshVisibilityBoundEnabled,
+            shadow.excludedGaussianId);
         if (hit.instanceIndex == InvalidIndex)
             return false;
 

src/libnoorray/Kernels/GaussianHit.cu

@@ -16,6 +16,15 @@ extern "C" __global__ void __anyhit__gaussian()
     const uint32_t instanceId = optixGetInstanceId();
     const uint32_t globalGaussianId = instanceId;
 
+    // Shadow rays start inside the Gaussian that produced the scattering
+    // event. Exclude that exact instance instead of relying on a world-space
+    // epsilon, which cannot cover differently scaled Gaussians robustly.
+    if (globalGaussianId == optixGetPayload_2())
+    {
+        optixIgnoreIntersection();
+        return;
+    }
+
     const float opacity = params.scene.gaussianOpacities[globalGaussianId];
     if (opacity <= 0.0f)
     {

src/libnoorray/Kernels/Shade.cu

@@ -142,6 +142,7 @@ NR_GPU_KERNEL void shadeKernel(const KernelParams params)
         ShadowWorkItem shadow{};
         shadow.tMin = 1.0f;
         shadow.tMax = 0.0f;
+        shadow.excludedGaussianId = InvalidIndex;
         if (mayWriteShadowQueue)
             params.queues.shadowQueue[index].tMax = 0.0f;
         PathState state = params.queues.pathStates[hit.sampleIndex];
@@ -269,6 +270,7 @@ NR_GPU_KERNEL void shadeKernel(const KernelParams params)
                             state.throughput * brdf * lightSample.radiance;
                         shadow.rngState     = shadowRng;
                         shadow.sampleIndex  = hit.sampleIndex;
+                        shadow.excludedGaussianId = gaussianId;
                         params.queues.shadowQueue[index] = shadow;
                     }
                 }

src/libnoorray/Raytracing/RayTraversal.h

@@ -53,7 +53,8 @@ NR_GPU inline RayHit intersectRay(
     const float tMax,
     const uint32_t sampleIndex = 0,
     const bool gaussianEnabled = true,
-    const bool meshVisibilityBoundEnabled = true)
+    const bool meshVisibilityBoundEnabled = true,
+    const uint32_t excludedGaussianId = InvalidIndex)
 {
     RayHit hit{};
 #if defined(NR_GPU_DEVICE_COMPILE)
@@ -90,7 +91,10 @@ NR_GPU inline RayHit intersectRay(
         const float gaussianTMax = meshHit.instanceIndex != InvalidIndex ? meshHit.t : tMax;
         uint32_t payload0 = sampleIndex;
         uint32_t payload1 = __float_as_uint(gaussianTMax);
-        uint32_t payload2 = InvalidIndex;
+        // Before acceptance payload 2 identifies a Gaussian that must not
+        // intersect this ray. Any-hit overwrites it only when another
+        // Gaussian is accepted.
+        uint32_t payload2 = excludedGaussianId;
         optixTraverse(
             accel,
             make_float3(origin.x, origin.y, origin.z),
@@ -105,12 +109,12 @@ NR_GPU inline RayHit intersectRay(
             0,
             payload0, payload1, payload2);
 
-        const uint32_t gaussianId = payload2;
-        if (gaussianId != InvalidIndex)
+        const float gaussianT = __uint_as_float(payload1);
+        if (gaussianT < gaussianTMax)
         {
             // Gaussian accepted via Russian roulette.
-            hit.instanceIndex = gaussianId;
-            hit.t = __uint_as_float(payload1);
+            hit.instanceIndex = payload2;
+            hit.t = gaussianT;
         }
         else if (meshHit.instanceIndex != InvalidIndex)
         {

src/libnoorray/Raytracing/Types.h

@@ -85,6 +85,7 @@ struct alignas(16) ShadowWorkItem
     SampledSpectrum contribution;
     RandomState rngState;
     uint32_t sampleIndex;
+    uint32_t excludedGaussianId;
 };
 
 struct WavefrontQueues