Skip to content

Commit

Permalink
merge spectral changes for other integrators
Browse files Browse the repository at this point in the history
  • Loading branch information
@cuteday
cuteday committed Nov 1, 2023
1 parent 6a42fa3 commit 6689098
Show file tree
Hide file tree
Showing 16 changed files with 195 additions and 347 deletions.
4 changes: 2 additions & 2 deletions CMakeLists.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -25,7 +25,7 @@ OPTION (KRR_BUILD_EXAMPLES

OPTION (KRR_RENDER_SPECTRAL
"Build the spectral variant of the renderer (WIP)."
OFF
ON
)

###############################################################################
Expand Down
2 changes: 0 additions & 2 deletions common/build/source.cmake
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
34 changes: 17 additions & 17 deletions src/misc/render/ppg/device.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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<SurfaceInteraction>();
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;
Expand All @@ -72,31 +74,28 @@ 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,
0, OPTIX_RAY_FLAG_NONE, (void*)&intr);
}

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);
Expand All @@ -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<SurfaceInteraction>();
prepareSurfaceInteraction(intr, hitInfo);
prepareSurfaceInteraction(intr, hitInfo, launchParams.colorSpace, lambda);
}

extern "C" __global__ void KRR_RT_AH(ShadowTr)() {
Expand All @@ -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;
Expand Down
47 changes: 20 additions & 27 deletions src/misc/render/ppg/guideditem.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,33 +17,32 @@ 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;
}

KRR_DEVICE void commit(STree *sdTree, float statisticalWeight, ESpatialFilter spatialFilter,
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();
Expand All @@ -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:
Expand Down Expand Up @@ -123,17 +119,14 @@ class GuidedPathStateBuffer : public SOA<GuidedPathState> {
GuidedPathStateBuffer() = default;
GuidedPathStateBuffer(int n, Allocator alloc) : SOA<GuidedPathState>(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;
Expand All @@ -148,10 +141,10 @@ class GuidedPathStateBuffer : public SOA<GuidedPathState> {
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;
}
}
Expand All @@ -161,7 +154,7 @@ class GuidedPathStateBuffer : public SOA<GuidedPathState> {
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,
Expand Down
9 changes: 4 additions & 5 deletions src/misc/render/ppg/guideditem.soa
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,8 +12,7 @@ flat HaltonSampler;
flat DTreeWrapper;

flat Vector3f;
flat Color3f;
flat Color;
flat SampledSpectrum;

struct Ray;
struct LightSampleContext;
Expand All @@ -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;
Expand Down
47 changes: 30 additions & 17 deletions src/misc/render/ppg/integrator.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}

Expand All @@ -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;
Expand All @@ -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);
Expand All @@ -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
Expand All @@ -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;
Expand All @@ -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 = {};
Expand Down Expand Up @@ -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(); });
Expand All @@ -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();
Expand All @@ -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);
});
Expand All @@ -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);
});
Expand Down Expand Up @@ -508,7 +521,7 @@ void PPGPathTracer::filterFrame(Film *image) {
image->getInternalBuffer().copy_to_host(reinterpret_cast<PixelData *>(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;
Expand Down
Loading

0 comments on commit 6689098

Please sign in to comment.