From 225ea05a765a1a9c6a7ada31ec402b8a77fc745a Mon Sep 17 00:00:00 2001 From: Matt Pharr Date: Tue, 17 Nov 2020 09:46:01 -0800 Subject: [PATCH] Update from book source. No functionality changes. Cleaned up some of the GPU *WorkItem definitions. --- src/pbrt/gpu/camera.cpp | 26 +++++++-------------- src/pbrt/gpu/media.cpp | 22 ++++++++---------- src/pbrt/gpu/optix.cu | 17 ++++++-------- src/pbrt/gpu/pathintegrator.cpp | 17 ++++++++------ src/pbrt/gpu/pathintegrator.h | 7 +++--- src/pbrt/gpu/samples.cpp | 38 +++++++++++++++--------------- src/pbrt/gpu/subsurface.cpp | 7 +++--- src/pbrt/gpu/surfscatter.cpp | 5 +++- src/pbrt/gpu/workitems.h | 41 +++++++++++---------------------- src/pbrt/gpu/workitems.soa | 19 +++++---------- src/pbrt/pbrt.soa | 5 ++++ 11 files changed, 89 insertions(+), 115 deletions(-) diff --git a/src/pbrt/gpu/camera.cpp b/src/pbrt/gpu/camera.cpp index 9d7d687..056397f 100644 --- a/src/pbrt/gpu/camera.cpp +++ b/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; if constexpr (!std::is_same_v && !std::is_same_v) GenerateCameraRays(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; diff --git a/src/pbrt/gpu/media.cpp b/src/pbrt/gpu/media.cpp index df40095..273e8d9 100644 --- a/src/pbrt/gpu/media.cpp +++ b/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", diff --git a/src/pbrt/gpu/optix.cu b/src/pbrt/gpu/optix.cu index b9604d3..bdeaa34 100644 --- a/src/pbrt/gpu/optix.cu +++ b/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(); diff --git a/src/pbrt/gpu/pathintegrator.cpp b/src/pbrt/gpu/pathintegrator.cpp index 9272bf7..7701ab2 100644 --- a/src/pbrt/gpu/pathintegrator.cpp +++ b/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 * diff --git a/src/pbrt/gpu/pathintegrator.h b/src/pbrt/gpu/pathintegrator.h index 258a576..7f81524 100644 --- a/src/pbrt/gpu/pathintegrator.h +++ b/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; + RayQueue *rayQueues[2]; + MediumTransitionQueue *mediumTransitionQueue = nullptr; MediumSampleQueue *mediumSampleQueue = nullptr; MediumScatterQueue *mediumScatterQueue = nullptr; diff --git a/src/pbrt/gpu/samples.cpp b/src/pbrt/gpu/samples.cpp index ea49138..12b0234 100644 --- a/src/pbrt/gpu/samples.cpp +++ b/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; + if constexpr (!std::is_same_v && + !std::is_same_v) + GenerateRaySamples(depth, sampleIndex); + }; + // Call the appropriate GenerateRaySamples specialization based on the + // Sampler's actual type. + sampler.DispatchCPU(generateSamples); +} + template 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(); 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; - if constexpr (!std::is_same_v && - !std::is_same_v) - GenerateRaySamples(depth, sampleIndex); - }; - // Call the appropriate GenerateRaySamples specialization based on the - // Sampler's actual type. - sampler.DispatchCPU(generateSamples); -} - } // namespace pbrt diff --git a/src/pbrt/gpu/subsurface.cpp b/src/pbrt/gpu/subsurface.cpp index dedb9f3..b0d48e7 100644 --- a/src/pbrt/gpu/subsurface.cpp +++ b/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( diff --git a/src/pbrt/gpu/surfscatter.cpp b/src/pbrt/gpu/surfscatter.cpp index 4749e56..9d5157b 100644 --- a/src/pbrt/gpu/surfscatter.cpp +++ b/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); diff --git a/src/pbrt/gpu/workitems.h b/src/pbrt/gpu/workitems.h index ad7c011..c9c7371 100644 --- a/src/pbrt/gpu/workitems.h +++ b/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 { // 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 { } 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 { 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 { 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 { 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; diff --git a/src/pbrt/gpu/workitems.soa b/src/pbrt/gpu/workitems.soa index d3d5a64..e20a1d7 100644 --- a/src/pbrt/gpu/workitems.soa +++ b/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; diff --git a/src/pbrt/pbrt.soa b/src/pbrt/pbrt.soa index 0bf24c4..9b03514 100644 --- a/src/pbrt/pbrt.soa +++ b/src/pbrt/pbrt.soa @@ -64,3 +64,8 @@ soa TabulatedBSSRDF { const BSSRDFTable *table; SampledSpectrum sigma_t, rho; }; + +soa LightSampleContext { + Point3fi pi; + Normal3f n, ns; +}; -- GitLab