GPU: update pixel sample radiance estimates due to shadow rays in Opt…

…iX code

This saves launching a separate kernel to do so, which isn't really
necessary and gives a few percent performance benefit on simple scenes.
(There is no meaningful benefit with more complex scenes, but it doesn't
hurt...)

src/pbrt/gpu/accel.cpp

@@ -1164,7 +1164,8 @@ void GPUAccel::IntersectClosest(
     cudaEventRecord(events.second);
 };
 
-void GPUAccel::IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue) const {
+void GPUAccel::IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue,
+                               SOA<PixelSampleState> *pixelSampleState) const {
     std::pair<cudaEvent_t, cudaEvent_t> events = GetProfilerEvents("Tracing shadow rays");
 
     cudaEventRecord(events.first);
@@ -1173,6 +1174,7 @@ void GPUAccel::IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue) cons
         RayIntersectParameters params;
         params.traversable = rootTraversable;
         params.shadowRayQueue = shadowRayQueue;
+        params.pixelSampleState = pixelSampleState;
 
         ParamBufferState &pbs = getParamBuffer(params);
 
@@ -1200,7 +1202,8 @@ void GPUAccel::IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue) cons
     cudaEventRecord(events.second);
 }
 
-void GPUAccel::IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue) const {
+void GPUAccel::IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue,
+                                 SOA<PixelSampleState> *pixelSampleState) const {
     std::pair<cudaEvent_t, cudaEvent_t> events = GetProfilerEvents("Tracing shadow Tr rays");
 
     cudaEventRecord(events.first);
@@ -1209,6 +1212,7 @@ void GPUAccel::IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue) co
         RayIntersectParameters params;
         params.traversable = rootTraversable;
         params.shadowRayQueue = shadowRayQueue;
+        params.pixelSampleState = pixelSampleState;
 
         ParamBufferState &pbs = getParamBuffer(params);
 

src/pbrt/gpu/accel.h

@@ -43,9 +43,11 @@ class GPUAccel {
         MaterialEvalQueue *universalEvalMaterialQueue,
         MediumSampleQueue *mediumSampleQueue, RayQueue *rayQueue, RayQueue *nextRayQueue) const;
 
-    void IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue) const;
+    void IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue,
+                         SOA<PixelSampleState> *pixelSampleState) const;
 
-    void IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue) const;
+    void IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue,
+                           SOA<PixelSampleState> *pixelSampleState) const;
 
     void IntersectOneRandom(int maxRays, SubsurfaceScatterQueue *subsurfaceScatterQueue) const;
 

src/pbrt/gpu/optix.cu

@@ -357,12 +357,12 @@ extern "C" __global__ void __raygen__shadow() {
                  sr.Ld[0], sr.Ld[1], sr.Ld[2], sr.Ld[3],
                  sr.uniPathPDF[0], sr.uniPathPDF[1], sr.uniPathPDF[2], sr.uniPathPDF[3],
                  sr.lightPathPDF[0], sr.lightPathPDF[1], sr.lightPathPDF[2], sr.lightPathPDF[3]);
+
+        SampledSpectrum Lpixel = params.pixelSampleState->L[sr.pixelIndex];
+        params.pixelSampleState->L[sr.pixelIndex] = Lpixel + Ld;
     } else {
         PBRT_DBG("Shadow ray was occluded\n");
-        Ld = SampledSpectrum(0.);
     }
-
-    params.shadowRayQueue->Ld[index] = Ld;
 }
 
 extern "C" __global__ void __miss__shadow() {
@@ -503,17 +503,17 @@ extern "C" __global__ void __raygen__shadow_Tr() {
              T_ray[2] / (sr.uniPathPDF * uniPathPDF + sr.lightPathPDF * lightPathPDF).Average(),
              T_ray[3] / (sr.uniPathPDF * uniPathPDF + sr.lightPathPDF * lightPathPDF).Average());
 
-    if (!T_ray)
-        Ld = SampledSpectrum(0.f);
-    else
+    if (T_ray) {
         // FIXME/reconcile: this takes lightPathPDF as input while
         // e.g. VolPathIntegrator::SampleLd() does not...
         Ld *= T_ray / (sr.uniPathPDF * uniPathPDF + sr.lightPathPDF * lightPathPDF).Average();
 
-    PBRT_DBG("Setting final Ld for shadow ray index %d pixel index %d = as %f %f %f %f\n",
-             index, sr.pixelIndex, Ld[0], Ld[1], Ld[2], Ld[3]);
+        PBRT_DBG("Setting final Ld for shadow ray index %d pixel index %d = as %f %f %f %f\n",
+                 index, sr.pixelIndex, Ld[0], Ld[1], Ld[2], Ld[3]);
 
-    params.shadowRayQueue->Ld[index] = Ld;
+        SampledSpectrum Lpixel = params.pixelSampleState->L[sr.pixelIndex];
+        params.pixelSampleState->L[sr.pixelIndex] = Lpixel + Ld;
+    }
 }
 
 extern "C" __global__ void __miss__shadow_Tr() {

src/pbrt/gpu/optix.h

@@ -61,6 +61,7 @@ struct RayIntersectParameters {
 
     // shadow rays
     ShadowRayQueue *shadowRayQueue;
+    SOA<PixelSampleState> *pixelSampleState;
 
     // Subsurface scattering...
     SubsurfaceScatterQueue *subsurfaceScatterQueue;

src/pbrt/gpu/pathintegrator.cpp

@@ -523,22 +523,9 @@ void GPUPathIntegrator::HandleRayFoundEmission(int depth) {
 
 void GPUPathIntegrator::TraceShadowRays(int depth) {
     if (haveMedia)
-        accel->IntersectShadowTr(maxQueueSize, shadowRayQueue);
+        accel->IntersectShadowTr(maxQueueSize, shadowRayQueue, &pixelSampleState);
     else
-        accel->IntersectShadow(maxQueueSize, shadowRayQueue);
-    // Add shadow ray radiance contributions to pixels
-    ForAllQueued(
-        "Incorporate shadow ray contribution", shadowRayQueue, maxQueueSize,
-        PBRT_GPU_LAMBDA(const ShadowRayWorkItem sr) {
-            if (!sr.Ld)
-                return;
-            SampledSpectrum Lpixel = pixelSampleState.L[sr.pixelIndex];
-
-            PBRT_DBG("Adding shadow ray Ld %f %f %f %f at pixel index %d \n", sr.Ld[0],
-                     sr.Ld[1], sr.Ld[2], sr.Ld[3], sr.pixelIndex);
-
-            pixelSampleState.L[sr.pixelIndex] = Lpixel + sr.Ld;
-        });
+        accel->IntersectShadow(maxQueueSize, shadowRayQueue, &pixelSampleState);
 
     // Reset shadow ray queue
     GPUDo(