diff --git a/CMakeLists.txt b/CMakeLists.txt index fa9b198e..58f4eb1b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -15,7 +15,7 @@ SET(KRR_RENDER_BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/src) OPTION (KRR_BUILD_STARLIGHT "Build some additional code implementation of fancy algorithms." - OFF + ON ) OPTION (KRR_BUILD_EXAMPLES @@ -25,7 +25,7 @@ OPTION (KRR_BUILD_EXAMPLES OPTION (KRR_RENDER_SPECTRAL "Build the spectral variant of the renderer (WIP)." - OFF + ON ) ############################################################################### diff --git a/common/build/source.cmake b/common/build/source.cmake index c93195ed..4e83f969 100644 --- a/common/build/source.cmake +++ b/common/build/source.cmake @@ -140,14 +140,12 @@ SET(KRR_SOURCE ${KRR_SOURCE} ${KRR_RENDER_SOURCE_DIR}/misc/render/ppg/integrator.cpp ${KRR_RENDER_SOURCE_DIR}/misc/render/ppg/treemanip.cpp - ${KRR_RENDER_SOURCE_DIR}/misc/render/ppg/medium.cpp ${KRR_RENDER_SOURCE_DIR}/misc/render/zero_guiding/integrator.cpp ${KRR_RENDER_SOURCE_DIR}/misc/render/zero_guiding/medium.cpp ) SET_SOURCE_FILES_PROPERTIES ( ${KRR_RENDER_SOURCE_DIR}/misc/render/ppg/integrator.cpp - ${KRR_RENDER_SOURCE_DIR}/misc/render/ppg/medium.cpp ${KRR_RENDER_SOURCE_DIR}/misc/render/zero_guiding/integrator.cpp ${KRR_RENDER_SOURCE_DIR}/misc/render/zero_guiding/medium.cpp PROPERTIES LANGUAGE CUDA diff --git a/src/misc/render/ppg/device.cu b/src/misc/render/ppg/device.cu index 110216d6..6dd96ade 100644 --- a/src/misc/render/ppg/device.cu +++ b/src/misc/render/ppg/device.cu @@ -30,23 +30,25 @@ KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, traceRay(traversable, ray, tMax, rayType, flags, u0, u1); } +KRR_DEVICE_FUNCTION int getRayId() { return optixGetLaunchIndex().x; } + KRR_DEVICE_FUNCTION RayWorkItem getRayWorkItem() { - int rayIndex(optixGetLaunchIndex().x); - DCHECK_LT(rayIndex, launchParams.currentRayQueue->size()); - return (*launchParams.currentRayQueue)[rayIndex]; + DCHECK_LT(getRayId(), launchParams.currentRayQueue->size()); + return (*launchParams.currentRayQueue)[getRayId()]; } KRR_DEVICE_FUNCTION ShadowRayWorkItem getShadowRayWorkItem() { - int rayIndex(optixGetLaunchIndex().x); - DCHECK_LT(rayIndex, launchParams.shadowRayQueue->size()); - return (*launchParams.shadowRayQueue)[rayIndex]; + DCHECK_LT(getRayId(), launchParams.shadowRayQueue->size()); + return (*launchParams.shadowRayQueue)[getRayId()]; } extern "C" __global__ void KRR_RT_CH(Closest)() { HitInfo hitInfo = getHitInfo(); SurfaceInteraction& intr = *getPRD(); RayWorkItem r = getRayWorkItem(); - prepareSurfaceInteraction(intr, hitInfo); + int pixelId = launchParams.currentRayQueue->pixelId[getRayId()]; + SampledWavelengths lambda = launchParams.pixelState->lambda[pixelId]; + prepareSurfaceInteraction(intr, hitInfo, launchParams.colorSpace, lambda); if (launchParams.mediumSampleQueue && r.ray.medium) { launchParams.mediumSampleQueue->push(r, intr, optixGetRayTmax()); return; @@ -72,8 +74,7 @@ extern "C" __global__ void KRR_RT_MS(Closest)() { } extern "C" __global__ void KRR_RT_RG(Closest)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.currentRayQueue->size()) return; + if (getRayId() >= launchParams.currentRayQueue->size()) return; RayWorkItem r = getRayWorkItem(); SurfaceInteraction intr = {}; traceRay(launchParams.traversable, r.ray, M_FLOAT_INF, @@ -81,22 +82,20 @@ extern "C" __global__ void KRR_RT_RG(Closest)() { } extern "C" __global__ void KRR_RT_AH(Shadow)() { - if (alphaKilled()) - optixIgnoreIntersection(); + if (alphaKilled()) optixIgnoreIntersection(); } extern "C" __global__ void KRR_RT_MS(Shadow)() { optixSetPayload_0(1); } extern "C" __global__ void KRR_RT_RG(Shadow)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); uint32_t visible{0}; traceRay(launchParams.traversable, r.ray, r.tMax, 1, OptixRayFlags( OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT | OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT), visible); if (visible) { - Color contrib = r.Ld / (r.pl + r.pu).mean(); + SampledSpectrum contrib = r.Ld / (r.pl + r.pu).mean(); launchParams.pixelState->addRadiance(r.pixelId, contrib); if (launchParams.enableTraining) launchParams.guidedState->recordRadiance(r.pixelId, contrib); @@ -105,8 +104,10 @@ extern "C" __global__ void KRR_RT_RG(Shadow)() { extern "C" __global__ void KRR_RT_CH(ShadowTr)() { HitInfo hitInfo = getHitInfo(); + int pixelId = launchParams.shadowRayQueue->pixelId[getRayId()]; + SampledWavelengths lambda = launchParams.pixelState->lambda[pixelId]; SurfaceInteraction &intr = *getPRD(); - prepareSurfaceInteraction(intr, hitInfo); + prepareSurfaceInteraction(intr, hitInfo, launchParams.colorSpace, lambda); } extern "C" __global__ void KRR_RT_AH(ShadowTr)() { @@ -116,8 +117,7 @@ extern "C" __global__ void KRR_RT_AH(ShadowTr)() { extern "C" __global__ void KRR_RT_MS(ShadowTr)() { optixSetPayload_2(1); } extern "C" __global__ void KRR_RT_RG(ShadowTr)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); SurfaceInteraction intr = {}; uint u0, u1; diff --git a/src/misc/render/ppg/guideditem.h b/src/misc/render/ppg/guideditem.h index 134a105b..5b6331cc 100644 --- a/src/misc/render/ppg/guideditem.h +++ b/src/misc/render/ppg/guideditem.h @@ -17,14 +17,14 @@ struct Vertex { DTreeWrapper* dTree; Vector3f dTreeVoxelSize; Ray ray; - Color3f throughput; - Color3f bsdfVal; - Color3f radiance; + SampledSpectrum throughput; + SampledSpectrum bsdfVal; + SampledSpectrum radiance; float wiPdf, bsdfPdf, dTreePdf; float wiMisWeight; // @addition VAPG bool isDelta; - KRR_DEVICE void record(const Color3f& r) { + KRR_DEVICE void record(const SampledSpectrum &r) { radiance += r; } @@ -32,18 +32,17 @@ struct Vertex { EDirectionalFilter directionalFilter, EBsdfSamplingFractionLoss bsdfSamplingFractionLoss, Sampler &sampler, EDistribution distribution = EDistribution::ERadiance, - const Color &pixelEstimate = {}) { - if (wiPdf <= 0 || isDelta) { - return; - } - Color3f localRadiance = 0; + const SampledSpectrum &pixelEstimate = {}) { + if (wiPdf <= 0 || isDelta) return; + + SampledSpectrum localRadiance(0); if (throughput[0] * wiPdf > 1e-4f) localRadiance[0] = radiance[0] / throughput[0]; if (throughput[1] * wiPdf > 1e-4f) localRadiance[1] = radiance[1] / throughput[1]; if (throughput[2] * wiPdf > 1e-4f) localRadiance[2] = radiance[2] / throughput[2]; - Color3f product = localRadiance * bsdfVal; + SampledSpectrum product = localRadiance * bsdfVal; /* @modified: VAPG */ float value = localRadiance.mean(); @@ -66,11 +65,8 @@ struct Vertex { break; } - DTreeRecord rec{ ray.dir, - value, product.mean(), - wiPdf, bsdfPdf, dTreePdf, - statisticalWeight, - isDelta }; + DTreeRecord rec{ray.dir, value, product.mean(), wiPdf, + bsdfPdf, dTreePdf, statisticalWeight, isDelta}; switch (spatialFilter) { case ESpatialFilter::ENearest: @@ -123,17 +119,14 @@ class GuidedPathStateBuffer : public SOA { GuidedPathStateBuffer() = default; GuidedPathStateBuffer(int n, Allocator alloc) : SOA(n, alloc) {} - KRR_DEVICE void incrementDepth(int pixelId, - Ray& ray, - DTreeWrapper* dTree, - Vector3f dTreeVoxelSize, - Color3f thp, - Color3f bsdfVal, - float wiPdf, float bsdfPdf, float dTreePdf, - bool isDelta = false) { + KRR_DEVICE void incrementDepth(int pixelId, Ray &ray, DTreeWrapper *dTree, + Vector3f dTreeVoxelSize, SampledSpectrum thp, + SampledSpectrum bsdfVal, float wiPdf, float bsdfPdf, + float dTreePdf, bool isDelta = false) { int depth = n_vertices[pixelId]; if (depth >= MAX_TRAIN_DEPTH) return; - vertices[depth].radiance[pixelId] = Color3f(0); /* Always remember to clear the radiance from last frame, as I forgotten... */ + /* Always remember to clear the radiance from last frame, as I forgotten... */ + vertices[depth].radiance[pixelId] = SampledSpectrum(0); vertices[depth].ray[pixelId] = ray; vertices[depth].dTree[pixelId] = dTree; vertices[depth].dTreeVoxelSize[pixelId] = dTreeVoxelSize; @@ -148,10 +141,10 @@ class GuidedPathStateBuffer : public SOA { n_vertices[pixelId] = 1 + depth; } - KRR_DEVICE void recordRadiance(int pixelId, Color3f L) { + KRR_DEVICE void recordRadiance(int pixelId, SampledSpectrum L) { int cur_depth = n_vertices[pixelId]; for (int i = 0; i < cur_depth; i++) { - Color3f prevRadiance = vertices[i].radiance[pixelId]; + SampledSpectrum prevRadiance = vertices[i].radiance[pixelId]; vertices[i].radiance[pixelId] = prevRadiance + L; } } @@ -161,7 +154,7 @@ class GuidedPathStateBuffer : public SOA { STree* sdTree, float statisticalWeight, ESpatialFilter spatialFilter, EDirectionalFilter directionalFilter, EBsdfSamplingFractionLoss bsdfSamplingFractionLoss, Sampler& sampler, - EDistribution distribution = EDistribution::ERadiance, const Color3f &pixelEstimate = {}) { + EDistribution distribution = EDistribution::ERadiance, const SampledSpectrum &pixelEstimate = {}) { for (int i = 0; i < n_vertices[pixelId]; i++) { Vertex v = vertices[i][pixelId]; v.commit(sdTree, statisticalWeight, diff --git a/src/misc/render/ppg/guideditem.soa b/src/misc/render/ppg/guideditem.soa index 9a8b60c2..33041537 100644 --- a/src/misc/render/ppg/guideditem.soa +++ b/src/misc/render/ppg/guideditem.soa @@ -12,8 +12,7 @@ flat HaltonSampler; flat DTreeWrapper; flat Vector3f; -flat Color3f; -flat Color; +flat SampledSpectrum; struct Ray; struct LightSampleContext; @@ -26,9 +25,9 @@ struct Vertex { DTreeWrapper* dTree; Vector3f dTreeVoxelSize; Ray ray; - Color3f throughput; - Color3f bsdfVal; - Color3f radiance; + SampledSpectrum throughput; + SampledSpectrum bsdfVal; + SampledSpectrum radiance; float wiPdf, bsdfPdf, dTreePdf; float wiMisWeight; bool isDelta; diff --git a/src/misc/render/ppg/integrator.cpp b/src/misc/render/ppg/integrator.cpp index a4243c2f..771c8bfe 100644 --- a/src/misc/render/ppg/integrator.cpp +++ b/src/misc/render/ppg/integrator.cpp @@ -78,12 +78,14 @@ void PPGPathTracer::traceClosest(int depth) { static LaunchParamsPPG params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); + params.colorSpace = KRR_DEFAULT_COLORSPACE; params.currentRayQueue = currentRayQueue(depth); params.missRayQueue = missRayQueue; params.hitLightRayQueue = hitLightRayQueue; params.scatterRayQueue = scatterRayQueue; params.nextRayQueue = nextRayQueue(depth); params.mediumSampleQueue = enableMedium ? mediumSampleQueue : nullptr; + params.pixelState = pixelState; backend->launch(params, "Closest", maxQueueSize, 1, 1); } @@ -92,6 +94,7 @@ void PPGPathTracer::traceShadow() { static LaunchParamsPPG params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); + params.colorSpace = KRR_DEFAULT_COLORSPACE; params.shadowRayQueue = shadowRayQueue; params.pixelState = pixelState; params.guidedState = guidedPathState; @@ -104,7 +107,8 @@ void PPGPathTracer::handleHit() { PROFILE("Process intersected rays"); ForAllQueued(hitLightRayQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const HitLightWorkItem & w){ - Color Le = w.light.L(w.p, w.n, w.uv, w.wo) * w.thp; + SampledSpectrum Le = + w.light.L(w.p, w.n, w.uv, w.wo, pixelState->lambda[w.pixelId]) * w.thp; if (enableNEE && w.depth && !(w.bsdfType & BSDF_SPECULAR)) { Interaction intr(w.p, w.wo, w.n, w.uv); float lightPdf = w.light.pdfLi(intr, w.ctx) * lightSampler.pdf(w.light); @@ -119,13 +123,14 @@ void PPGPathTracer::handleMiss() { PROFILE("Process escaped rays"); const rt::SceneData& sceneData = mScene->mSceneRT->getSceneData(); ForAllQueued(missRayQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const MissRayWorkItem & w) { - Color L = {}; + SampledSpectrum L = {}; + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; Interaction intr(w.ray.origin); for (const rt::InfiniteLight &light : sceneData.infiniteLights) { if (enableNEE && w.depth && !(w.bsdfType & BSDF_SPECULAR)) { float lightPdf = light.pdfLi(intr, w.ctx) * lightSampler.pdf(&light); - L += light.Li(w.ray.dir) / (w.pu + w.pl * lightPdf).mean(); - } else L += light.Li(w.ray.dir) / w.pu.mean(); + L += light.Li(w.ray.dir, lambda) / (w.pu + w.pl * lightPdf).mean(); + } else L += light.Li(w.ray.dir, lambda) / w.pu.mean(); } pixelState->addRadiance(w.pixelId, w.thp * L); guidedPathState->recordRadiance(w.pixelId, w.thp * L); // ppg @@ -143,6 +148,7 @@ void PPGPathTracer::handleIntersections() { const SurfaceInteraction& intr = w.intr; BSDFType bsdfType = intr.getBsdfType(); Vector3f woLocal = intr.toLocal(intr.wo); + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; /* Statistics for mixed bsdf-guided sampling */ float bsdfPdf, dTreePdf; @@ -151,13 +157,13 @@ void PPGPathTracer::handleIntersections() { if (enableNEE && (bsdfType & BSDF_SMOOTH)) { SampledLight sampledLight = lightSampler.sample(sampler.get1D()); Light light = sampledLight.light; - LightSample ls = light.sampleLi(sampler.get2D(), { intr.p, intr.n }); + LightSample ls = light.sampleLi(sampler.get2D(), {intr.p, intr.n}, lambda); Ray shadowRay = intr.spawnRayTo(ls.intr); Vector3f wiWorld = normalize(shadowRay.dir); Vector3f wiLocal = intr.toLocal(wiWorld); float lightPdf = sampledLight.pdf * ls.pdf; - Color bsdfVal = BxDF::f(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); + SampledSpectrum bsdfVal = BxDF::f(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); float bsdfPdf = light.isDeltaLight() ? 0 : BxDF::pdf(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); if (lightPdf > 0 && bsdfVal.any()) { ShadowRayWorkItem sw = {}; @@ -219,6 +225,7 @@ void PPGPathTracer::generateScatterRays() { void PPGPathTracer::render(RenderContext *context) { if (!mScene || !maxQueueSize) return; PROFILE("PPG Path Tracer"); + CHECK_LOG(samplesPerPixel == 1, "Only 1spp/frame is supported for PPG!"); for (int sampleId = 0; sampleId < samplesPerPixel; sampleId++) { // [STEP#1] generate camera / primary rays GPUCall(KRR_DEVICE_LAMBDA() { currentRayQueue(0)->reset(); }); @@ -239,11 +246,6 @@ void PPGPathTracer::render(RenderContext *context) { }); // [STEP#2.1] find closest intersections, filling in scatterRayQueue and hitLightQueue traceClosest(depth); - // [STEP#2.2] sample medium interaction, and optionally sample in-volume scattering events - if (enableMedium) { - sampleMediumInteraction(depth); - sampleMediumScattering(depth); - } // [STEP#2.2] handle hit and missed rays, contribute to pixels handleHit(); handleMiss(); @@ -259,11 +261,16 @@ void PPGPathTracer::render(RenderContext *context) { // write results of the current frame... CudaRenderTarget frameBuffer = context->getColorTexture()->getCudaRenderTarget(); GPUParallelFor(maxQueueSize, KRR_DEVICE_LAMBDA(int pixelId) { - Color3f L = Color3f(pixelState->L[pixelId]) / samplesPerPixel; - if (enableClamp) L = clamp(L, 0.f, clampMax); - m_image->put(Color4f(L, 1.f), pixelId); +#if KRR_RENDER_SPECTRAL + RGB L = pixelState->L[pixelId].toRGB(pixelState->lambda[pixelId], + *KRR_DEFAULT_COLORSPACE_GPU) / samplesPerPixel; +#else + RGB L = RGB(pixelState->L[pixelId]) / samplesPerPixel; +#endif + if (enableClamp) L.clamp(0, clampMax); + m_image->put(RGBA(L, 1.f), pixelId); if (m_renderMode == RenderMode::Interactive) - frameBuffer.write(Color4f(L, 1), pixelId); + frameBuffer.write(RGBA(L, 1), pixelId); else if (m_renderMode == RenderMode::Offline) frameBuffer.write(m_image->getPixel(pixelId), pixelId); }); @@ -286,9 +293,15 @@ void PPGPathTracer::endFrame(RenderContext* context) { PROFILE("Training SD-Tree"); GPUParallelFor(maxQueueSize, KRR_DEVICE_LAMBDA(int pixelId){ Sampler sampler = &pixelState->sampler[pixelId]; - Color pixelEstimate = 0.5f; + SampledSpectrum pixelEstimate(0.5); if (m_isBuilt && m_distribution == EDistribution::EFull) +#if KRR_RENDER_SPECTRAL + pixelEstimate = RGBUnboundedSpectrum(m_pixelEstimate->getPixel(pixelId).head<3>(), + *KRR_DEFAULT_COLORSPACE_GPU) + .sample(pixelState->lambda[pixelId]); +#else pixelEstimate = m_pixelEstimate->getPixel(pixelId).head<3>(); +#endif guidedPathState->commitAll(pixelId, m_sdTree, 1.f, m_spatialFilter, m_directionalFilter, m_bsdfSamplingFractionLoss, sampler, m_distribution, pixelEstimate); }); @@ -508,7 +521,7 @@ void PPGPathTracer::filterFrame(Film *image) { image->getInternalBuffer().copy_to_host(reinterpret_cast(data), n_pixels); constexpr int FILTER_ITERATIONS = 3; - constexpr int SPECTRUM_SAMPLES = Color4f::dim; + constexpr int SPECTRUM_SAMPLES = nSpectrumSamples; for (int i = 0, j = size[0] * size[1]; i < j; ++i) { int isBlack = true; diff --git a/src/misc/render/ppg/integrator.h b/src/misc/render/ppg/integrator.h index 63ea62b8..a2b780b2 100644 --- a/src/misc/render/ppg/integrator.h +++ b/src/misc/render/ppg/integrator.h @@ -34,8 +34,6 @@ class PPGPathTracer : public WavefrontPathTracer{ void handleMiss(); void handleIntersections(); void generateScatterRays(); - void sampleMediumInteraction(int depth); - void sampleMediumScattering(int depth); void traceClosest(int depth); void traceShadow(); diff --git a/src/misc/render/ppg/medium.cpp b/src/misc/render/ppg/medium.cpp deleted file mode 100644 index e2ebc040..00000000 --- a/src/misc/render/ppg/medium.cpp +++ /dev/null @@ -1,159 +0,0 @@ -#include -#include - -#include "device/cuda.h" -#include "integrator.h" -#include "guideditem.h" -#include "render/media.h" -#include "render/profiler/profiler.h" -#include "render/wavefront/wavefront.h" - -KRR_NAMESPACE_BEGIN - -void PPGPathTracer::sampleMediumInteraction(int depth) { - PROFILE("Sample medium interaction"); - ForAllQueued( - mediumSampleQueue, maxQueueSize, - KRR_DEVICE_LAMBDA(const MediumSampleWorkItem &w) { - /* Each ray is either absorped or scattered after some null collisions... */ - Color L(0); - Ray ray = w.ray; - Color thp = w.thp; - Color pu = w.pu, pl = w.pl; - Sampler sampler = &pixelState->sampler[w.pixelId]; - SampledChannel channel = pixelState->channel[w.pixelId]; - Medium medium = ray.medium; - bool scattered{false}; - - Color T_maj = sampleT_maj( - ray, w.tMax, sampler, channel, - [&](Vector3f p, MediumProperties mp, Color sigma_maj, Color T_maj) -> bool { - if (w.depth < maxDepth && mp.Le.any()) { - float pr = sigma_maj[channel] * T_maj[channel]; - Color pe = pu * sigma_maj * T_maj / pr; - if (pe.any()) L += thp * mp.sigma_a * T_maj * mp.Le / (pr * pe.mean()); - } - /* [STEP.2] Sample a type of the three scattering events */ - float pAbsorb = mp.sigma_a[channel] / sigma_maj[channel]; - float pScatter = mp.sigma_s[channel] / sigma_maj[channel]; - float pNull = max(0.f, 1.f - pAbsorb - pScatter); - int mode = sampleDiscrete({pAbsorb, pScatter, pNull}, sampler.get1D()); - if (mode == 0) { // Absorbed (goodbye) - thp = 0; // Will not continue - return false; - } else if (mode == 1) { // Real scattering - float pr = T_maj[channel] * mp.sigma_s[channel]; - thp *= T_maj * mp.sigma_s / pr; - pu *= T_maj * mp.sigma_s / pr; - if (thp.any() && pu.any()) - mediumScatterQueue->push(p, thp, pu, -ray.dir, ray.time, ray.medium, - mp.phase, w.depth, w.pixelId); - scattered = true; // Continue on another direction - return false; - } else { // Null-collision - Color sigma_n = (sigma_maj - mp.sigma_a - mp.sigma_s).cwiseMax(0); - float pr = T_maj[channel] * sigma_n[channel]; - thp *= T_maj * sigma_n / pr; - if (pr == 0) thp = 0; - pu *= T_maj * sigma_n / pr; - pl *= T_maj * sigma_maj / pr; - return thp.any() && pu.any(); - } - }); - - // Add any contribution along this volumetric path... - if (L.any()) pixelState->addRadiance(w.pixelId, L); - - if (!scattered && thp.any()) { - /* Update statistics... */ - thp *= T_maj / T_maj[channel]; - pu *= T_maj / T_maj[channel]; - pl *= T_maj / T_maj[channel]; - } - - // Return if scattered or no throughput or max depth reached - if (scattered || !thp.any() || !pu.any() || w.depth == maxDepth) return; - - // [Grand Survival] There are three cases needed to handle... - // [I am free...] The ray escaped from our sight... - if (w.tMax == M_FLOAT_INF) { - /* [CHECK] Is it correct to use w.bsdfType here? */ - missRayQueue->push(ray, w.ctx, thp, pu, pl, w.bsdfType, w.depth, w.pixelId); - return; - } - - // [You can not see me] The surface do not pocess a material (usually an interface?) - if (!w.intr.material) { - /* Just let it go (use *argument* _depth_ here (not w.depth). ) */ - nextRayQueue(depth)->push(w.intr.spawnRayTowards(ray.dir), w.ctx, thp, pu, pl, - w.depth, w.pixelId, w.bsdfType); - return; - } - - // [Moon landing] We finally reached the surface...! - if (w.intr.light) { - // Handle contribution for light hit - /* The light is sampled from the last vertex on surface, so use its context! */ - hitLightRayQueue->push(w.intr, w.ctx, w.bsdfType, w.depth, w.pixelId, thp, pu, pl); - } - - // [Tomorrow is another day] Next surface scattering event...! - scatterRayQueue->push(w.intr, thp, pu, w.depth, w.pixelId); - }, - gpContext->cudaStream); -} - -void PPGPathTracer::sampleMediumScattering(int depth) { - PROFILE("Sample medium scattering"); - ForAllQueued( - mediumScatterQueue, maxQueueSize, - KRR_DEVICE_LAMBDA(const MediumScatterWorkItem &w) { - const Vector3f &wo = w.wo; - Sampler sampler = &pixelState->sampler[w.pixelId]; - LightSampleContext ctx{w.p, Vector3f::Zero()}; - // [PART-A] Sample direct lighting with ShadowRayTr - if (enableNEE) { - SampledLight sampledLight = lightSampler.sample(sampler.get1D()); - Light light = sampledLight.light; - LightSample ls = light.sampleLi(sampler.get2D(), ctx); - Ray shadowRay = Interaction(w.p, w.time, w.medium).spawnRayTo(ls.intr); - Vector3f wi = shadowRay.dir.normalized(); - Color thp = w.thp * w.phase.p(wo, wi); - float lightPdf = sampledLight.pdf * ls.pdf; - float phasePdf = light.isDeltaLight() ? 0 : w.phase.pdf(wo, wi); - - Color Ld = thp * ls.L; - if (Ld.any() && lightPdf > 0) { - ShadowRayWorkItem sw = {}; - sw.ray = shadowRay; - sw.pl = w.pu * lightPdf; - sw.pu = w.pu * phasePdf; - sw.Ld = Ld; - sw.pixelId = w.pixelId; - sw.tMax = 1; - shadowRayQueue->push(sw); - } - } - - // [PART-B] Sample indirect lighting with scattering function - PhaseFunctionSample ps = w.phase.sample(wo, sampler.get2D()); - Color thp = w.thp * ps.p / ps.pdf; - // Russian roulette - float rrProb = (thp / w.pu.mean()).maxCoeff(); - if (w.depth >= 1 && rrProb < 1) { - if (w.depth >= 1 && sampler.get1D() >= rrProb) return; - thp /= rrProb; - } - - Ray ray{w.p, ps.wi, w.time, w.medium}; - if (!thp.isZero() && !thp.hasNaN()) - /* [NOTE] We need to multiply P_path by w.pu. While the PDF of P_light and P_bsdf is - the same until this vertex, the channel-wise PDF along the path may be - different. */ - nextRayQueue(depth)->push(ray, ctx, thp, w.pu, w.pu / ps.pdf, w.depth + 1, - w.pixelId, BSDF_SMOOTH); - }, - gpContext->cudaStream); -} - -KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/misc/render/ppg/ppg.h b/src/misc/render/ppg/ppg.h index df445219..8880a377 100644 --- a/src/misc/render/ppg/ppg.h +++ b/src/misc/render/ppg/ppg.h @@ -15,6 +15,7 @@ typedef struct { bool enableTraining; PixelStateBuffer* pixelState; + const RGBColorSpace *colorSpace; rt::SceneData sceneData; GuidedPathStateBuffer* guidedState; OptixTraversableHandle traversable; diff --git a/src/misc/render/zero_guiding/device.cu b/src/misc/render/zero_guiding/device.cu index e1c390e5..99af9b19 100644 --- a/src/misc/render/zero_guiding/device.cu +++ b/src/misc/render/zero_guiding/device.cu @@ -30,23 +30,25 @@ KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, traceRay(traversable, ray, tMax, rayType, flags, u0, u1); } +KRR_DEVICE_FUNCTION int getRayId() { return optixGetLaunchIndex().x; } + KRR_DEVICE_FUNCTION RayWorkItem getRayWorkItem() { - int rayIndex(optixGetLaunchIndex().x); - DCHECK_LT(rayIndex, launchParams.currentRayQueue->size()); - return (*launchParams.currentRayQueue)[rayIndex]; + DCHECK_LT(getRayId(), launchParams.currentRayQueue->size()); + return (*launchParams.currentRayQueue)[getRayId()]; } KRR_DEVICE_FUNCTION ShadowRayWorkItem getShadowRayWorkItem() { - int rayIndex(optixGetLaunchIndex().x); - DCHECK_LT(rayIndex, launchParams.shadowRayQueue->size()); - return (*launchParams.shadowRayQueue)[rayIndex]; + DCHECK_LT(getRayId(), launchParams.shadowRayQueue->size()); + return (*launchParams.shadowRayQueue)[getRayId()]; } extern "C" __global__ void KRR_RT_CH(Closest)() { HitInfo hitInfo = getHitInfo(); SurfaceInteraction& intr = *getPRD(); RayWorkItem r = getRayWorkItem(); - prepareSurfaceInteraction(intr, hitInfo); + int pixelId = launchParams.currentRayQueue->pixelId[getRayId()]; + SampledWavelengths lambda = launchParams.pixelState->lambda[pixelId]; + prepareSurfaceInteraction(intr, hitInfo, launchParams.colorSpace, lambda); if (launchParams.mediumSampleQueue && r.ray.medium) { launchParams.mediumSampleQueue->push(r, intr, optixGetRayTmax()); return; @@ -72,8 +74,7 @@ extern "C" __global__ void KRR_RT_MS(Closest)() { } extern "C" __global__ void KRR_RT_RG(Closest)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.currentRayQueue->size()) return; + if (getRayId() >= launchParams.currentRayQueue->size()) return; RayWorkItem r = getRayWorkItem(); SurfaceInteraction intr = {}; traceRay(launchParams.traversable, r.ray, M_FLOAT_INF, @@ -81,22 +82,20 @@ extern "C" __global__ void KRR_RT_RG(Closest)() { } extern "C" __global__ void KRR_RT_AH(Shadow)() { - if (alphaKilled()) - optixIgnoreIntersection(); + if (alphaKilled()) optixIgnoreIntersection(); } extern "C" __global__ void KRR_RT_MS(Shadow)() { optixSetPayload_0(1); } extern "C" __global__ void KRR_RT_RG(Shadow)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); uint32_t visible{0}; traceRay(launchParams.traversable, r.ray, r.tMax, 1, OptixRayFlags( OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT | OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT), visible); if (visible) { - Color contrib = r.Ld / (r.pl + r.pu).mean(); + SampledSpectrum contrib = r.Ld / (r.pl + r.pu).mean(); launchParams.pixelState->addRadiance(r.pixelId, contrib); if (launchParams.enableTraining) launchParams.guidedState->recordRadiance(r.pixelId, contrib); @@ -105,8 +104,10 @@ extern "C" __global__ void KRR_RT_RG(Shadow)() { extern "C" __global__ void KRR_RT_CH(ShadowTr)() { HitInfo hitInfo = getHitInfo(); + int pixelId = launchParams.shadowRayQueue->pixelId[getRayId()]; + SampledWavelengths lambda = launchParams.pixelState->lambda[pixelId]; SurfaceInteraction &intr = *getPRD(); - prepareSurfaceInteraction(intr, hitInfo); + prepareSurfaceInteraction(intr, hitInfo, launchParams.colorSpace, lambda); } extern "C" __global__ void KRR_RT_AH(ShadowTr)() { @@ -116,8 +117,7 @@ extern "C" __global__ void KRR_RT_AH(ShadowTr)() { extern "C" __global__ void KRR_RT_MS(ShadowTr)() { optixSetPayload_2(1); } extern "C" __global__ void KRR_RT_RG(ShadowTr)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); SurfaceInteraction intr = {}; uint u0, u1; diff --git a/src/misc/render/zero_guiding/guideditem.h b/src/misc/render/zero_guiding/guideditem.h index da826de5..f23d13b5 100644 --- a/src/misc/render/zero_guiding/guideditem.h +++ b/src/misc/render/zero_guiding/guideditem.h @@ -17,14 +17,14 @@ struct Vertex { DTreeWrapper* dTree; Vector3f dTreeVoxelSize; Ray ray; - Color3f throughput; - Color3f bsdfVal; - Color3f radiance; + SampledSpectrum throughput; + SampledSpectrum bsdfVal; + SampledSpectrum radiance; float wiPdf, bsdfPdf, dTreePdf; float wiMisWeight; // @addition VAPG bool isDelta; - KRR_DEVICE void record(const Color3f& r) { + KRR_DEVICE void record(const SampledSpectrum &r) { radiance += r; } @@ -32,18 +32,17 @@ struct Vertex { EDirectionalFilter directionalFilter, EBsdfSamplingFractionLoss bsdfSamplingFractionLoss, Sampler &sampler, EDistribution distribution = EDistribution::ERadiance, - const Color &pixelEstimate = {}) { - if (wiPdf <= 0 || isDelta) { - return; - } - Color3f localRadiance = 0; + const SampledSpectrum &pixelEstimate = {}) { + if (wiPdf <= 0 || isDelta) return; + + SampledSpectrum localRadiance(0); if (throughput[0] * wiPdf > 1e-4f) localRadiance[0] = radiance[0] / throughput[0]; if (throughput[1] * wiPdf > 1e-4f) localRadiance[1] = radiance[1] / throughput[1]; if (throughput[2] * wiPdf > 1e-4f) localRadiance[2] = radiance[2] / throughput[2]; - Color3f product = localRadiance * bsdfVal; + SampledSpectrum product = localRadiance * bsdfVal; /* @modified: VAPG */ float value = localRadiance.mean(); @@ -66,11 +65,8 @@ struct Vertex { break; } - DTreeRecord rec{ ray.dir, - value, product.mean(), - wiPdf, bsdfPdf, dTreePdf, - statisticalWeight, - isDelta }; + DTreeRecord rec{ray.dir, value, product.mean(), wiPdf, + bsdfPdf, dTreePdf, statisticalWeight, isDelta}; switch (spatialFilter) { case ESpatialFilter::ENearest: @@ -108,7 +104,7 @@ struct GuidedPathState { uint n_vertices{}; }; -#include "ppg/guideditem_soa.h" +#include "zero_guiding/guideditem_soa.h" class GuidedRayQueue : public WorkQueue { public: @@ -123,17 +119,14 @@ class GuidedPathStateBuffer : public SOA { GuidedPathStateBuffer() = default; GuidedPathStateBuffer(int n, Allocator alloc) : SOA(n, alloc) {} - KRR_DEVICE void incrementDepth(int pixelId, - Ray& ray, - DTreeWrapper* dTree, - Vector3f dTreeVoxelSize, - Color3f thp, - Color3f bsdfVal, - float wiPdf, float bsdfPdf, float dTreePdf, - bool isDelta = false) { + KRR_DEVICE void incrementDepth(int pixelId, Ray &ray, DTreeWrapper *dTree, + Vector3f dTreeVoxelSize, SampledSpectrum thp, + SampledSpectrum bsdfVal, float wiPdf, float bsdfPdf, + float dTreePdf, bool isDelta = false) { int depth = n_vertices[pixelId]; if (depth >= MAX_TRAIN_DEPTH) return; - vertices[depth].radiance[pixelId] = Color3f(0); /* Always remember to clear the radiance from last frame, as I forgotten... */ + /* Always remember to clear the radiance from last frame, as I forgotten... */ + vertices[depth].radiance[pixelId] = SampledSpectrum(0); vertices[depth].ray[pixelId] = ray; vertices[depth].dTree[pixelId] = dTree; vertices[depth].dTreeVoxelSize[pixelId] = dTreeVoxelSize; @@ -148,10 +141,10 @@ class GuidedPathStateBuffer : public SOA { n_vertices[pixelId] = 1 + depth; } - KRR_DEVICE void recordRadiance(int pixelId, Color3f L) { + KRR_DEVICE void recordRadiance(int pixelId, SampledSpectrum L) { int cur_depth = n_vertices[pixelId]; for (int i = 0; i < cur_depth; i++) { - Color3f prevRadiance = vertices[i].radiance[pixelId]; + SampledSpectrum prevRadiance = vertices[i].radiance[pixelId]; vertices[i].radiance[pixelId] = prevRadiance + L; } } @@ -161,7 +154,7 @@ class GuidedPathStateBuffer : public SOA { STree* sdTree, float statisticalWeight, ESpatialFilter spatialFilter, EDirectionalFilter directionalFilter, EBsdfSamplingFractionLoss bsdfSamplingFractionLoss, Sampler& sampler, - EDistribution distribution = EDistribution::ERadiance, const Color3f &pixelEstimate = {}) { + EDistribution distribution = EDistribution::ERadiance, const SampledSpectrum &pixelEstimate = {}) { for (int i = 0; i < n_vertices[pixelId]; i++) { Vertex v = vertices[i][pixelId]; v.commit(sdTree, statisticalWeight, diff --git a/src/misc/render/zero_guiding/guideditem.soa b/src/misc/render/zero_guiding/guideditem.soa index 9a8b60c2..33041537 100644 --- a/src/misc/render/zero_guiding/guideditem.soa +++ b/src/misc/render/zero_guiding/guideditem.soa @@ -12,8 +12,7 @@ flat HaltonSampler; flat DTreeWrapper; flat Vector3f; -flat Color3f; -flat Color; +flat SampledSpectrum; struct Ray; struct LightSampleContext; @@ -26,9 +25,9 @@ struct Vertex { DTreeWrapper* dTree; Vector3f dTreeVoxelSize; Ray ray; - Color3f throughput; - Color3f bsdfVal; - Color3f radiance; + SampledSpectrum throughput; + SampledSpectrum bsdfVal; + SampledSpectrum radiance; float wiPdf, bsdfPdf, dTreePdf; float wiMisWeight; bool isDelta; diff --git a/src/misc/render/zero_guiding/integrator.cpp b/src/misc/render/zero_guiding/integrator.cpp index 310c3281..8497d044 100644 --- a/src/misc/render/zero_guiding/integrator.cpp +++ b/src/misc/render/zero_guiding/integrator.cpp @@ -77,12 +77,14 @@ void ZeroGuidingPT::traceClosest(int depth) { static LaunchParamsZero params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); + params.colorSpace = KRR_DEFAULT_COLORSPACE; params.currentRayQueue = currentRayQueue(depth); params.missRayQueue = missRayQueue; params.hitLightRayQueue = hitLightRayQueue; params.scatterRayQueue = scatterRayQueue; params.nextRayQueue = nextRayQueue(depth); params.mediumSampleQueue = enableMedium ? mediumSampleQueue : nullptr; + params.pixelState = pixelState; backend->launch(params, "Closest", maxQueueSize, 1, 1); } @@ -91,6 +93,7 @@ void ZeroGuidingPT::traceShadow() { static LaunchParamsZero params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); + params.colorSpace = KRR_DEFAULT_COLORSPACE; params.shadowRayQueue = shadowRayQueue; params.pixelState = pixelState; params.guidedState = guidedPathState; @@ -103,7 +106,8 @@ void ZeroGuidingPT::handleHit() { PROFILE("Process intersected rays"); ForAllQueued(hitLightRayQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const HitLightWorkItem & w){ - Color Le = w.light.L(w.p, w.n, w.uv, w.wo) * w.thp; + SampledSpectrum Le = + w.light.L(w.p, w.n, w.uv, w.wo, pixelState->lambda[w.pixelId]) * w.thp; if (enableNEE && w.depth && !(w.bsdfType & BSDF_SPECULAR)) { Interaction intr(w.p, w.wo, w.n, w.uv); float lightPdf = w.light.pdfLi(intr, w.ctx) * lightSampler.pdf(w.light); @@ -118,13 +122,14 @@ void ZeroGuidingPT::handleMiss() { PROFILE("Process escaped rays"); const rt::SceneData& sceneData = mScene->mSceneRT->getSceneData(); ForAllQueued(missRayQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const MissRayWorkItem & w) { - Color L = {}; + SampledSpectrum L = {}; + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; Interaction intr(w.ray.origin); for (const rt::InfiniteLight &light : sceneData.infiniteLights) { if (enableNEE && w.depth && !(w.bsdfType & BSDF_SPECULAR)) { float lightPdf = light.pdfLi(intr, w.ctx) * lightSampler.pdf(&light); - L += light.Li(w.ray.dir) / (w.pu + w.pl * lightPdf).mean(); - } else L += light.Li(w.ray.dir) / w.pu.mean(); + L += light.Li(w.ray.dir, lambda) / (w.pu + w.pl * lightPdf).mean(); + } else L += light.Li(w.ray.dir, lambda) / w.pu.mean(); } pixelState->addRadiance(w.pixelId, w.thp * L); guidedPathState->recordRadiance(w.pixelId, w.thp * L); // ppg @@ -142,6 +147,7 @@ void ZeroGuidingPT::handleIntersections() { const SurfaceInteraction& intr = w.intr; BSDFType bsdfType = intr.getBsdfType(); Vector3f woLocal = intr.toLocal(intr.wo); + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; /* Statistics for mixed bsdf-guided sampling */ float bsdfPdf, dTreePdf; @@ -150,13 +156,13 @@ void ZeroGuidingPT::handleIntersections() { if (enableNEE && (bsdfType & BSDF_SMOOTH)) { SampledLight sampledLight = lightSampler.sample(sampler.get1D()); Light light = sampledLight.light; - LightSample ls = light.sampleLi(sampler.get2D(), { intr.p, intr.n }); + LightSample ls = light.sampleLi(sampler.get2D(), {intr.p, intr.n}, lambda); Ray shadowRay = intr.spawnRayTo(ls.intr); Vector3f wiWorld = normalize(shadowRay.dir); Vector3f wiLocal = intr.toLocal(wiWorld); float lightPdf = sampledLight.pdf * ls.pdf; - Color bsdfVal = BxDF::f(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); + SampledSpectrum bsdfVal = BxDF::f(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); float bsdfPdf = light.isDeltaLight() ? 0 : BxDF::pdf(intr, woLocal, wiLocal, (int) intr.sd.bsdfType); if (lightPdf > 0 && bsdfVal.any()) { ShadowRayWorkItem sw = {}; @@ -258,11 +264,16 @@ void ZeroGuidingPT::render(RenderContext *context) { // write results of the current frame... CudaRenderTarget frameBuffer = context->getColorTexture()->getCudaRenderTarget(); GPUParallelFor(maxQueueSize, KRR_DEVICE_LAMBDA(int pixelId) { - Color3f L = Color3f(pixelState->L[pixelId]) / samplesPerPixel; - if (enableClamp) L = clamp(L, 0.f, clampMax); - m_image->put(Color4f(L, 1.f), pixelId); +#if KRR_RENDER_SPECTRAL + RGB L = pixelState->L[pixelId].toRGB(pixelState->lambda[pixelId], + *KRR_DEFAULT_COLORSPACE_GPU) / samplesPerPixel; +#else + RGB L = RGB(pixelState->L[pixelId]) / samplesPerPixel; +#endif + if (enableClamp) L.clamp(0, clampMax); + m_image->put(RGBA(L, 1.f), pixelId); if (m_renderMode == RenderMode::Interactive) - frameBuffer.write(Color4f(L, 1), pixelId); + frameBuffer.write(RGBA(L, 1), pixelId); else if (m_renderMode == RenderMode::Offline) frameBuffer.write(m_image->getPixel(pixelId), pixelId); }); @@ -285,9 +296,15 @@ void ZeroGuidingPT::endFrame(RenderContext* context) { PROFILE("Training SD-Tree"); GPUParallelFor(maxQueueSize, KRR_DEVICE_LAMBDA(int pixelId){ Sampler sampler = &pixelState->sampler[pixelId]; - Color pixelEstimate = 0.5f; + SampledSpectrum pixelEstimate(0.5); if (m_isBuilt && m_distribution == EDistribution::EFull) +#if KRR_RENDER_SPECTRAL + pixelEstimate = RGBUnboundedSpectrum(m_pixelEstimate->getPixel(pixelId).head<3>(), + *KRR_DEFAULT_COLORSPACE_GPU) + .sample(pixelState->lambda[pixelId]); +#else pixelEstimate = m_pixelEstimate->getPixel(pixelId).head<3>(); +#endif guidedPathState->commitAll(pixelId, m_sdTree, 1.f, m_spatialFilter, m_directionalFilter, m_bsdfSamplingFractionLoss, sampler, m_distribution, pixelEstimate); }); diff --git a/src/misc/render/zero_guiding/medium.cpp b/src/misc/render/zero_guiding/medium.cpp index e78b064b..8bd114df 100644 --- a/src/misc/render/zero_guiding/medium.cpp +++ b/src/misc/render/zero_guiding/medium.cpp @@ -16,49 +16,50 @@ void ZeroGuidingPT::sampleMediumInteraction(int depth) { mediumSampleQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const MediumSampleWorkItem &w) { /* Each ray is either absorped or scattered after some null collisions... */ - Color L(0); + SampledSpectrum L(0); Ray ray = w.ray; - Color thp = w.thp; - Color pu = w.pu, pl = w.pl; - Sampler sampler = &pixelState->sampler[w.pixelId]; - SampledChannel channel = pixelState->channel[w.pixelId]; - Medium medium = ray.medium; + SampledSpectrum thp = w.thp; + SampledSpectrum pu = w.pu, pl = w.pl; + Sampler sampler = &pixelState->sampler[w.pixelId]; + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; + int channel = lambda.mainIndex(); + Medium medium = ray.medium; bool scattered{false}; - Color T_maj = sampleT_maj( - ray, w.tMax, sampler, channel, - [&](Vector3f p, MediumProperties mp, Color sigma_maj, Color T_maj) -> bool { + SampledSpectrum T_maj = sampleT_maj(ray, w.tMax, sampler, lambda, + [&](Vector3f p, MediumProperties mp, SampledSpectrum sigma_maj, + SampledSpectrum T_maj) -> bool { if (w.depth < maxDepth && mp.Le.any()) { float pr = sigma_maj[channel] * T_maj[channel]; - Color pe = pu * sigma_maj * T_maj / pr; + SampledSpectrum pe = pu * sigma_maj * T_maj / pr; if (pe.any()) L += thp * mp.sigma_a * T_maj * mp.Le / (pr * pe.mean()); } /* [STEP.2] Sample a type of the three scattering events */ - float pAbsorb = mp.sigma_a[channel] / sigma_maj[channel]; - float pScatter = mp.sigma_s[channel] / sigma_maj[channel]; - float pNull = max(0.f, 1.f - pAbsorb - pScatter); - int mode = sampleDiscrete({pAbsorb, pScatter, pNull}, sampler.get1D()); - if (mode == 0) { // Absorbed (goodbye) - thp = 0; // Will not continue + float pAbsorb = mp.sigma_a[channel] / sigma_maj[channel]; + float pScatter = mp.sigma_s[channel] / sigma_maj[channel]; + float pNull = max(0.f, 1.f - pAbsorb - pScatter); + int mode = sampleDiscrete({pAbsorb, pScatter, pNull}, sampler.get1D()); + if (mode == 0) { // Absorbed (goodbye) + thp = SampledSpectrum::Zero(); // Will not continue return false; - } else if (mode == 1) { // Real scattering + } else if (mode == 1) { // Real scattering float pr = T_maj[channel] * mp.sigma_s[channel]; thp *= T_maj * mp.sigma_s / pr; pu *= T_maj * mp.sigma_s / pr; if (thp.any() && pu.any()) mediumScatterQueue->push(p, thp, pu, -ray.dir, ray.time, ray.medium, mp.phase, w.depth, w.pixelId); - scattered = true; // Continue on another direction + scattered = true; // Continue on another direction return false; - } else { // Null-collision - Color sigma_n = (sigma_maj - mp.sigma_a - mp.sigma_s).cwiseMax(0); - float pr = T_maj[channel] * sigma_n[channel]; + } else { // Null-collision + SampledSpectrum sigma_n = (sigma_maj - mp.sigma_a - mp.sigma_s).cwiseMax(0); + float pr = T_maj[channel] * sigma_n[channel]; thp *= T_maj * sigma_n / pr; - if (pr == 0) thp = 0; + if (pr == 0) thp = SampledSpectrum::Zero(); pu *= T_maj * sigma_n / pr; pl *= T_maj * sigma_maj / pr; return thp.any() && pu.any(); - } + } }); // Add any contribution along this volumetric path... @@ -83,7 +84,7 @@ void ZeroGuidingPT::sampleMediumInteraction(int depth) { } // [You can not see me] The surface do not pocess a material (usually an interface?) - if (!w.intr.material) { + if (!w.intr.material) { /* Just let it go (use *argument* _depth_ here (not w.depth). ) */ nextRayQueue(depth)->push(w.intr.spawnRayTowards(ray.dir), w.ctx, thp, pu, pl, w.depth, w.pixelId, w.bsdfType); @@ -94,13 +95,13 @@ void ZeroGuidingPT::sampleMediumInteraction(int depth) { if (w.intr.light) { // Handle contribution for light hit /* The light is sampled from the last vertex on surface, so use its context! */ - hitLightRayQueue->push(w.intr, w.ctx, w.bsdfType, w.depth, w.pixelId, thp, pu, pl); + hitLightRayQueue->push(w.intr, w.ctx, w.bsdfType, w.depth, w.pixelId, thp, + pu, pl); } // [Tomorrow is another day] Next surface scattering event...! - scatterRayQueue->push(w.intr, thp, pu, w.depth, w.pixelId); - }, - gpContext->cudaStream); + scatterRayQueue->push(w.intr, thp, pu, w.depth, w.pixelId); + }, gpContext->cudaStream); } void ZeroGuidingPT::sampleMediumScattering(int depth) { @@ -108,21 +109,22 @@ void ZeroGuidingPT::sampleMediumScattering(int depth) { ForAllQueued( mediumScatterQueue, maxQueueSize, KRR_DEVICE_LAMBDA(const MediumScatterWorkItem &w) { - const Vector3f &wo = w.wo; - Sampler sampler = &pixelState->sampler[w.pixelId]; + const Vector3f &wo = w.wo; + Sampler sampler = &pixelState->sampler[w.pixelId]; + SampledWavelengths lambda = pixelState->lambda[w.pixelId]; LightSampleContext ctx{w.p, Vector3f::Zero()}; // [PART-A] Sample direct lighting with ShadowRayTr if (enableNEE) { SampledLight sampledLight = lightSampler.sample(sampler.get1D()); Light light = sampledLight.light; - LightSample ls = light.sampleLi(sampler.get2D(), ctx); + LightSample ls = light.sampleLi(sampler.get2D(), ctx, lambda); Ray shadowRay = Interaction(w.p, w.time, w.medium).spawnRayTo(ls.intr); Vector3f wi = shadowRay.dir.normalized(); - Color thp = w.thp * w.phase.p(wo, wi); + SampledSpectrum thp = w.thp * w.phase.p(wo, wi); float lightPdf = sampledLight.pdf * ls.pdf; float phasePdf = light.isDeltaLight() ? 0 : w.phase.pdf(wo, wi); - - Color Ld = thp * ls.L; + + SampledSpectrum Ld = thp * ls.L; if (Ld.any() && lightPdf > 0) { ShadowRayWorkItem sw = {}; sw.ray = shadowRay; @@ -134,26 +136,23 @@ void ZeroGuidingPT::sampleMediumScattering(int depth) { shadowRayQueue->push(sw); } } - + // [PART-B] Sample indirect lighting with scattering function PhaseFunctionSample ps = w.phase.sample(wo, sampler.get2D()); - Color thp = w.thp * ps.p / ps.pdf; + SampledSpectrum thp = w.thp * ps.p / ps.pdf; // Russian roulette float rrProb = (thp / w.pu.mean()).maxCoeff(); if (w.depth >= 1 && rrProb < 1) { if (w.depth >= 1 && sampler.get1D() >= rrProb) return; thp /= rrProb; } - + Ray ray{w.p, ps.wi, w.time, w.medium}; - if (!thp.isZero() && !thp.hasNaN()) - /* [NOTE] We need to multiply P_path by w.pu. While the PDF of P_light and P_bsdf is - the same until this vertex, the channel-wise PDF along the path may be - different. */ - nextRayQueue(depth)->push(ray, ctx, thp, w.pu, w.pu / ps.pdf, w.depth + 1, - w.pixelId, BSDF_SMOOTH); - }, - gpContext->cudaStream); + if (!thp.isZero() && !thp.hasNaN()) + /* [NOTE] We need to multiply P_path by w.pu. While the PDF of P_light and P_bsdf is + the same until this vertex, the channel-wise PDF along the path may be different. */ + nextRayQueue(depth)->push(ray, ctx, thp, w.pu, w.pu / ps.pdf, w.depth + 1, w.pixelId, BSDF_SMOOTH); + }, gpContext->cudaStream); } KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/misc/render/zero_guiding/zeroguiding.h b/src/misc/render/zero_guiding/zeroguiding.h index 8a4356cb..a65a50e7 100644 --- a/src/misc/render/zero_guiding/zeroguiding.h +++ b/src/misc/render/zero_guiding/zeroguiding.h @@ -15,6 +15,7 @@ typedef struct { bool enableTraining; PixelStateBuffer* pixelState; + const RGBColorSpace *colorSpace; rt::SceneData sceneData; GuidedPathStateBuffer* guidedState; OptixTraversableHandle traversable; diff --git a/src/render/wavefront/device.cu b/src/render/wavefront/device.cu index a5f11a6b..521cb2e8 100644 --- a/src/render/wavefront/device.cu +++ b/src/render/wavefront/device.cu @@ -32,14 +32,13 @@ KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, KRR_DEVICE_FUNCTION int getRayId() { return optixGetLaunchIndex().x; } KRR_DEVICE_FUNCTION RayWorkItem getRayWorkItem() { - DCHECK_LT(getRayIndex(), launchParams.currentRayQueue->size()); + DCHECK_LT(getRayId(), launchParams.currentRayQueue->size()); return (*launchParams.currentRayQueue)[getRayId()]; } KRR_DEVICE_FUNCTION ShadowRayWorkItem getShadowRayWorkItem() { - int rayIndex(optixGetLaunchIndex().x); - DCHECK_LT(rayIndex, launchParams.shadowRayQueue->size()); - return (*launchParams.shadowRayQueue)[rayIndex]; + DCHECK_LT(getRayId(), launchParams.shadowRayQueue->size()); + return (*launchParams.shadowRayQueue)[getRayId()]; } extern "C" __global__ void KRR_RT_CH(Closest)() { @@ -76,8 +75,7 @@ extern "C" __global__ void KRR_RT_MS(Closest)() { } extern "C" __global__ void KRR_RT_RG(Closest)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.currentRayQueue->size()) return; + if (getRayId() >= launchParams.currentRayQueue->size()) return; RayWorkItem r = getRayWorkItem(); SurfaceInteraction intr = {}; traceRay(launchParams.traversable, r.ray, M_FLOAT_INF, @@ -91,8 +89,7 @@ extern "C" __global__ void KRR_RT_AH(Shadow)() { extern "C" __global__ void KRR_RT_MS(Shadow)() { optixSetPayload_0(1); } extern "C" __global__ void KRR_RT_RG(Shadow)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); uint32_t visible{0}; traceRay(launchParams.traversable, r.ray, r.tMax, 1, @@ -116,8 +113,7 @@ extern "C" __global__ void KRR_RT_AH(ShadowTr)() { extern "C" __global__ void KRR_RT_MS(ShadowTr)() { optixSetPayload_2(1); } extern "C" __global__ void KRR_RT_RG(ShadowTr)() { - uint rayIndex(optixGetLaunchIndex().x); - if (rayIndex >= launchParams.shadowRayQueue->size()) return; + if (getRayId() >= launchParams.shadowRayQueue->size()) return; ShadowRayWorkItem r = getShadowRayWorkItem(); SurfaceInteraction intr = {}; uint u0, u1;