Fix GPU build

src/pbrt/gpu/media.cpp

@@ -168,8 +168,8 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                     PBRT_DBG("Adding ray to escapedRayQueue pixel index %d depth %d\n",
                              w.pixelIndex, depth);
                     escapedRayQueue->Push(EscapedRayWorkItem{
-                        T_hat, uniPathPDF, lightPathPDF, lambda, ray.o, ray.d,
-                        w.prevIntrCtx, (int)w.isSpecularBounce, w.pixelIndex});
+                        ray.o, ray.d, lambda, w.pixelIndex, (int)w.isSpecularBounce,
+                        T_hat, uniPathPDF, lightPathPDF, w.prevIntrCtx});
                 }
             }
 
@@ -191,8 +191,9 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                     "depth %d\n",
                     w.pixelIndex, depth);
                 hitAreaLightQueue->Push(HitAreaLightWorkItem{
-                    w.areaLight, lambda, T_hat, uniPathPDF, lightPathPDF, Point3f(w.pi),
-                    w.n, w.uv, -ray.d, w.prevIntrCtx, w.isSpecularBounce, w.pixelIndex});
+                    w.areaLight, Point3f(w.pi), w.n, w.uv, -ray.d, lambda, T_hat,
+                    uniPathPDF, lightPathPDF, w.prevIntrCtx, w.isSpecularBounce,
+                    w.pixelIndex});
             }
 
             FloatTextureHandle displacement = material.GetDisplacement();
@@ -210,8 +211,8 @@ void GPUPathIntegrator::SampleMediumInteraction(int depth) {
                 using Material = typename std::remove_reference_t<decltype(*ptr)>;
                 q->Push<MaterialEvalWorkItem<Material>>(MaterialEvalWorkItem<Material>{
                     ptr, w.pi, w.n, w.ns, w.dpdus, w.dpdvs, w.dndus, w.dndvs, w.uv,
-                    lambda, w.anyNonSpecularBounces, T_hat, uniPathPDF, -ray.d, ray.time,
-                    w.etaScale, w.mediumInterface, w.pixelIndex});
+                    lambda, w.anyNonSpecularBounces, -ray.d, w.pixelIndex, T_hat,
+                    uniPathPDF, w.mediumInterface, w.etaScale, ray.time});
             };
             material.Dispatch(enqueue);
         });

src/pbrt/gpu/optix.cu

@@ -124,8 +124,8 @@ extern "C" __global__ void __raygen__findClosest() {
             PBRT_DBG("Adding ray to escapedRayQueue ray index %d pixel index %d\n", rayIndex,
                      r.pixelIndex);
             params.escapedRayQueue->Push(EscapedRayWorkItem{
-                r.T_hat, r.uniPathPDF, r.lightPathPDF, r.lambda, ray.o, ray.d, r.prevIntrCtx,
-                (int)r.isSpecularBounce, r.pixelIndex});
+                ray.o, ray.d, r.lambda, r.pixelIndex, (int)r.isSpecularBounce,
+                r.T_hat, r.uniPathPDF, r.lightPathPDF, r.prevIntrCtx});
         }
     }
 }
@@ -198,7 +198,7 @@ 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.nextRayQueue->PushIndirect(
+        params.nextRayQueue->PushIndirectRay(
             newRay, r.prevIntrCtx, r.T_hat, r.uniPathPDF, r.lightPathPDF, r.lambda,
             r.etaScale, r.isSpecularBounce, r.anyNonSpecularBounces, r.pixelIndex);
         return;
@@ -210,8 +210,9 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
         Ray ray = r.ray;
         // TODO: intr.wo == -ray.d?
         params.hitAreaLightQueue->Push(HitAreaLightWorkItem{
-            intr.areaLight, r.lambda, r.T_hat, r.uniPathPDF, r.lightPathPDF, intr.p(), intr.n,
-            intr.uv, intr.wo, r.prevIntrCtx, (int)r.isSpecularBounce, r.pixelIndex});
+            intr.areaLight, intr.p(), intr.n, intr.uv, intr.wo, r.lambda,
+            r.T_hat, r.uniPathPDF, r.lightPathPDF, r.prevIntrCtx,
+            (int)r.isSpecularBounce, r.pixelIndex});
     }
 
     FloatTextureHandle displacement = material.GetDisplacement();
@@ -229,9 +230,9 @@ static __forceinline__ __device__ void ProcessClosestIntersection(
         q->Push(MaterialEvalWorkItem<Material>{
             ptr, intr.pi, intr.n, intr.shading.n,
             intr.shading.dpdu, intr.shading.dpdv, intr.shading.dndu, intr.shading.dndv,
-            intr.uv, r.lambda, r.anyNonSpecularBounces,
-            r.T_hat, r.uniPathPDF, intr.wo, intr.time, r.etaScale,
-            getPayload<ClosestHitContext>()->mediumInterface, r.pixelIndex});
+                intr.uv, r.lambda, r.anyNonSpecularBounces, intr.wo, r.pixelIndex,
+                r.T_hat, r.uniPathPDF, getPayload<ClosestHitContext>()->mediumInterface,
+                r.etaScale, intr.time});
     };
     material.Dispatch(enqueue);
 

src/pbrt/gpu/pathintegrator.cpp

@@ -476,7 +476,7 @@ void GPUPathIntegrator::HandleRayFoundEmission(int depth) {
         "Handle emitters hit by indirect rays", hitAreaLightQueue, maxQueueSize,
         PBRT_GPU_LAMBDA(const HitAreaLightWorkItem w, int index) {
             // Find emitted radiance from surface that ray hit
-            SampledSpectrum Le = areaLight.L(w.p, w.n, w.uv, w.wo, w.lambda);
+            SampledSpectrum Le = w.areaLight.L(w.p, w.n, w.uv, w.wo, w.lambda);
             if (!Le)
                 return;
             PBRT_DBG("Got Le %f %f %f %f from hit area light at depth %d\n", Le[0], Le[1],

src/pbrt/gpu/workitems.h

@@ -291,7 +291,6 @@ struct MaterialEvalWorkItem {
     int pixelIndex;
     SampledSpectrum T_hat, uniPathPDF;
     MediumInterface mediumInterface;
-    Float time;
     Float etaScale;
     Float time;
 };