Update from book source.

No functionality changes.
Cleaned up some of the GPU *WorkItem definitions.

src/pbrt/gpu/camera.cpp

@@ -17,20 +17,14 @@ namespace pbrt {
 
 // GPUPathIntegrator Camera Ray Methods
 void GPUPathIntegrator::GenerateCameraRays(int y0, int sampleIndex) {
+    // Define _generateRays_ lambda function
     auto generateRays = [=](auto sampler) {
         using Sampler = std::remove_reference_t<decltype(*sampler)>;
         if constexpr (!std::is_same_v<Sampler, MLTSampler> &&
                       !std::is_same_v<Sampler, DebugMLTSampler>)
             GenerateCameraRays<Sampler>(y0, sampleIndex);
     };
-    // Somewhat surprisingly, GenerateCameraRays() is specialized on the
-    // type of the Sampler being used and not on, say, the Camera.  By
-    // specializing on the sampler type, the particular Sampler used can be
-    // stack allocated (rather than living in global memory), which in turn
-    // allows its state to be stored in registers in the
-    // GenerateCameraRays() kernel. There's little benefit from
-    // specializing on the Camera since its state is read-only and shared
-    // among all of the threads, so caches well in practice.
+
     sampler.DispatchCPU(generateRays);
 }
 
@@ -39,18 +33,15 @@ void GPUPathIntegrator::GenerateCameraRays(int y0, int sampleIndex) {
     RayQueue *rayQueue = CurrentRayQueue(0);
     GPUParallelFor(
         "Generate Camera rays", maxQueueSize, PBRT_GPU_LAMBDA(int pixelIndex) {
-            // Initialize _pPixel_ and test against pixel bounds
-            Vector2i resolution = film.PixelBounds().Diagonal();
+            // Enqueue camera ray and set pixel state for sample
+            // Compute pixel coordinates for _pixelIndex_
             Bounds2i pixelBounds = film.PixelBounds();
-
-            Point2i pPixel(pixelBounds.pMin.x + int(pixelIndex) % resolution.x,
-                           pixelBounds.pMin.y + y0 + int(pixelIndex) / resolution.x);
+            int xResolution = pixelBounds.pMax.x - pixelBounds.pMin.x;
+            Point2i pPixel(pixelBounds.pMin.x + pixelIndex % xResolution,
+                           y0 + pixelIndex / xResolution);
             pixelSampleState.pPixel[pixelIndex] = pPixel;
 
-            // If we've split the image into multiple spans of scanlines,
-            // then in the final pass, we may have a few more threads
-            // launched than there are remaining pixels. Bail out without
-            // enqueuing a ray if so.
+            // Test pixel coordinates against pixel bounds
             if (!InsideExclusive(pPixel, pixelBounds))
                 return;
 
@@ -72,7 +63,6 @@ void GPUPathIntegrator::GenerateCameraRays(int y0, int sampleIndex) {
                 camera.GenerateRay(cameraSample, lambda);
 
             // Initialize remainder of _PixelSampleState_ for ray
-            // Initialize the rest of the pixel sample's state.
             pixelSampleState.L[pixelIndex] = SampledSpectrum(0.f);
             pixelSampleState.lambda[pixelIndex] = lambda;
             pixelSampleState.filterWeight[pixelIndex] = cameraSample.weight;

src/pbrt/gpu/media.cpp

@@ -167,8 +167,8 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                     PBRT_DBG("Adding ray to escapedRayQueue pixel index %d depth %d\n",
                              ms.pixelIndex, depth);
                     escapedRayQueue->Push(EscapedRayWorkItem{
-                        beta, uniPathPDF, lightPathPDF, lambda, ray.o, ray.d, ms.piPrev,
-                        ms.nPrev, ms.nsPrev, (int)ms.isSpecularBounce, ms.pixelIndex});
+                        beta, uniPathPDF, lightPathPDF, lambda, ray.o, ray.d,
+                        ms.prevIntrCtx, (int)ms.isSpecularBounce, ms.pixelIndex});
                 }
             }
 
@@ -178,12 +178,12 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                 intr.mediumInterface = &ms.mediumInterface;
                 Ray newRay = intr.SpawnRay(ray.d);
                 mediumTransitionQueue->Push(MediumTransitionWorkItem{
-                    newRay, lambda, beta, uniPathPDF, lightPathPDF, ms.piPrev, ms.nPrev,
-                    ms.nsPrev, ms.isSpecularBounce, ms.anyNonSpecularBounces, ms.etaScale,
+                    newRay, lambda, beta, uniPathPDF, lightPathPDF, ms.prevIntrCtx,
+                    ms.isSpecularBounce, ms.anyNonSpecularBounces, ms.etaScale,
                     ms.pixelIndex});
 #if 0
                 // WHY NOT THIS?
-                rayQueues[(depth + 1) & 1]->PushIndirect(newRay, ms.piPrev, ms.nPrev, ms.nsPrev,
+                rayQueues[(depth + 1) & 1]->PushIndirect(newRay, ms.prevIntrCtx,
                                                          beta, uniPathPDF, lightPathPDF, lambda, ms.etaScale,
                                                          ms.isSpecularBounce, ms.anyNonSpecularBounces,
                                                          ms.pixelIndex);
@@ -196,11 +196,10 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                     "Ray hit an area light: adding to hitAreaLightQueue pixel index %d "
                     "depth %d\n",
                     ms.pixelIndex, depth);
-                // TODO: intr.wo == -ray.d?
                 hitAreaLightQueue->Push(HitAreaLightWorkItem{
                     ms.areaLight, lambda, beta, uniPathPDF, lightPathPDF, Point3f(ms.pi),
-                    ms.n, ms.uv, -ray.d, ms.piPrev, ray.d, ms.nPrev, ms.nsPrev,
-                    ms.isSpecularBounce, ms.pixelIndex});
+                    ms.n, ms.uv, -ray.d, ms.prevIntrCtx, ms.isSpecularBounce,
+                    ms.pixelIndex});
             }
 
             FloatTextureHandle displacement = material.GetDisplacement();
@@ -312,8 +311,7 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
             bool anyNonSpecularBounces = true;
 
             // Spawn indirect ray.
-            nextRayQueue->PushIndirect(ray, Point3fi(ms.p), Normal3f(0, 0, 0),
-                                       Normal3f(0, 0, 0), beta, uniPathPDF, lightPathPDF,
+            nextRayQueue->PushIndirect(ray, ctx, beta, uniPathPDF, lightPathPDF,
                                        ms.lambda, ms.etaScale, isSpecularBounce,
                                        anyNonSpecularBounces, ms.pixelIndex);
             PBRT_DBG("Enqueuing indirect medium ray at depth %d pixel index %d\n",
@@ -332,8 +330,8 @@ void GPUPathIntegrator::HandleMediumTransitions(int depth) {
             // Why not? Basically boils down to current indirect enqueue (and other
             // places?))
             // TODO: figure this out...
-            rayQueue->PushIndirect(mt.ray, mt.piPrev, mt.nPrev, mt.nsPrev, mt.beta,
-                                   mt.uniPathPDF, mt.lightPathPDF, mt.lambda, mt.etaScale,
+            rayQueue->PushIndirect(mt.ray, mt.prevIntrCtx, mt.beta, 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",

src/pbrt/gpu/optix.cu

@@ -117,15 +117,15 @@ extern "C" __global__ void __raygen__findClosest() {
                 r.ray.o.x, r.ray.o.y, r.ray.o.z, r.ray.d.x, r.ray.d.y, r.ray.d.z,
                 r.beta[0], r.beta[1], r.beta[2], r.beta[3]);
             params.mediumSampleQueue->Push(r.ray, Infinity, r.lambda, r.beta, r.uniPathPDF,
-                                           r.lightPathPDF, r.pixelIndex, r.piPrev,
-                                           r.nPrev, r.nsPrev, r.isSpecularBounce,
+                                           r.lightPathPDF, r.pixelIndex, r.prevIntrCtx,
+                                           r.isSpecularBounce,
                                            r.anyNonSpecularBounces, r.etaScale);
         } else if (params.escapedRayQueue) {
             PBRT_DBG("Adding ray to escapedRayQueue ray index %d pixel index %d\n", rayIndex,
                      r.pixelIndex);
             params.escapedRayQueue->Push(EscapedRayWorkItem{
-                r.beta, r.uniPathPDF, r.lightPathPDF, r.lambda, ray.o, ray.d, r.piPrev, r.nPrev,
-                r.nsPrev, (int)r.isSpecularBounce, r.pixelIndex});
+                r.beta, r.uniPathPDF, r.lightPathPDF, r.lambda, ray.o, ray.d, r.prevIntrCtx,
+                (int)r.isSpecularBounce, r.pixelIndex});
         }
     }
 }
@@ -166,9 +166,7 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
                                  r.uniPathPDF,
                                  r.lightPathPDF,
                                  r.pixelIndex,
-                                 r.piPrev,
-                                 r.nPrev,
-                                 r.nsPrev,
+                                 r.prevIntrCtx,
                                  r.isSpecularBounce,
                                  r.anyNonSpecularBounces,
                                  r.etaScale,
@@ -201,7 +199,7 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
                  rayIndex, r.pixelIndex);
         Ray newRay = intr.SpawnRay(r.ray.d);
         params.mediumTransitionQueue->Push(MediumTransitionWorkItem{
-            newRay, r.lambda, r.beta, r.uniPathPDF, r.lightPathPDF, r.piPrev, r.nPrev, r.nsPrev,
+            newRay, r.lambda, r.beta, r.uniPathPDF, r.lightPathPDF, r.prevIntrCtx,
             r.isSpecularBounce, r.anyNonSpecularBounces, r.etaScale, r.pixelIndex});
         return;
     }
@@ -213,8 +211,7 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
         // TODO: intr.wo == -ray.d?
         params.hitAreaLightQueue->Push(HitAreaLightWorkItem{
             intr.areaLight, r.lambda, r.beta, r.uniPathPDF, r.lightPathPDF, intr.p(), intr.n,
-            intr.uv, intr.wo, r.piPrev, ray.d, r.nPrev, r.nsPrev,
-            (int)r.isSpecularBounce, r.pixelIndex});
+            intr.uv, intr.wo, r.prevIntrCtx, (int)r.isSpecularBounce, r.pixelIndex});
     }
 
     FloatTextureHandle displacement = material.GetDisplacement();

src/pbrt/gpu/pathintegrator.cpp

@@ -330,7 +330,9 @@ void GPUPathIntegrator::Render() {
          ++sampleIndex) {
         // Render image for sample _sampleIndex_
         LOG_VERBOSE("Starting to submit work for sample %d", sampleIndex);
-        for (int y0 = 0; y0 < resolution.y; y0 += scanlinesPerPass) {
+        Bounds2i pixelBounds = film.PixelBounds();
+        for (int y0 = pixelBounds.pMin.y; y0 < pixelBounds.pMax.y;
+             y0 += scanlinesPerPass) {
             // Generate camera rays for current scanline range
             RayQueue *cameraRayQueue = CurrentRayQueue(0);
             GPUDo(
@@ -344,11 +346,13 @@ void GPUPathIntegrator::Render() {
                 "Update camera ray stats",
                 PBRT_GPU_LAMBDA() { stats->cameraRays += cameraRayQueue->Size(); });
 
+            // Trace rays and estimate radiance up to maximum ray depth
             for (int depth = 0; true; ++depth) {
                 // Reset ray queues before tracing rays
-                RayQueue *resetQueue = NextRayQueue(depth);
+                RayQueue *nextQueue = NextRayQueue(depth);
                 GPUDo(
                     "Reset queues before tracing rays", PBRT_GPU_LAMBDA() {
+                        nextQueue->Reset();
                         // Reset queues before tracing next batch of rays
                         mediumTransitionQueue->Reset();
                         if (mediumSampleQueue)
@@ -356,8 +360,6 @@ void GPUPathIntegrator::Render() {
                         if (mediumScatterQueue)
                             mediumScatterQueue->Reset();
 
-                        resetQueue->Reset();
-
                         if (escapedRayQueue)
                             escapedRayQueue->Reset();
 
@@ -402,6 +404,7 @@ void GPUPathIntegrator::Render() {
                 if (haveSubsurface)
                     SampleSubsurface(depth);
             }
+
             UpdateFilm();
             // Copy updated film pixels to buffer for display
             if (!Options->displayServer.empty())
@@ -457,7 +460,7 @@ void GPUPathIntegrator::HandleEscapedRays(int depth) {
             if (depth == 0 || er.specularBounce) {
                 L = er.beta * Le / er.uniPathPDF.Average();
             } else {
-                LightSampleContext ctx(er.piPrev, er.nPrev, er.nsPrev);
+                LightSampleContext ctx = er.prevIntrCtx;
 
                 Float lightChoicePDF = lightSampler.PDF(ctx, envLight);
                 Float lightPDF = lightChoicePDF *
@@ -495,9 +498,9 @@ void GPUPathIntegrator::HandleRayFoundEmission(int depth) {
             if (depth == 0 || he.isSpecularBounce) {
                 L = he.beta * Le / he.uniPathPDF.Average();
             } else {
-                Vector3f wi = he.rayd;
+                Vector3f wi = -he.wo;
 
-                LightSampleContext ctx(he.piPrev, he.nPrev, he.nsPrev);
+                LightSampleContext ctx = he.prevIntrCtx;
 
                 Float lightChoicePDF = lightSampler.PDF(ctx, areaLight);
                 Float lightPDF = lightChoicePDF *

src/pbrt/gpu/pathintegrator.h

@@ -71,7 +71,6 @@ class GPUPathIntegrator {
     RayQueue *NextRayQueue(int depth) { return rayQueues[(depth + 1) & 1]; }
 
     // GPUPathIntegrator Member Variables
-    // Various properties of the scene
     bool initializeVisibleSurface;
     bool haveSubsurface;
     bool haveMedia;
@@ -91,8 +90,6 @@ class GPUPathIntegrator {
     };
     Stats *stats;
 
-    int scanlinesPerPass, maxQueueSize;
-
     FilterHandle filter;
     FilmHandle film;
     SamplerHandle sampler;
@@ -100,10 +97,12 @@ class GPUPathIntegrator {
     LightHandle envLight;
     LightSamplerHandle lightSampler;
 
-    RayQueue *rayQueues[2] = {nullptr, nullptr};
+    int scanlinesPerPass, maxQueueSize;
 
     SOA<PixelSampleState> pixelSampleState;
 
+    RayQueue *rayQueues[2];
+
     MediumTransitionQueue *mediumTransitionQueue = nullptr;
     MediumSampleQueue *mediumSampleQueue = nullptr;
     MediumScatterQueue *mediumScatterQueue = nullptr;

src/pbrt/gpu/samples.cpp

@@ -12,32 +12,41 @@
 namespace pbrt {
 
 // GPUPathIntegrator Sampler Methods
+void GPUPathIntegrator::GenerateRaySamples(int depth, int sampleIndex) {
+    auto generateSamples = [=](auto sampler) {
+        using Sampler = std::remove_reference_t<decltype(*sampler)>;
+        if constexpr (!std::is_same_v<Sampler, MLTSampler> &&
+                      !std::is_same_v<Sampler, DebugMLTSampler>)
+            GenerateRaySamples<Sampler>(depth, sampleIndex);
+    };
+    // Call the appropriate GenerateRaySamples specialization based on the
+    // Sampler's actual type.
+    sampler.DispatchCPU(generateSamples);
+}
+
 template <typename Sampler>
 void GPUPathIntegrator::GenerateRaySamples(int depth, int sampleIndex) {
     std::string desc = std::string("Generate ray samples - ") + Sampler::Name();
     RayQueue *rayQueue = CurrentRayQueue(depth);
-
     ForAllQueued(
         desc.c_str(), rayQueue, maxQueueSize,
         PBRT_GPU_LAMBDA(const RayWorkItem w, int index) {
-            // Figure out how many dimensions have been consumed so far: 5
-            // are used for the initial camera sample and then either 7 or
-            // 10 per ray, depending on whether there's subsurface
-            // scattering.
+            // Generate samples for ray segment at current sample index
+            // Find first sample dimension
             int dimension = 5 + 7 * depth;
             if (haveSubsurface)
                 dimension += 3 * depth;
 
-            // Initialize a Sampler
+            // Initialize _Sampler_ for pixel, sample index, and dimension
             Sampler pixelSampler = *sampler.Cast<Sampler>();
             Point2i pPixel = pixelSampleState.pPixel[w.pixelIndex];
             pixelSampler.StartPixelSample(pPixel, sampleIndex, dimension);
 
-            // Generate the samples for the ray and store them with it in
-            // the ray queue.
+            // Initialize _RaySamples_ structure with sample values
             RaySamples rs;
             rs.direct.u = pixelSampler.Get2D();
             rs.direct.uc = pixelSampler.Get1D();
+            // Initialize indirect and possibly subsurface samples in _rs_
             rs.indirect.u = pixelSampler.Get2D();
             rs.indirect.uc = pixelSampler.Get1D();
             rs.indirect.rr = pixelSampler.Get1D();
@@ -47,20 +56,9 @@ void GPUPathIntegrator::GenerateRaySamples(int depth, int sampleIndex) {
                 rs.subsurface.u = pixelSampler.Get2D();
             }
 
+            // Store _RaySamples_ in pixel sample state
             pixelSampleState.samples[w.pixelIndex] = rs;
         });
 }
 
-void GPUPathIntegrator::GenerateRaySamples(int depth, int sampleIndex) {
-    auto generateSamples = [=](auto sampler) {
-        using Sampler = std::remove_reference_t<decltype(*sampler)>;
-        if constexpr (!std::is_same_v<Sampler, MLTSampler> &&
-                      !std::is_same_v<Sampler, DebugMLTSampler>)
-            GenerateRaySamples<Sampler>(depth, sampleIndex);
-    };
-    // Call the appropriate GenerateRaySamples specialization based on the
-    // Sampler's actual type.
-    sampler.DispatchCPU(generateSamples);
-}
-
 }  // namespace pbrt

src/pbrt/gpu/subsurface.cpp

@@ -124,9 +124,10 @@ void GPUPathIntegrator::SampleSubsurface(int depth) {
                         // possible...
                         bool anyNonSpecularBounces = true;
 
-                        nextRayQueue->PushIndirect(ray, intr.pi, intr.n, intr.ns, beta,
-                                                   uniPathPDF, lightPathPDF, lambda,
-                                                   etaScale, bsdfSample->IsSpecular(),
+                        LightSampleContext ctx(intr.pi, intr.n, intr.ns);
+                        nextRayQueue->PushIndirect(ray, ctx, beta, uniPathPDF,
+                                                   lightPathPDF, lambda, etaScale,
+                                                   bsdfSample->IsSpecular(),
                                                    anyNonSpecularBounces, s.pixelIndex);
 
                         PBRT_DBG(

src/pbrt/gpu/surfscatter.cpp

@@ -173,7 +173,10 @@ void GPUPathIntegrator::EvaluateMaterialAndBSDF(TextureEvaluator texEval,
                             !bsdfSample->IsSpecular() || me.anyNonSpecularBounces;
 
                         // Spawn indriect ray.
-                        nextRayQueue->PushIndirect(ray, me.pi, me.n, ns, beta, uniPathPDF,
+                        LightSampleContext ctx(
+                            me.pi, me.n,
+                            ns);  // Note: slightly different than context below. Problem?
+                        nextRayQueue->PushIndirect(ray, ctx, beta, uniPathPDF,
                                                    lightPathPDF, lambda, etaScale,
                                                    bsdfSample->IsSpecular(),
                                                    anyNonSpecularBounces, me.pixelIndex);

src/pbrt/gpu/workitems.h

@@ -20,6 +20,7 @@ namespace pbrt {
 
 // RaySamples Definition
 struct RaySamples {
+    // RaySamples Public Members
     struct {
         Point2f u;
         Float uc;
@@ -28,7 +29,6 @@ struct RaySamples {
         Float uc, rr;
         Point2f u;
     } indirect;
-
     bool haveSubsurface;
     struct {
         Float uc;
@@ -99,11 +99,11 @@ struct SOA<RaySamples> {
 
 // PixelSampleState Definition
 struct PixelSampleState {
-    Float filterWeight;
+    // PixelSampleState Public Members
     Point2i pPixel;
+    Float filterWeight;
     SampledWavelengths lambda;
-    SampledSpectrum L;
-    SampledSpectrum cameraRayWeight;
+    SampledSpectrum L, cameraRayWeight;
     VisibleSurface visibleSurface;
     RaySamples samples;
 };
@@ -114,9 +114,7 @@ struct RayWorkItem {
     int pixelIndex;
     SampledWavelengths lambda;
     SampledSpectrum beta, uniPathPDF, lightPathPDF;
-    Point3fi piPrev;
-    Normal3f nPrev;
-    Normal3f nsPrev;
+    LightSampleContext prevIntrCtx;
     Float etaScale;
     int isSpecularBounce;
     int anyNonSpecularBounces;
@@ -128,8 +126,7 @@ struct EscapedRayWorkItem {
     SampledWavelengths lambda;
     Point3f rayo;
     Vector3f rayd;
-    Point3fi piPrev;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int specularBounce;
     int pixelIndex;
 };
@@ -146,9 +143,7 @@ struct HitAreaLightWorkItem {
     Normal3f n;
     Point2f uv;
     Vector3f wo;
-    Point3fi piPrev;
-    Vector3f rayd;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     int pixelIndex;
 };
@@ -256,8 +251,7 @@ struct MediumTransitionWorkItem {
     Ray ray;
     SampledWavelengths lambda;
     SampledSpectrum beta, uniPathPDF, lightPathPDF;
-    Point3fi piPrev;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     int anyNonSpecularBounces;
     Float etaScale;
@@ -277,9 +271,7 @@ struct MediumSampleWorkItem {
     SampledSpectrum uniPathPDF;
     SampledSpectrum lightPathPDF;
     int pixelIndex;
-    Point3fi piPrev;
-    Normal3f nPrev;
-    Normal3f nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     int anyNonSpecularBounces;
     Float etaScale;
@@ -335,9 +327,8 @@ class RayQueue : public WorkQueue<RayWorkItem> {
     }
 
     PBRT_CPU_GPU
-    int PushIndirect(const Ray &ray, const Point3fi &piPrev, const Normal3f &nPrev,
-                     const Normal3f &nsPrev, const SampledSpectrum &beta,
-                     const SampledSpectrum &uniPathPDF,
+    int PushIndirect(const Ray &ray, const LightSampleContext &prevIntrCtx,
+                     const SampledSpectrum &beta, const SampledSpectrum &uniPathPDF,
                      const SampledSpectrum &lightPathPDF,
                      const SampledWavelengths &lambda, Float etaScale,
                      bool isSpecularBounce, bool anyNonSpecularBounces, int pixelIndex) {
@@ -345,9 +336,7 @@ class RayQueue : public WorkQueue<RayWorkItem> {
         DCHECK(!ray.HasNaN());
         this->ray[index] = ray;
         this->pixelIndex[index] = pixelIndex;
-        this->piPrev[index] = piPrev;
-        this->nPrev[index] = nPrev;
-        this->nsPrev[index] = nsPrev;
+        this->prevIntrCtx[index] = prevIntrCtx;
         this->beta[index] = beta;
         this->uniPathPDF[index] = uniPathPDF;
         this->lightPathPDF[index] = lightPathPDF;
@@ -434,7 +423,7 @@ class MediumSampleQueue : public WorkQueue<MediumSampleWorkItem> {
     PBRT_CPU_GPU
     int Push(Ray ray, Float tMax, SampledWavelengths lambda, SampledSpectrum beta,
              SampledSpectrum uniPathPDF, SampledSpectrum lightPathPDF, int pixelIndex,
-             Point3fi piPrev, Normal3f nPrev, Normal3f nsPrev, int isSpecularBounce,
+             LightSampleContext prevIntrCtx, int isSpecularBounce,
              int anyNonSpecularBounces, Float etaScale) {
         int index = AllocateEntry();
         this->ray[index] = ray;
@@ -444,9 +433,7 @@ class MediumSampleQueue : public WorkQueue<MediumSampleWorkItem> {
         this->uniPathPDF[index] = uniPathPDF;
         this->lightPathPDF[index] = lightPathPDF;
         this->pixelIndex[index] = pixelIndex;
-        this->piPrev[index] = piPrev;
-        this->nPrev[index] = nPrev;
-        this->nsPrev[index] = nsPrev;
+        this->prevIntrCtx[index] = prevIntrCtx;
         this->isSpecularBounce[index] = isSpecularBounce;
         this->anyNonSpecularBounces[index] = anyNonSpecularBounces;
         this->etaScale[index] = etaScale;

src/pbrt/gpu/workitems.soa

@@ -11,6 +11,7 @@ flat MediumHandle;
 flat int;
 
 soa BSDF;
+soa LightSampleContext;
 soa MediumInterface;
 soa Normal3f;
 soa Point2f;
@@ -41,9 +42,7 @@ soa RayWorkItem {
     int pixelIndex;
     SampledWavelengths lambda;
     SampledSpectrum beta, uniPathPDF, lightPathPDF;
-    Point3fi piPrev;
-    Normal3f nPrev;
-    Normal3f nsPrev;
+    LightSampleContext prevIntrCtx;
     Float etaScale;
     int isSpecularBounce;
     int anyNonSpecularBounces;
@@ -54,8 +53,7 @@ soa EscapedRayWorkItem {
     SampledWavelengths lambda;
     Point3f rayo;
     Vector3f rayd;
-    Point3fi piPrev;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int specularBounce;
     int pixelIndex;
 };
@@ -68,9 +66,7 @@ soa HitAreaLightWorkItem {
     Normal3f n;
     Point2f uv;
     Vector3f wo;
-    Point3fi piPrev;
-    Vector3f rayd;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     int pixelIndex;
 };
@@ -118,8 +114,7 @@ soa MediumTransitionWorkItem {
     Ray ray;
     SampledWavelengths lambda;
     SampledSpectrum beta, uniPathPDF, lightPathPDF;
-    Point3fi piPrev;
-    Normal3f nPrev, nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     int anyNonSpecularBounces;
     Float etaScale;
@@ -139,9 +134,7 @@ soa MediumSampleWorkItem {
     Normal3f n;
     Vector3f wo;
     Point2f uv;
-    Point3fi piPrev;
-    Normal3f nPrev;
-    Normal3f nsPrev;
+    LightSampleContext prevIntrCtx;
     int isSpecularBounce;
     MaterialHandle material;
     Normal3f ns;

src/pbrt/pbrt.soa

@@ -64,3 +64,8 @@ soa TabulatedBSSRDF {
     const BSSRDFTable *table;
     SampledSpectrum sigma_t, rho;
 };
+
+soa LightSampleContext {
+    Point3fi pi;
+    Normal3f n, ns;
+};