From e1aabe1e32006290f6bc19557b89fb5909e3f173 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E8=9E=8D=E9=9B=AA=E6=97=B6=E5=88=86?= Date: Fri, 10 Nov 2023 20:42:06 +0800 Subject: [PATCH] use LaunchParameter instead of LaunchParamsIntegrator --- .gitmodules | 14 ++++++------- src/core/camera.h | 23 ++++++++++++++------- src/core/device/optix.h | 18 +++++++++------- src/core/device/scene.h | 3 +++ src/main/kiraray.cpp | 2 +- src/misc/render/ppg/device.cu | 2 +- src/misc/render/ppg/integrator.cpp | 4 ++-- src/misc/render/ppg/ppg.h | 7 +++++-- src/misc/render/zero_guiding/device.cu | 2 +- src/misc/render/zero_guiding/integrator.cpp | 4 ++-- src/misc/render/zero_guiding/zeroguiding.h | 7 +++++-- src/render/bdpt/bdpt.h | 4 +++- src/render/bdpt/integrator.cpp | 2 +- src/render/bdpt/integrator.h | 2 +- src/render/passes/gbuffer/device.cu | 2 +- src/render/passes/gbuffer/device.h | 7 +++++-- src/render/passes/gbuffer/gbuffer.cpp | 2 +- src/render/passes/gbuffer/gbuffer.h | 2 +- src/render/path/device.cu | 2 +- src/render/path/path.h | 18 +++++++++------- src/render/path/pathtracer.h | 2 +- src/render/wavefront/device.cu | 2 +- src/render/wavefront/integrator.cpp | 4 ++-- src/render/wavefront/wavefront.h | 12 ++++++----- 24 files changed, 87 insertions(+), 60 deletions(-) diff --git a/.gitmodules b/.gitmodules index 4c0a79a8..8761465a 100644 --- a/.gitmodules +++ b/.gitmodules @@ -1,23 +1,23 @@ [submodule "src/ext/assimp"] path = src/ext/assimp - url = https://github.com/assimp/assimp.git + url = git@github.com:assimp/assimp.git [submodule "src/ext/glfw"] path = src/ext/glfw - url = https://github.com/glfw/glfw.git + url = git@github.com:glfw/glfw.git [submodule "src/ext/pybind11"] path = src/ext/pybind11 - url = https://github.com/pybind/pybind11.git + url = git@github.com:pybind/pybind11.git branch = stable [submodule "src/ext/pbrtparser"] path = src/ext/pbrtparser - url = https://github.com/cuteday/pbrt-parser.git + url = git@github.com:cuteday/pbrt-parser.git branch = cute [submodule "src/core/math/3rdparty/eigen"] path = src/core/math/3rdparty/eigen - url = https://github.com/cuteday/Eigen.git + url = git@github.com:cuteday/Eigen.git [submodule "src/ext/nvrhi"] path = src/ext/nvrhi - url = https://github.com/cuteday/nvrhi.git + url = git@github.com:cuteday/nvrhi.git [submodule "src/ext/openvdb"] path = src/ext/openvdb - url = https://github.com/cuteday/openvdb_win.git + url = git@github.com:cuteday/openvdb_win.git \ No newline at end of file diff --git a/src/core/camera.h b/src/core/camera.h index 6ee0c83a..f5ceed46 100644 --- a/src/core/camera.h +++ b/src/core/camera.h @@ -25,22 +25,29 @@ class Camera { Vector3f u{ 1, 0, 0 }; // camera right [dependent to aspect ratio] Vector3f v{ 0, 1, 0 }; // camera up [dependent to aspect ratio] Vector3f w{ 0, 0, -1 }; // camera forward + + Medium medium{nullptr}; // the ray is inside the medium KRR_CALLABLE Ray getRay(Vector2i pixel, Vector2i frameSize, Sampler& sampler) { Ray ray; /* 1. Statified sample on the film plane (within the fragment) */ - Vector2f p = (Vector2f)pixel + Vector2f(0.5f) + sampler.get2D(); // uniform sample + box filter - Vector2f ndc = Vector2f(2 * p) / Vector2f(frameSize) + Vector2f(-1.f); // ndc in [-1, 1]^2 - if (lensRadius > 0) { /*Thin lens*/ + Vector2f p = + (Vector2f) pixel + Vector2f(0.5f) + sampler.get2D(); // uniform sample + box filter + Vector2f ndc = + Vector2f(2 * p) / Vector2f(frameSize) + Vector2f(-1.f); // ndc in [-1, 1]^2 + if (lensRadius > 0) { /*Thin lens*/ /* 2. Sample the lens (uniform) */ Vector3f focalPoint = pos + ndc[0] * u + ndc[1] * v + w; - Vector2f apertureSample = lensRadius > M_EPSILON ? uniformSampleDisk(sampler.get2D()) : Vector2f::Zero(); - ray.origin = pos + lensRadius * (apertureSample[0] * normalize(u) + apertureSample[1] * normalize(v)); - ray.dir = normalize(focalPoint - ray.origin); - } else { /*Pin hole*/ + Vector2f apertureSample = + lensRadius > M_EPSILON ? uniformSampleDisk(sampler.get2D()) : Vector2f::Zero(); + ray.origin = pos + lensRadius * (apertureSample[0] * normalize(u) + + apertureSample[1] * normalize(v)); + ray.dir = normalize(focalPoint - ray.origin); + } else { /*Pin hole*/ ray.origin = pos; - ray.dir = normalize(ndc[0] * u + ndc[1] * v + w); + ray.dir = normalize(ndc[0] * u + ndc[1] * v + w); } + ray.medium = medium; return ray; } diff --git a/src/core/device/optix.h b/src/core/device/optix.h index 75ad1509..5caa3d09 100644 --- a/src/core/device/optix.h +++ b/src/core/device/optix.h @@ -90,17 +90,19 @@ class OptixBackend { void initialize(const OptixInitializeParameters& params); void setScene(std::shared_ptr _scene); - template - void launch(const T& parameters, string entryPoint, int width, - int height, int depth = 1) { + template + void launch(const LaunchParameters& parameters, string entryPoint, + int width, int height, int depth = 1) { if (height * width * depth == 0) return; - static T *launchParams{nullptr}; - if (!launchParams) cudaMalloc(&launchParams, sizeof(T)); + static LaunchParameters *launchParams{nullptr}; + if (!launchParams) cudaMalloc(&launchParams, sizeof(LaunchParameters)); if (!entryPoints.count(entryPoint)) Log(Fatal, "The entrypoint %s is not initialized!", entryPoint.c_str()); - cudaMemcpyAsync(launchParams, ¶meters, sizeof(T), cudaMemcpyHostToDevice, cudaStream); - OPTIX_CHECK(optixLaunch(optixPipeline, cudaStream, CUdeviceptr(launchParams), - sizeof(T), &SBT[entryPoints[entryPoint]], width, height, depth)); + cudaMemcpyAsync(launchParams, ¶meters, sizeof(LaunchParameters), + cudaMemcpyHostToDevice, cudaStream); + OPTIX_CHECK(optixLaunch(optixPipeline, cudaStream, CUdeviceptr(launchParams), + sizeof(LaunchParameters), &SBT[entryPoints[entryPoint]], + width, height, depth)); } std::shared_ptr getScene() const { return scene; } diff --git a/src/core/device/scene.h b/src/core/device/scene.h index a2f3da46..7bcba790 100644 --- a/src/core/device/scene.h +++ b/src/core/device/scene.h @@ -90,4 +90,7 @@ class RTScene { std::shared_ptr mOptixScene; }; +template +struct LaunchParameters {}; + KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/main/kiraray.cpp b/src/main/kiraray.cpp index fd2872f8..c52542e0 100644 --- a/src/main/kiraray.cpp +++ b/src/main/kiraray.cpp @@ -14,7 +14,7 @@ extern "C" int main(int argc, char *argv[]) { "Switch to Release build for normal performance!"); #endif - string configFile = "common/configs/example_cbox.json"; + string configFile = "common/configs/example_volbox.json"; if (argc < 2){ Log(Warning, "No config file specified, using default config file: %s", configFile.c_str()); } else { diff --git a/src/misc/render/ppg/device.cu b/src/misc/render/ppg/device.cu index ce466de4..8fdd049f 100644 --- a/src/misc/render/ppg/device.cu +++ b/src/misc/render/ppg/device.cu @@ -8,7 +8,7 @@ using namespace krr; KRR_NAMESPACE_BEGIN -extern "C" __constant__ LaunchParamsPPG launchParams; +extern "C" __constant__ LaunchParameters launchParams; template KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, diff --git a/src/misc/render/ppg/integrator.cpp b/src/misc/render/ppg/integrator.cpp index d999ec2a..6760cab8 100644 --- a/src/misc/render/ppg/integrator.cpp +++ b/src/misc/render/ppg/integrator.cpp @@ -75,7 +75,7 @@ void PPGPathTracer::initialize() { void PPGPathTracer::traceClosest(int depth) { PROFILE("Trace intersect rays"); - static LaunchParamsPPG params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; @@ -91,7 +91,7 @@ void PPGPathTracer::traceClosest(int depth) { void PPGPathTracer::traceShadow() { PROFILE("Trace shadow rays"); - static LaunchParamsPPG params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; diff --git a/src/misc/render/ppg/ppg.h b/src/misc/render/ppg/ppg.h index 8880a377..810d1054 100644 --- a/src/misc/render/ppg/ppg.h +++ b/src/misc/render/ppg/ppg.h @@ -4,7 +4,10 @@ KRR_NAMESPACE_BEGIN -typedef struct { +class PPGPathTracer; + +template <> +struct LaunchParameters { RayQueue* currentRayQueue; RayQueue* nextRayQueue; ShadowRayQueue* shadowRayQueue; @@ -19,6 +22,6 @@ typedef struct { rt::SceneData sceneData; GuidedPathStateBuffer* guidedState; OptixTraversableHandle traversable; -} LaunchParamsPPG; +}; KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/misc/render/zero_guiding/device.cu b/src/misc/render/zero_guiding/device.cu index 0650bbf1..a9225df6 100644 --- a/src/misc/render/zero_guiding/device.cu +++ b/src/misc/render/zero_guiding/device.cu @@ -8,7 +8,7 @@ using namespace krr; KRR_NAMESPACE_BEGIN -extern "C" __constant__ LaunchParamsZero launchParams; +extern "C" __constant__ LaunchParameters launchParams; template KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, diff --git a/src/misc/render/zero_guiding/integrator.cpp b/src/misc/render/zero_guiding/integrator.cpp index ec1e8803..ff3646b2 100644 --- a/src/misc/render/zero_guiding/integrator.cpp +++ b/src/misc/render/zero_guiding/integrator.cpp @@ -74,7 +74,7 @@ void ZeroGuidingPT::initialize() { void ZeroGuidingPT::traceClosest(int depth) { PROFILE("Trace intersect rays"); - static LaunchParamsZero params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; @@ -90,7 +90,7 @@ void ZeroGuidingPT::traceClosest(int depth) { void ZeroGuidingPT::traceShadow() { PROFILE("Trace shadow rays"); - static LaunchParamsZero params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; diff --git a/src/misc/render/zero_guiding/zeroguiding.h b/src/misc/render/zero_guiding/zeroguiding.h index a65a50e7..d33b7693 100644 --- a/src/misc/render/zero_guiding/zeroguiding.h +++ b/src/misc/render/zero_guiding/zeroguiding.h @@ -4,7 +4,10 @@ KRR_NAMESPACE_BEGIN -typedef struct { +class ZeroGuidingPT; + +template <> +struct LaunchParameters { RayQueue* currentRayQueue; RayQueue* nextRayQueue; ShadowRayQueue* shadowRayQueue; @@ -19,6 +22,6 @@ typedef struct { rt::SceneData sceneData; GuidedPathStateBuffer* guidedState; OptixTraversableHandle traversable; -} LaunchParamsZero; +}; KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/render/bdpt/bdpt.h b/src/render/bdpt/bdpt.h index 70fd6c22..abd93307 100644 --- a/src/render/bdpt/bdpt.h +++ b/src/render/bdpt/bdpt.h @@ -10,8 +10,10 @@ KRR_NAMESPACE_BEGIN using namespace shader; +class BDPTIntegrator; -struct LaunchParamsBDPT { +template <> +struct LaunchParameters { uint frameID{ 0 }; Vector2i fbSize = Vector2i::Zero(); bool debugOutput = false; diff --git a/src/render/bdpt/integrator.cpp b/src/render/bdpt/integrator.cpp index 9392e411..049b6e2e 100644 --- a/src/render/bdpt/integrator.cpp +++ b/src/render/bdpt/integrator.cpp @@ -9,7 +9,7 @@ KRR_NAMESPACE_BEGIN extern "C" char BDPT_PTX[]; namespace { - LaunchParamsBDPT *launchParamsDevice{}; + LaunchParameters *launchParamsDevice{}; } BDPTIntegrator::BDPTIntegrator() { diff --git a/src/render/bdpt/integrator.h b/src/render/bdpt/integrator.h index 3e274871..d2082bd7 100644 --- a/src/render/bdpt/integrator.h +++ b/src/render/bdpt/integrator.h @@ -33,7 +33,7 @@ class BDPTIntegrator: public RenderPass{ OptixPipeline pipeline; OptixModule module; - LaunchParamsBDPT launchParams; + LaunchParameters launchParams; }; KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/render/passes/gbuffer/device.cu b/src/render/passes/gbuffer/device.cu index 9c71cabc..90a05a77 100644 --- a/src/render/passes/gbuffer/device.cu +++ b/src/render/passes/gbuffer/device.cu @@ -7,7 +7,7 @@ using namespace krr; KRR_NAMESPACE_BEGIN -extern "C" __constant__ LaunchParamsGBuffer launchParams; +extern "C" __constant__ LaunchParameters launchParams; template KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, float tMax, diff --git a/src/render/passes/gbuffer/device.h b/src/render/passes/gbuffer/device.h index 263a6816..22f43cc6 100644 --- a/src/render/passes/gbuffer/device.h +++ b/src/render/passes/gbuffer/device.h @@ -6,13 +6,16 @@ KRR_NAMESPACE_BEGIN -typedef struct { +class GBufferPass; + +template <> +struct LaunchParameters { Vector2i frameSize; size_t frameIndex; rt::SceneData sceneData; Camera::CameraData cameraData; OptixTraversableHandle traversable; -} LaunchParamsGBuffer; +}; KRR_NAMESPACE_END \ No newline at end of file diff --git a/src/render/passes/gbuffer/gbuffer.cpp b/src/render/passes/gbuffer/gbuffer.cpp index c0d670b2..14a67db2 100644 --- a/src/render/passes/gbuffer/gbuffer.cpp +++ b/src/render/passes/gbuffer/gbuffer.cpp @@ -23,7 +23,7 @@ void GBufferPass::setScene(Scene::SharedPtr scene) { void GBufferPass::render(RenderContext* context) { PROFILE("GBuffer drawing"); - static LaunchParamsGBuffer launchParams = {}; + static LaunchParameters launchParams = {}; launchParams.frameIndex = getFrameIndex(); launchParams.frameSize = getFrameSize(); launchParams.cameraData = mScene->getCamera()->getCameraData(); diff --git a/src/render/passes/gbuffer/gbuffer.h b/src/render/passes/gbuffer/gbuffer.h index 6dd53a8c..981572e1 100644 --- a/src/render/passes/gbuffer/gbuffer.h +++ b/src/render/passes/gbuffer/gbuffer.h @@ -20,7 +20,7 @@ class GBufferPass : public RenderPass { private: OptixBackend::SharedPtr mOptixBackend; - LaunchParamsGBuffer mLaunchParams; + LaunchParameters mLaunchParams; bool mEnableDepth{}; bool mEnableDiffuse{}; diff --git a/src/render/path/device.cu b/src/render/path/device.cu index e42d6565..9e2b229c 100644 --- a/src/render/path/device.cu +++ b/src/render/path/device.cu @@ -10,7 +10,7 @@ KRR_NAMESPACE_BEGIN using namespace shader; using namespace rt; -extern "C" __constant__ LaunchParamsPT launchParams; +extern "C" __constant__ LaunchParameters launchParams; template KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, float tMax, diff --git a/src/render/path/path.h b/src/render/path/path.h index d3140ed2..cfdeb89a 100644 --- a/src/render/path/path.h +++ b/src/render/path/path.h @@ -24,26 +24,28 @@ const string shaderProgramNames[RAY_TYPE_COUNT] = { "ShadowTransmittanceRay" }; -struct LaunchParamsPT { - uint frameID{ 0 }; +class MegakernelPathTracer; + +template <> +struct LaunchParameters { + uint frameID{0}; Vector2i fbSize = Vector2i::Zero(); // per pixel debugging output - bool debugOutput = false; - bool NEE = false; // enable next event estimation (and MIS) - Vector2i debugPixel = {960, 540}; + bool debugOutput = false; + bool NEE = false; // enable next event estimation (and MIS) + Vector2i debugPixel = {960, 540}; int maxDepth = 10; float probRR = 0.8; float clampThreshold = 1e4f; // clamp max radiance contrib per frame int spp = 1; int lightSamples = 1; - // scene + // scene Camera::CameraData camera; LightSampler lightSampler; rt::SceneData sceneData; - const RGBColorSpace *colorSpace; CudaRenderTarget colorBuffer; - OptixTraversableHandle traversable{ 0 }; + OptixTraversableHandle traversable{0}; }; struct PathData { diff --git a/src/render/path/pathtracer.h b/src/render/path/pathtracer.h index 6708423f..b5fbf097 100644 --- a/src/render/path/pathtracer.h +++ b/src/render/path/pathtracer.h @@ -27,7 +27,7 @@ class MegakernelPathTracer: public RenderPass{ private: OptixBackend::SharedPtr optixBackend; - LaunchParamsPT launchParams; + LaunchParameters launchParams; friend void to_json(json &j, const MegakernelPathTracer &p) { j = json{ diff --git a/src/render/wavefront/device.cu b/src/render/wavefront/device.cu index ef7252ae..e94759e3 100644 --- a/src/render/wavefront/device.cu +++ b/src/render/wavefront/device.cu @@ -8,7 +8,7 @@ using namespace krr; KRR_NAMESPACE_BEGIN -extern "C" __constant__ LaunchParams launchParams; +extern "C" __constant__ LaunchParameters launchParams; template KRR_DEVICE_FUNCTION void traceRay(OptixTraversableHandle traversable, Ray ray, diff --git a/src/render/wavefront/integrator.cpp b/src/render/wavefront/integrator.cpp index 67d36d0e..66523746 100644 --- a/src/render/wavefront/integrator.cpp +++ b/src/render/wavefront/integrator.cpp @@ -43,7 +43,7 @@ void WavefrontPathTracer::initialize() { void WavefrontPathTracer::traceClosest(int depth) { PROFILE("Trace intersect rays"); // Telling whether volume rendering is enabled by mediumSampleQueue is null? - static LaunchParams params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; @@ -59,7 +59,7 @@ void WavefrontPathTracer::traceClosest(int depth) { void WavefrontPathTracer::traceShadow() { PROFILE("Trace shadow rays"); - static LaunchParams params = {}; + static LaunchParameters params = {}; params.traversable = backend->getRootTraversable(); params.sceneData = backend->getSceneData(); params.colorSpace = KRR_DEFAULT_COLORSPACE; diff --git a/src/render/wavefront/wavefront.h b/src/render/wavefront/wavefront.h index e14e95cc..178733d3 100644 --- a/src/render/wavefront/wavefront.h +++ b/src/render/wavefront/wavefront.h @@ -13,10 +13,12 @@ KRR_NAMESPACE_BEGIN class PixelStateBuffer; +class WavefrontPathTracer; using namespace shader; -typedef struct { +template <> +struct LaunchParameters { RayQueue* currentRayQueue; RayQueue* nextRayQueue; ShadowRayQueue* shadowRayQueue; @@ -29,7 +31,7 @@ typedef struct { const RGBColorSpace *colorSpace; rt::SceneData sceneData; OptixTraversableHandle traversable; -} LaunchParams; +}; template KRR_HOST_DEVICE Spectrum sampleT_maj(Ray ray, float tMax, Sampler sampler, @@ -78,12 +80,12 @@ KRR_HOST_DEVICE Spectrum sampleT_maj(Ray ray, float tMax, Sampler sampler, } template -KRR_HOST_DEVICE void traceTransmittance(ShadowRayWorkItem sr, const SurfaceInteraction& intr, +KRR_HOST_DEVICE void traceTransmittance(const ShadowRayWorkItem& sr, const SurfaceInteraction& intr, PixelStateBuffer *pixelState, TraceFunc trace) { + Ray ray = sr.ray; SampledWavelengths lambda = pixelState->lambda[sr.pixelId]; - int channel = lambda.mainIndex(); Sampler sampler = &pixelState->sampler[sr.pixelId]; - Ray ray = sr.ray; + int channel = lambda.mainIndex(); float tMax = sr.tMax; Vector3f pLight = ray(tMax);