GPU: copy SOA<PixelSampleState> in RayIntersectParameters

This fixes the debug build on Windows on the GPU. (Release crashes with OptiX complaining about malformed PTX.)

The issue is essentially the same as why *this is copied in GPU lambdas
rather than being passed as a pointer; we are accessing the GPUPathIntegrator
in read-only fashion from the CPU during rendering and with unified
memory on Windows, it isn't allowed to concurrently access it on the GPU.

This also fits with the data point that 82ace326 is when things first started crashing.

Issues #41, #48, #72, #89 , and #96.

src/pbrt/gpu/accel.cpp

@@ -1189,7 +1189,7 @@ void GPUAccel::IntersectShadow(int maxRays, ShadowRayQueue *shadowRayQueue,
         RayIntersectParameters params;
         params.traversable = rootTraversable;
         params.shadowRayQueue = shadowRayQueue;
-        params.pixelSampleState = pixelSampleState;
+        params.pixelSampleState = *pixelSampleState;
 
         ParamBufferState &pbs = getParamBuffer(params);
 
@@ -1227,7 +1227,7 @@ void GPUAccel::IntersectShadowTr(int maxRays, ShadowRayQueue *shadowRayQueue,
         RayIntersectParameters params;
         params.traversable = rootTraversable;
         params.shadowRayQueue = shadowRayQueue;
-        params.pixelSampleState = pixelSampleState;
+        params.pixelSampleState = *pixelSampleState;
 
         ParamBufferState &pbs = getParamBuffer(params);
 

src/pbrt/gpu/optix.cu

@@ -356,8 +356,8 @@ extern "C" __global__ void __raygen__shadow() {
                  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;
+        SampledSpectrum Lpixel = params.pixelSampleState.L[sr.pixelIndex];
+        params.pixelSampleState.L[sr.pixelIndex] = Lpixel + Ld;
     } else {
         PBRT_DBG("Shadow ray was occluded\n");
     }
@@ -506,8 +506,8 @@ extern "C" __global__ void __raygen__shadow_Tr() {
         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]);
 
-        SampledSpectrum Lpixel = params.pixelSampleState->L[sr.pixelIndex];
-        params.pixelSampleState->L[sr.pixelIndex] = Lpixel + Ld;
+        SampledSpectrum Lpixel = params.pixelSampleState.L[sr.pixelIndex];
+        params.pixelSampleState.L[sr.pixelIndex] = Lpixel + Ld;
     }
 }
 

src/pbrt/gpu/optix.h

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