From fdcaa3a8e2f1d112de2cf40032969a188cdf301a Mon Sep 17 00:00:00 2001 From: Matt Pharr Date: Sat, 17 Apr 2021 11:45:31 -0700 Subject: [PATCH] GPU: copy SOA 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 82ace32 is when things first started crashing. Issues #41, #48, #72, #89 , and #96. --- src/pbrt/gpu/accel.cpp | 4 ++-- src/pbrt/gpu/optix.cu | 8 ++++---- src/pbrt/gpu/optix.h | 2 +- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/src/pbrt/gpu/accel.cpp b/src/pbrt/gpu/accel.cpp index a865af7..fbe8e13 100644 --- a/src/pbrt/gpu/accel.cpp +++ b/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); diff --git a/src/pbrt/gpu/optix.cu b/src/pbrt/gpu/optix.cu index f176717..031b9ed 100644 --- a/src/pbrt/gpu/optix.cu +++ b/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; } } diff --git a/src/pbrt/gpu/optix.h b/src/pbrt/gpu/optix.h index f4b8320..7d74b3c 100644 --- a/src/pbrt/gpu/optix.h +++ b/src/pbrt/gpu/optix.h @@ -61,7 +61,7 @@ struct RayIntersectParameters { // shadow rays ShadowRayQueue *shadowRayQueue; - SOA *pixelSampleState; + SOA pixelSampleState; // Subsurface scattering... SubsurfaceScatterQueue *subsurfaceScatterQueue; -- GitLab