GPU: get rid of MediumTransitionQueue (and kernel)

Instead, just push those rays onto the ray queue for the next ray depth.
This gives up to a 5% speed up for scenes with participating media.

src/pbrt/gpu/accel.cpp

@@ -1120,8 +1120,8 @@ void GPUAccel::IntersectClosest(
     int maxRays, EscapedRayQueue *escapedRayQueue, HitAreaLightQueue *hitAreaLightQueue,
     MaterialEvalQueue *basicEvalMaterialQueue,
     MaterialEvalQueue *universalEvalMaterialQueue,
-    MediumTransitionQueue *mediumTransitionQueue, MediumSampleQueue *mediumSampleQueue,
-    RayQueue *rayQueue) const {
+    MediumSampleQueue *mediumSampleQueue,
+    RayQueue *rayQueue, RayQueue *nextRayQueue) const {
     std::pair<cudaEvent_t, cudaEvent_t> events =
         GetProfilerEvents("Tracing closest hit rays");
 
@@ -1131,11 +1131,11 @@ void GPUAccel::IntersectClosest(
         RayIntersectParameters params;
         params.traversable = rootTraversable;
         params.rayQueue = rayQueue;
+        params.nextRayQueue = nextRayQueue;
         params.escapedRayQueue = escapedRayQueue;
         params.hitAreaLightQueue = hitAreaLightQueue;
         params.basicEvalMaterialQueue = basicEvalMaterialQueue;
         params.universalEvalMaterialQueue = universalEvalMaterialQueue;
-        params.mediumTransitionQueue = mediumTransitionQueue;
         params.mediumSampleQueue = mediumSampleQueue;
 
         ParamBufferState &pbs = getParamBuffer(params);

src/pbrt/gpu/accel.h

@@ -41,8 +41,7 @@ class GPUAccel {
         int maxRays, EscapedRayQueue *escapedRayQueue,
         HitAreaLightQueue *hitAreaLightQueue, MaterialEvalQueue *basicEvalMaterialQueue,
         MaterialEvalQueue *universalEvalMaterialQueue,
-        MediumTransitionQueue *mediumTransitionQueue,
-        MediumSampleQueue *mediumSampleQueue, RayQueue *rayQueue) const;
+        MediumSampleQueue *mediumSampleQueue, RayQueue *rayQueue, RayQueue *nextRayQueue) const;
 
     void IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue) const;
 

src/pbrt/gpu/media.cpp

@@ -36,6 +36,7 @@ static inline void rescale(SampledSpectrum &T_hat, SampledSpectrum &lightPathPDF
 
 // GPUPathIntegrator Participating Media Methods
 void GPUPathIntegrator::SampleMediumInteraction(int depth) {
+    RayQueue *nextRayQueue = NextRayQueue(depth);
     ForAllQueued(
         "Sample medium interaction", mediumSampleQueue, maxQueueSize,
         PBRT_GPU_LAMBDA(MediumSampleWorkItem w, int index) {
@@ -177,17 +178,10 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                 Interaction intr(w.pi, w.n);
                 intr.mediumInterface = &w.mediumInterface;
                 Ray newRay = intr.SpawnRay(ray.d);
-                mediumTransitionQueue->Push(MediumTransitionWorkItem{
-                    newRay, lambda, T_hat, uniPathPDF, lightPathPDF, w.prevIntrCtx,
-                    w.isSpecularBounce, w.anyNonSpecularBounces, w.etaScale,
-                    w.pixelIndex});
-#if 0
-                // WHY NOT THIS?
-                rayQueues[(depth + 1) & 1]->PushIndirect(newRay, w.prevIntrCtx,
-                                                         T_hat, uniPathPDF, lightPathPDF, lambda, w.etaScale,
-                                                         w.isSpecularBounce, w.anyNonSpecularBounces,
-                                                         w.pixelIndex);
-#endif
+                nextRayQueue->PushIndirect(newRay, w.prevIntrCtx, T_hat, uniPathPDF,
+                                           lightPathPDF, lambda, w.etaScale,
+                                           w.isSpecularBounce, w.anyNonSpecularBounces,
+                                           w.pixelIndex);
                 return;
             }
 
@@ -226,7 +220,6 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
         return;
 
     RayQueue *currentRayQueue = CurrentRayQueue(depth);
-    RayQueue *nextRayQueue = NextRayQueue(depth);
 
     using PhaseFunction = HGPhaseFunction;
     std::string desc = std::string("Sample direct/indirect - Henyey Greenstein");
@@ -318,24 +311,4 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
         });
 }
 
-void GPUPathIntegrator::HandleMediumTransitions(int depth) {
-    RayQueue *rayQueue = NextRayQueue(depth);
-
-    ForAllQueued(
-        "Handle medium transitions", mediumTransitionQueue, maxQueueSize,
-        PBRT_GPU_LAMBDA(MediumTransitionWorkItem mt, int index) {
-            // Have to do this here, later, since we can't be writing into
-            // the other ray queue in optix closest hit.  (Wait--really?
-            // Why not? Basically boils down to current indirect enqueue (and other
-            // places?))
-            // TODO: figure this out...
-            rayQueue->PushIndirect(mt.ray, mt.prevIntrCtx, mt.T_hat, mt.uniPathPDF,
-                                   mt.lightPathPDF, mt.lambda, mt.etaScale,
-                                   mt.isSpecularBounce, mt.anyNonSpecularBounces,
-                                   mt.pixelIndex);
-            PBRT_DBG("Enqueuied ray after medium transition at depth %d pixel index %d",
-                     depth + 1, mt.pixelIndex);
-        });
-}
-
 }  // namespace pbrt

src/pbrt/gpu/optix.cu

@@ -198,9 +198,9 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
         PBRT_DBG("Enqueuing into medium transition queue: ray index %d pixel index %d \n",
                  rayIndex, r.pixelIndex);
         Ray newRay = intr.SpawnRay(r.ray.d);
-        params.mediumTransitionQueue->Push(MediumTransitionWorkItem{
-            newRay, r.lambda, r.T_hat, r.uniPathPDF, r.lightPathPDF, r.prevIntrCtx,
-            r.isSpecularBounce, r.anyNonSpecularBounces, r.etaScale, r.pixelIndex});
+        params.nextRayQueue->PushIndirect(
+            newRay, r.prevIntrCtx, r.T_hat, r.uniPathPDF, r.lightPathPDF, r.lambda,
+            r.etaScale, r.isSpecularBounce, r.anyNonSpecularBounces, r.pixelIndex);
         return;
     }
 

src/pbrt/gpu/optix.h

@@ -53,10 +53,10 @@ struct RayIntersectParameters {
     RayQueue *rayQueue;
 
     // closest hit
+    RayQueue *nextRayQueue;
     EscapedRayQueue *escapedRayQueue;
     HitAreaLightQueue *hitAreaLightQueue;
     MaterialEvalQueue *basicEvalMaterialQueue, *universalEvalMaterialQueue;
-    MediumTransitionQueue *mediumTransitionQueue;
     MediumSampleQueue *mediumSampleQueue;
 
     // shadow rays

src/pbrt/gpu/pathintegrator.cpp

@@ -228,8 +228,6 @@ GPUPathIntegrator::GPUPathIntegrator(Allocator alloc, const ParsedScene &scene)
         pstd::MakeConstSpan(&haveUniversalEvalMaterial[1],
                             haveUniversalEvalMaterial.size() - 1));
 
-    // Always allocate this, even if no media
-    mediumTransitionQueue = alloc.new_object<MediumTransitionQueue>(maxQueueSize, alloc);
     if (haveMedia) {
         mediumSampleQueue = alloc.new_object<MediumSampleQueue>(maxQueueSize, alloc);
         mediumScatterQueue = alloc.new_object<MediumScatterQueue>(maxQueueSize, alloc);
@@ -354,7 +352,6 @@ void GPUPathIntegrator::Render() {
                     "Reset queues before tracing rays", PBRT_GPU_LAMBDA() {
                         nextQueue->Reset();
                         // Reset queues before tracing next batch of rays
-                        mediumTransitionQueue->Reset();
                         if (mediumSampleQueue)
                             mediumSampleQueue->Reset();
                         if (mediumScatterQueue)
@@ -378,8 +375,8 @@ void GPUPathIntegrator::Render() {
                 // Find closest intersections along active rays
                 accel->IntersectClosest(maxQueueSize, escapedRayQueue, hitAreaLightQueue,
                                         basicEvalMaterialQueue,
-                                        universalEvalMaterialQueue, mediumTransitionQueue,
-                                        mediumSampleQueue, CurrentRayQueue(depth));
+                                        universalEvalMaterialQueue, mediumSampleQueue,
+                                        CurrentRayQueue(depth), NextRayQueue(depth));
 
                 if (depth > 0) {
                     // As above, with the indexing...
@@ -399,7 +396,6 @@ void GPUPathIntegrator::Render() {
                 EvaluateMaterialsAndBSDFs(depth);
                 // Do immediately so that we have space for shadow rays for subsurface..
                 TraceShadowRays(depth);
-                HandleMediumTransitions(depth);
                 if (haveSubsurface)
                     SampleSubsurface(depth);
             }

src/pbrt/gpu/pathintegrator.h

@@ -42,7 +42,6 @@ class GPUPathIntegrator {
 
     void TraceShadowRays(int depth);
     void SampleMediumInteraction(int depth);
-    void HandleMediumTransitions(int depth);
     void SampleSubsurface(int depth);
 
     void HandleEscapedRays(int depth);
@@ -103,7 +102,6 @@ class GPUPathIntegrator {
 
     RayQueue *rayQueues[2];
 
-    MediumTransitionQueue *mediumTransitionQueue = nullptr;
     MediumSampleQueue *mediumSampleQueue = nullptr;
     MediumScatterQueue *mediumScatterQueue = nullptr;
 

src/pbrt/gpu/workitems.h

@@ -248,21 +248,6 @@ struct SubsurfaceScatterWorkItem {
     int pixelIndex;
 };
 
-// MediumTransitionWorkItem Definition
-struct MediumTransitionWorkItem {
-    Ray ray;
-    SampledWavelengths lambda;
-    SampledSpectrum T_hat, uniPathPDF, lightPathPDF;
-    LightSampleContext prevIntrCtx;
-    int isSpecularBounce;
-    int anyNonSpecularBounces;
-    Float etaScale;
-    int pixelIndex;
-};
-
-// MediumTransitionQueue Definition
-using MediumTransitionQueue = WorkQueue<MediumTransitionWorkItem>;
-
 // MediumSampleWorkItem Definition
 struct MediumSampleWorkItem {
     // Both enqueue types (have mtl and no hit)

src/pbrt/gpu/workitems.soa

@@ -110,17 +110,6 @@ soa SubsurfaceScatterWorkItem {
     SubsurfaceInteraction ssi;
 };
 
-soa MediumTransitionWorkItem {
-    Ray ray;
-    SampledWavelengths lambda;
-    SampledSpectrum T_hat, uniPathPDF, lightPathPDF;
-    LightSampleContext prevIntrCtx;
-    int isSpecularBounce;
-    int anyNonSpecularBounces;
-    Float etaScale;
-    int pixelIndex;
-};
-
 soa MediumSampleWorkItem {
     Ray ray;
     Float tMax;