diff --git a/devices/rtx/apps/tests/api/CMakeLists.txt b/devices/rtx/apps/tests/api/CMakeLists.txt index 1c4736642..0e2a494ad 100644 --- a/devices/rtx/apps/tests/api/CMakeLists.txt +++ b/devices/rtx/apps/tests/api/CMakeLists.txt @@ -79,6 +79,7 @@ if (VISRTX_ENABLE_MDL_SUPPORT) target_compile_definitions( TestPbrSpecularDefault PRIVATE VISRTX_TEST_MDL_WRAPPER) make_test(TestEmissiveMdlLight) + make_test(TestEmissionDescriptorPolicy) make_test(TestMdlPipelineRebuild) make_test(TestMdlMaterialRegistryRelease) make_test(TestEmissiveMaterialParity) diff --git a/devices/rtx/apps/tests/api/TestEmissionDescriptorPolicy.cpp b/devices/rtx/apps/tests/api/TestEmissionDescriptorPolicy.cpp new file mode 100644 index 000000000..8a21edb85 --- /dev/null +++ b/devices/rtx/apps/tests/api/TestEmissionDescriptorPolicy.cpp @@ -0,0 +1,177 @@ +/* + * Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: BSD-3-Clause + */ + +// Device-level checks of the emission descriptor registration policy (ADR +// 0007): a raw `mdl` surface becomes a next-event-sampled Geometry Light iff +// its folded descriptor is non-null AND faithfully NEE-evaluable — diffuse EDF, +// radiant- exitance mode, provably non-negative, no geometric-state dependence. +// Each case is asserted via the world's `numLightInstances` property (a +// zero-radiance or forward-only emitter is pixel-identical to no light, so the +// count is the only observable seam). Emission excluded here still renders via +// the forward path, unbiased — this test asserts registration, not radiance. + +// anari_cpp +#define ANARI_EXTENSION_UTILITY_IMPL +#include +#include +// VisRTX +#include +// std +#include +#include +#include + +using vec3 = std::array; + +static void statusFunc(const void *, + ANARIDevice, + ANARIObject, + ANARIDataType, + ANARIStatusSeverity severity, + ANARIStatusCode, + const char *message) +{ + if (severity <= ANARI_SEVERITY_WARNING) + fprintf(stderr, "[anari] %s\n", message); +} + +// A control: constant diffuse radiant-exitance emission — registers. +static const char *POLICY_DIFFUSE = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material diffuse_emit() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(8.0) * math::PI))); +)mdl"; + +// Non-diffuse EDF: not in faithfulSet — described, not registered. +static const char *POLICY_SPOT = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material spot_emit() = material( + surface: material_surface( + emission: material_emission( + emission: df::spot_edf(exponent: 1.0), + intensity: color(8.0) * math::PI))); +)mdl"; + +// Provably-negative emission: sign gate excludes it (its all-negative +// next-event contribution would be dropped by the shadow epsilon gate while the +// forward deposit is MIS-downweighted — bias). Forward-only, unbiased. +static const char *POLICY_NEGATIVE = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material negative_emit() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(-8.0) * math::PI))); +)mdl"; + +// Intensity reads state::normal — a geometric-state quantity the synthetic hit +// fabricates. Faithful EDF kind, but unfaithful integrand: not registered. +static const char *POLICY_STATE = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +import ::state::*; +export material state_emit() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(math::length(state::normal())) * math::PI))); +)mdl"; + +// Power intensity mode: not faithfully handled — described, not registered. +static const char *POLICY_POWER = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material power_emit() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(8.0) * math::PI, + mode: intensity_power))); +)mdl"; + +static bool checkLightCount( + ANARIDevice device, const char *source, const char *name, uint32_t expected) +{ + const std::array pos = {vec3{-0.5f, 1.5f, -0.5f}, + vec3{0.5f, 1.5f, -0.5f}, + vec3{0.5f, 1.5f, 0.5f}, + vec3{-0.5f, 1.5f, 0.5f}}; + const std::array, 2> idx = { + std::array{0, 1, 2}, std::array{0, 2, 3}}; + + auto geom = anari::newObject(device, "triangle"); + anari::setParameterArray1D(device, geom, "vertex.position", pos.data(), 4); + anari::setParameterArray1D(device, geom, "primitive.index", idx.data(), 2); + anari::commitParameters(device, geom); + + auto mat = anari::newObject(device, "mdl"); + anari::setParameter(device, mat, "sourceType", "code"); + anari::setParameter(device, mat, "source", source); + anari::setParameter(device, mat, "materialName", name); + anari::commitParameters(device, mat); + + auto surface = anari::newObject(device); + anari::setAndReleaseParameter(device, surface, "geometry", geom); + anari::setAndReleaseParameter(device, surface, "material", mat); + anari::commitParameters(device, surface); + + auto world = anari::newObject(device); + anari::setParameterArray1D(device, world, "surface", &surface, 1); + anari::release(device, surface); + anari::commitParameters(device, world); + + uint32_t count = ~0u; + const bool found = + anari::getProperty(device, world, "numLightInstances", count, ANARI_WAIT); + anari::release(device, world); + + printf("numLightInstances(%s)=%u (expected %u)\n", name, count, expected); + if (!found) { + fprintf(stderr, "FAIL: world has no numLightInstances property\n"); + return false; + } + if (count != expected) { + fprintf(stderr, + "FAIL: %s expected %u light(s), got %u\n", + name, + expected, + count); + return false; + } + return true; +} + +int main() +{ + auto device = makeVisRTXDevice(statusFunc); + if (!device) { + fprintf(stderr, "FAIL: could not create VisRTX device\n"); + return 1; + } + + bool ok = true; + // Control: a faithful diffuse emitter registers. + ok = checkLightCount(device, POLICY_DIFFUSE, "diffuse_emit", 1) && ok; + // Each faithfulness gate excludes its case (described, forward-only). + ok = checkLightCount(device, POLICY_SPOT, "spot_emit", 0) && ok; + ok = checkLightCount(device, POLICY_NEGATIVE, "negative_emit", 0) && ok; + ok = checkLightCount(device, POLICY_STATE, "state_emit", 0) && ok; + ok = checkLightCount(device, POLICY_POWER, "power_emit", 0) && ok; + + anari::release(device, device); + + if (!ok) { + fprintf(stderr, "TestEmissionDescriptorPolicy FAILED\n"); + return 1; + } + printf("TestEmissionDescriptorPolicy passed\n"); + return 0; +} diff --git a/devices/rtx/apps/tests/api/TestEmissiveMdlLight.cpp b/devices/rtx/apps/tests/api/TestEmissiveMdlLight.cpp index 4d2e4899d..574ef7fb4 100644 --- a/devices/rtx/apps/tests/api/TestEmissiveMdlLight.cpp +++ b/devices/rtx/apps/tests/api/TestEmissiveMdlLight.cpp @@ -585,10 +585,12 @@ int main() ok = checkLightCount(device, MDL_EMISSIVE, "emissive", 1) && ok; ok = checkLightCount(device, MDL_DARK, "dark", 0) && ok; - // Textured intensity with an UNBOUND texture still classifies as a light — - // the documented conservative over-inclusion (unbiased; it merely wastes a - // pick slot while rendering black). - ok = checkLightCount(device, MDL_TEXTURED, "emissive_tex", 1) && ok; + // Textured intensity with an UNBOUND texture is NOT a light: the MDL lookup + // is invalid and folds to 0 (ADR 0007 A5), so the surface emits nothing. This + // supersedes ADR-0006's conservative over-inclusion (which registered a + // zero-radiance light). A BOUND emissive texture still registers — see the + // textured-uniform/checker/away radiance parity above. + ok = checkLightCount(device, MDL_TEXTURED, "emissive_tex", 0) && ok; ok = checkLightCount(device, MDL_NONFINITE, "nonfinite", 0) && ok; ok = checkWrapperLightCount(device, 5.f, 1) && ok; ok = checkWrapperLightCount(device, 0.f, 0) && ok; diff --git a/devices/rtx/apps/tests/unit/CMakeLists.txt b/devices/rtx/apps/tests/unit/CMakeLists.txt index fdf289742..1e4e1966b 100644 --- a/devices/rtx/apps/tests/unit/CMakeLists.txt +++ b/devices/rtx/apps/tests/unit/CMakeLists.txt @@ -42,3 +42,20 @@ target_compile_definitions(${PROJECT_NAME} PRIVATE GLM_ENABLE_EXPERIMENTAL) target_compile_features(${PROJECT_NAME} PRIVATE cxx_std_17) add_test(NAME rtx::LightPickPower COMMAND ${PROJECT_NAME}) + +# classifyEmission lives in libmdl, a standalone static lib with no CUDA/OptiX/GPU +# (the MDL SDK is dlopen'd at runtime), so the classifier is testable on the host. +# Gated on MDL support, since libmdl only exists then. +if (VISRTX_ENABLE_MDL_SUPPORT) + project(rtx_test_MdlEmissionClassifier LANGUAGES CXX) + + add_executable(${PROJECT_NAME} test_MdlEmissionClassifier.cpp) + # libmdl's Core.h includes but links fmt PRIVATE, so a consumer + # that only links libmdl must pull fmt itself. + target_link_libraries(${PROJECT_NAME} PRIVATE libmdl fmt::fmt) + target_compile_features(${PROJECT_NAME} PRIVATE cxx_std_17) + + add_test(NAME rtx::MdlEmissionClassifier COMMAND ${PROJECT_NAME}) + # The test SKIPs (not fails) when the MDL SDK shared library is not loadable. + set_tests_properties(rtx::MdlEmissionClassifier PROPERTIES SKIP_RETURN_CODE 77) +endif() diff --git a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp new file mode 100644 index 000000000..b6a64b1cc --- /dev/null +++ b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp @@ -0,0 +1,500 @@ +/* + * Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. + * SPDX-License-Identifier: BSD-3-Clause + */ + +// Device-free unit tests for the MDL emission classifier (ADR 0007). libmdl is +// a standalone static library (no CUDA/OptiX/GPU), so the static IR pass and +// the descriptor fold are exercised on the host against inline .mdl snippets — +// the same compile flow the device runs in +// MaterialRegistry::acquireMaterialFromCode, minus the device. Only the MDL SDK +// shared library must be discoverable at runtime; if it is not, the test SKIPs +// (return code 77, wired in CMake). +// +// runIR covers the owned IR (topology, semantics, deps, shared temporaries); +// runFold covers the three-valued descriptor fold (verdict, edfKinds, +// magnitude, mode, sign, geometric-state dependence) against a fake value +// source. + +#include "libmdl/Core.h" +#include "libmdl/EmissionDescriptor.h" +#include "libmdl/EmissionFold.h" +#include "libmdl/EmissionIR.h" +#include "libmdl/source_name_utils.h" + +#include + +#include +#include +#include +#include +#include +#include +#include + +using visrtx::libmdl::Core; +using visrtx::libmdl::makeInlineModuleName; + +namespace { + +int g_failures = 0; + +#define CHECK(cond) \ + do { \ + if (!(cond)) { \ + std::printf("FAIL %s:%d %s\n", __FILE__, __LINE__, #cond); \ + ++g_failures; \ + } \ + } while (0) + +constexpr int kSkipReturnCode = 77; + +bool approxEqual(float a, float b, float tol = 1e-4f) +{ + return std::fabs(a - b) <= tol * std::max(1.0f, std::fabs(b)); +} + +// Compile `source`'s `material emissive(...)`, keeping it alive via `keepAlive` +// (the IR builder and fold need the compiled material live during the call). +void compileMaterial(Core &core, + mi::neuraylib::ITransaction *txn, + std::string_view source, + mi::base::Handle &keepAlive) +{ + using mi::base::make_handle; + const auto moduleName = makeInlineModuleName(source); + auto module = make_handle(core.loadModuleFromString(moduleName, source, txn)); + if (!module.is_valid_interface()) { + std::printf("FAIL could not load inline module\n"); + ++g_failures; + return; + } + auto fnDef = + make_handle(core.getFunctionDefinition(module.get(), "emissive", txn)); + if (!fnDef.is_valid_interface()) { + std::printf("FAIL could not find material 'emissive'\n"); + ++g_failures; + return; + } + keepAlive = make_handle(core.getCompiledMaterial(fnDef.get())); + if (!keepAlive.is_valid_interface()) { + std::printf("FAIL could not compile material\n"); + ++g_failures; + } +} + +// intensity = color(2.0) * math::PI folds to a body-literal constant: emitted +// radiance = intensity / PI = 2.0 per channel. +const char *kConstLiteral = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(2.0) * math::PI))); +)mdl"; + +// A parameter with a default stays symbolic under class compilation, so the +// intensity does not fold; the single-factor walk records a Parameter recipe +// with scale = PI * (1/PI) = 1. +const char *kParamDriven = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive(color value = color(3.0)) = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: value * math::PI))); +)mdl"; + +// Textured intensity: tex::lookup_color(tex) with tex a symbolic parameter → +// Texture recipe. +const char *kTextured = R"mdl(mdl 1.6; +import ::df::*; +import ::tex::*; +export material emissive(uniform texture_2d tex = texture_2d()) = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: tex::lookup_color(tex: tex, coord: float2(0.0))))); +)mdl"; + +// No emission: emission defaults to edf() (a constant invalid-df), not a +// df::diffuse_edf direct call. +const char *kNoEmission = R"mdl(mdl 1.6; +import ::df::*; +export material emissive() = material( + surface: material_surface( + scattering: df::diffuse_reflection_bsdf())); +)mdl"; + +// Diffuse EDF but power intensity mode — the classifier only handles +// radiant-exitance, so this is rejected. +const char *kPowerMode = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(2.0) * math::PI, + mode: intensity_power))); +)mdl"; + +// A let-shared subexpression: `k` is referenced twice, so class compilation +// stores it as a single temporary. The IR must resolve both references to the +// same node. +const char *kSharedTemporary = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive(color value = color(2.0)) = let { + color k = value * math::PI; +} in material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: k + k))); +)mdl"; + +using visrtx::libmdl::buildEmissionIR; +using visrtx::libmdl::ConstantKind; +using visrtx::libmdl::EmissionIR; +using visrtx::libmdl::EmissionNodeKind; +using Semantic = visrtx::libmdl::Semantic; + +const visrtx::libmdl::EmissionNode &node(const EmissionIR &ir, int i) +{ + return ir.nodes[std::size_t(i)]; +} + +bool hasParamNamed(const EmissionIR &ir, const char *name) +{ + for (const auto &n : ir.nodes) + if (n.parameterName == name) + return true; + return false; +} + +void runIR(Core &core, mi::neuraylib::ITransaction *txn) +{ + using mi::base::make_handle; + mi::base::Handle keepAlive; + + auto buildIR = [&](std::string_view src) { + compileMaterial(core, txn, src, keepAlive); // reuse the compile flow + return buildEmissionIR(keepAlive.get(), txn); + }; + + { + auto ir = buildIR(kConstLiteral); + CHECK(!ir.empty()); + CHECK(ir.surface.edfRoot >= 0); + CHECK(node(ir, ir.surface.edfRoot).kind == EmissionNodeKind::Call); + CHECK(node(ir, ir.surface.edfRoot).semantic + == Semantic::DS_INTRINSIC_DF_DIFFUSE_EDF); + CHECK(ir.surface.intensityRoot >= 0); + CHECK(ir.emissionDeps.empty()); // body-literal: no argument deps + } + + { + auto ir = buildIR(kParamDriven); + CHECK(node(ir, ir.surface.edfRoot).semantic + == Semantic::DS_INTRINSIC_DF_DIFFUSE_EDF); + CHECK(!ir.emissionDeps.empty()); + CHECK(hasParamNamed(ir, "value")); + } + + { + auto ir = buildIR(kTextured); + bool foundTexture = false; + for (const auto &n : ir.nodes) { + if (n.kind == EmissionNodeKind::Texture && n.parameterName == "tex") + foundTexture = true; + } + CHECK(foundTexture); + CHECK(hasParamNamed(ir, "tex")); + CHECK(!ir.emissionDeps.empty()); + } + + { + auto ir = buildIR(kNoEmission); + CHECK(ir.surface.edfRoot >= 0); + CHECK(node(ir, ir.surface.edfRoot).kind == EmissionNodeKind::Constant); + CHECK(node(ir, ir.surface.edfRoot).constantKind == ConstantKind::InvalidDf); + } + + { + // k + k: the two operands of the addition must resolve to the SAME node + // index (shared temporary), proving CSE-identity is preserved. + auto ir = buildIR(kSharedTemporary); + const auto &intensity = node(ir, ir.surface.intensityRoot); + CHECK(intensity.kind == EmissionNodeKind::Call); + CHECK(intensity.operands.size() == 2); + if (intensity.operands.size() == 2) + CHECK(intensity.operands[0] == intensity.operands[1]); + } +} + +// Purely negative constant emission: nonzero (an emitter) but sign is Unknown, +// so the policy must not register it (its all-negative NEE term would be +// dropped while the forward deposit is MIS-downweighted). +const char *kNegativeConst = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(-2.0) * math::PI))); +)mdl"; + +// Intensity reads state::normal (via length) — a geometric-state quantity the +// synthetic hit fabricates. Must set dependsOnGeometricState. +const char *kStateNormal = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +import ::state::*; +export material emissive() = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: color(math::length(state::normal())) * math::PI))); +)mdl"; + +// A textured lookup whose COORD reads state::position (not texture_coordinate) +// is unfaithful at the synthetic hit — must set dependsOnGeometricState even +// though a texture_coordinate-driven lookup does not. +const char *kTexturedStateCoord = R"mdl(mdl 1.6; +import ::df::*; +import ::tex::*; +import ::state::*; +export material emissive(uniform texture_2d tex = texture_2d()) = material( + surface: material_surface( + emission: material_emission( + emission: df::diffuse_edf(), + intensity: tex::lookup_color( + tex: tex, + coord: float2(state::position().x, state::position().y))))); +)mdl"; + +using visrtx::libmdl::EmissionSign; +using visrtx::libmdl::EmissionValueSource; +using visrtx::libmdl::EmissionVerdict; +using visrtx::libmdl::foldEmissionDescriptor; +using visrtx::libmdl::IntensityMode; +using visrtx::libmdl::NullValueSource; +using visrtx::libmdl::ResourceStats; + +// Map-backed value source for the fold vectors, keyed by parameter name. +struct MapValueSource : EmissionValueSource +{ + std::map> colors; + std::map resByParam; + + bool color(const std::string &name, std::array &o) const override + { + auto it = colors.find(name); + if (it == colors.end()) + return false; + o = it->second; + return true; + } + bool boolean(const std::string &, bool &) const override + { + return false; + } + bool resourceByName(const std::string &, ResourceStats &) const override + { + return false; + } + bool resourceByParam(const std::string &name, ResourceStats &o) const override + { + auto it = resByParam.find(name); + if (it == resByParam.end()) + return false; + o = it->second; + return true; + } +}; + +// A normalized_mix of two diffuse EDFs. The EDF leaves nest inside the +// df_component[] array, so the fold must deep-scan to see them: the mix must +// NOT fold to ProvablyNull, and its kind must be Diffuse. +const char *kMixDiffuse = R"mdl(mdl 1.6; +import ::df::*; +import ::math::*; +export material emissive() = material( + surface: material_surface( + emission: material_emission( + emission: df::normalized_mix( + df::edf_component[]( + df::edf_component(0.5, df::diffuse_edf()), + df::edf_component(0.5, df::diffuse_edf()))), + intensity: color(2.0) * math::PI))); +)mdl"; + +void runFold(Core &core, mi::neuraylib::ITransaction *txn) +{ + mi::base::Handle keepAlive; + NullValueSource none; + + auto irOf = [&](std::string_view src) { + compileMaterial(core, txn, src, keepAlive); + return buildEmissionIR(keepAlive.get(), txn); + }; + + { // constant literal: radiance = 2.0 per channel + auto ir = irOf(kConstLiteral); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyEmissive); + CHECK(hasKind(d.edfKinds, visrtx::libmdl::EdfKind::Diffuse)); + CHECK(d.mode == IntensityMode::RadiantExitance); + CHECK(d.sign == EmissionSign::ProvablyNonnegative); + CHECK(!d.dependsOnGeometricState); + CHECK(approxEqual(d.magnitude[0], 2.0f)); + } + + { // parameter with a known live value ⇒ magnitude tracks it + auto ir = irOf(kParamDriven); + MapValueSource vs; + vs.colors[std::string("value")] = {3.0f, 3.0f, 3.0f}; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyEmissive); + CHECK(d.sign == EmissionSign::ProvablyNonnegative); + CHECK(approxEqual(d.magnitude[0], 3.0f)); + } + + { // parameter with unknown value ⇒ Unknown verdict, unit-proxy magnitude + auto ir = irOf(kParamDriven); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.verdict == EmissionVerdict::Unknown); + CHECK(approxEqual(d.magnitude[0], 1.0f)); + } + + { // no emission ⇒ ProvablyNull + auto ir = irOf(kNoEmission); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyNull); + } + + { // power mode ⇒ described, mode Power + auto ir = irOf(kPowerMode); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.mode == IntensityMode::Power); + } + + { // bound texture, nonzero mean ⇒ Unknown verdict, magnitude from + // meanPositive + auto ir = irOf(kTextured); + MapValueSource vs; + ResourceStats s; + s.valid = true; + s.maxAbs = {5.0f, 5.0f, 5.0f}; + s.meanPositive = {4.0f, 4.0f, 4.0f}; + s.minValue = {0.0f, 0.0f, 0.0f}; + s.transferPreservesZero = true; + vs.resByParam[std::string("tex")] = s; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.verdict == EmissionVerdict::Unknown); + CHECK(d.sign == EmissionSign::ProvablyNonnegative); + CHECK(!d.dependsOnGeometricState); + CHECK(approxEqual(d.magnitude[0], 4.0f * 0.31830988f)); + } + + { // all-black bound texture (maxAbs 0, T(0)=0) ⇒ ProvablyNull + auto ir = irOf(kTextured); + MapValueSource vs; + ResourceStats s; + s.valid = true; + s.maxAbs = {0.0f, 0.0f, 0.0f}; + s.meanPositive = {0.0f, 0.0f, 0.0f}; + s.minValue = {0.0f, 0.0f, 0.0f}; + s.transferPreservesZero = true; + vs.resByParam[std::string("tex")] = s; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyNull); + } + + { // unbound texture ⇒ lookup folds to 0 ⇒ ProvablyNull + auto ir = irOf(kTextured); + MapValueSource vs; + ResourceStats s; + s.valid = false; + vs.resByParam[std::string("tex")] = s; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyNull); + } + + { // negative constant: emitter but sign Unknown ⇒ not registerable + auto ir = irOf(kNegativeConst); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.verdict == EmissionVerdict::ProvablyEmissive); + CHECK(d.sign == EmissionSign::Unknown); + CHECK(approxEqual(d.magnitude[0], 0.0f)); // meanPositive of a negative is 0 + } + + { // intensity reads state::normal ⇒ dependsOnGeometricState + auto ir = irOf(kStateNormal); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.dependsOnGeometricState); + } + + { // texture coord reads state::position ⇒ dependsOnGeometricState + auto ir = irOf(kTexturedStateCoord); + MapValueSource vs; + ResourceStats s; + s.valid = true; + s.maxAbs = {5.0f, 5.0f, 5.0f}; + s.meanPositive = {4.0f, 4.0f, 4.0f}; + s.minValue = {0.0f, 0.0f, 0.0f}; + vs.resByParam[std::string("tex")] = s; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.dependsOnGeometricState); + } + + { // normalized_mix of diffuse EDFs: leaves nest in the component array, so a + // shallow scan would wrongly prove Null. Must stay emissive + Diffuse. + auto ir = irOf(kMixDiffuse); + auto d = foldEmissionDescriptor(ir, none).surface; + CHECK(d.verdict != EmissionVerdict::ProvablyNull); + CHECK(hasKind(d.edfKinds, visrtx::libmdl::EdfKind::Diffuse)); + } + + { // non-finite live argument value ⇒ sign fails closed (never registers an + // inf/NaN Pick Power); the light still arrives via the forward path. + auto ir = irOf(kParamDriven); + MapValueSource vs; + const float inf = std::numeric_limits::infinity(); + vs.colors[std::string("value")] = {inf, inf, inf}; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.sign == EmissionSign::Unknown); + } +} + +} // namespace + +int main() +{ + using mi::base::make_handle; + + try { + Core core; + auto scope = core.createScope("MdlEmissionClassifierTestScope"); + auto txn = make_handle(core.createTransaction(scope)); + runIR(core, txn.get()); + runFold(core, txn.get()); + txn->commit(); + core.removeScope(scope); + } catch (const std::exception &e) { + std::printf("SKIP MDL SDK unavailable: %s\n", e.what()); + return kSkipReturnCode; + } + + if (g_failures) { + std::printf("%d check(s) failed\n", g_failures); + return 1; + } + std::printf("all checks passed\n"); + return 0; +} diff --git a/devices/rtx/device/CMakeLists.txt b/devices/rtx/device/CMakeLists.txt index a6472dc91..118d45ec9 100644 --- a/devices/rtx/device/CMakeLists.txt +++ b/devices/rtx/device/CMakeLists.txt @@ -187,6 +187,7 @@ set(SOURCES sampler/Image3D.cpp sampler/PrimitiveSampler.cpp sampler/Sampler.cpp + sampler/TextureStats.cu sampler/TransformSampler.cpp sampler/UnknownSampler.cpp diff --git a/devices/rtx/device/gpu/gpu_objects.h b/devices/rtx/device/gpu/gpu_objects.h index 5d3c42ea5..88a5dc17c 100644 --- a/devices/rtx/device/gpu/gpu_objects.h +++ b/devices/rtx/device/gpu/gpu_objects.h @@ -507,8 +507,8 @@ struct MaterialGPUData bool emissionIsConstant{false}; // Emission is not provably zero (constant, sampler, or attribute bound). The - // hit-side Geometry Light MIS gate; kept in sync by the material's own commit, - // so it is never stale. + // hit-side Geometry Light MIS gate; kept in sync by the material's own + // commit, so it is never stale. bool emissionIsSampleable{false}; // Mean emitted radiance, sizing the Geometry Light Pick Power on both the @@ -808,8 +808,8 @@ struct InstanceLightGPUData mat4 xfm; // Transform for this light instance // For a Geometry Light: index into WorldGPUData::surfaceInstances[] of the // surface instance it was synthesized from, so the NEE sampler can evaluate - // emission against the REAL instance (instance-uniform attributes resolve like - // the path-hit deposit). -1 for authored/HDRI lights. + // emission against the REAL instance (instance-uniform attributes resolve + // like the path-hit deposit). -1 for authored/HDRI lights. DeviceObjectIndex surfaceInstanceIndex = -1; }; @@ -835,7 +835,10 @@ struct WorldGPUData // Normalized cumulative Pick Power over lightInstances (length // numLightInstances, last entry == 1); pick a slot with inverseSampleCDF. The // slot's discrete pick probability is lightPickCdf[i]-lightPickCdf[i-1]. - const float *lightPickCdf; + // Double so a dim light's mass (below float epsilon of the total) is + // preserved as a nonzero interval — else it is unselectable while the + // hit-side pNee is still positive, biasing the deposit's MIS weight. + const double *lightPickCdf; float totalLightPower; // sum of un-normalized instance Pick Powers float hdriPower; // subset sum over HDRI instances (env-MIS pick probability) float sceneRadius; // bounding-sphere radius, for the ambient term's power diff --git a/devices/rtx/device/gpu/lightPickPower.h b/devices/rtx/device/gpu/lightPickPower.h index 875ae530b..9b3d2fd7c 100644 --- a/devices/rtx/device/gpu/lightPickPower.h +++ b/devices/rtx/device/gpu/lightPickPower.h @@ -76,8 +76,7 @@ VISRTX_HOST_DEVICE float lightPickPower( * sceneCrossSection; case LightType::POINT: // Isotropic point light: total flux = 4π · intensity. - return detail::pickLuminance(ld.color) * ld.point.intensity * 2.0f - * kTwoPi; + return detail::pickLuminance(ld.color) * ld.point.intensity * 2.0f * kTwoPi; case LightType::SPHERE: { // Lambertian sphere: flux = L · area · π, area = 4πr². const float area = kTwoPi * 2.0f * ld.sphere.radius * ld.sphere.radius @@ -85,9 +84,10 @@ VISRTX_HOST_DEVICE float lightPickPower( return detail::pickLuminance(ld.color) * ld.sphere.intensity * area * kPi; } case LightType::RECT: { - // Lambertian rectangle: flux = L · area · π, doubled if it emits both sides. - const float area = detail::affineAreaScale(xfm) - / glm::max(ld.rect.oneOverArea, 1e-8f); + // Lambertian rectangle: flux = L · area · π, doubled if it emits both + // sides. + const float area = + detail::affineAreaScale(xfm) / glm::max(ld.rect.oneOverArea, 1e-8f); const float sides = float(ld.rect.side.front + ld.rect.side.back); return detail::pickLuminance(ld.color) * ld.rect.intensity * area * kPi * sides; @@ -99,8 +99,8 @@ VISRTX_HOST_DEVICE float lightPickPower( } case LightType::RING: { // Lambertian disk annulus; the cone falloff is ignored for the estimate. - const float area = detail::affineAreaScale(xfm) - / glm::max(ld.ring.oneOverArea, 1e-8f); + const float area = + detail::affineAreaScale(xfm) / glm::max(ld.ring.oneOverArea, 1e-8f); return detail::pickLuminance(ld.color) * ld.ring.intensity * area * kPi; } case LightType::HDRI: @@ -109,6 +109,8 @@ VISRTX_HOST_DEVICE float lightPickPower( return detail::pickLuminance(ld.color) * ld.hdri.scale * sceneCrossSection; case LightType::GEOMETRY: { // Double-sided Lambertian surface: flux = L · area · π, doubled for sides. + // This diffuse-only assumption is in lockstep with kFaithfulSet + // (material/EmissionPolicy.h); growing that set requires generalizing this. const float area = ld.geometry.area * detail::affineAreaScale(xfm); return detail::pickLuminance(ld.geometry.radiance) * area * kPi * 2.0f; } diff --git a/devices/rtx/device/gpu/sampleLight.h b/devices/rtx/device/gpu/sampleLight.h index 09f602c4b..2b767199b 100644 --- a/devices/rtx/device/gpu/sampleLight.h +++ b/devices/rtx/device/gpu/sampleLight.h @@ -160,8 +160,7 @@ VISRTX_DEVICE LightSample sampleSphereLight( // to account for the transform's effect on surface area (determinant of // jacobian) Currently assumes uniform scaling or no scaling of the light // geometry - float areaPdf = - 1.f / (4.f * kPi * ld.sphere.radius * ld.sphere.radius); + float areaPdf = 1.f / (4.f * kPi * ld.sphere.radius * ld.sphere.radius); ls.pdf = areaPdf * pow2(ls.dist) / cosTheta; } else { // Back-facing surface element contributes no light @@ -326,11 +325,14 @@ VISRTX_DEVICE LightSample sampleSpotLight( return ls; } -VISRTX_DEVICE int inverseSampleCDF(const float *cdf, int size, float u) +// Binary search for the first index i such that cdf[i] >= u (cub::LowerBound): +// inverse transform sampling of a discrete cumulative distribution. Templated +// on the CDF element type so the light-pick CDF (double, to preserve dim-light +// masses below float epsilon) and the float HDRI/primitive CDFs share one path. +template +VISRTX_DEVICE int inverseSampleCDF(const T *cdf, int size, float u) { - // Binary search for the first index i such that cdf[i] >= u (cub::LowerBound): - // inverse transform sampling of a discrete distribution over a cumulative CDF. - return cub::LowerBound(cdf, size, u); + return cub::LowerBound(cdf, size, T(u)); } // The three object-space vertex indices of triangle primID, indexed or soup. @@ -342,10 +344,11 @@ VISRTX_DEVICE uvec3 triangleIndices( // Solid-angle pdf of uniform-in-object-area sampling of a Geometry Light, at a // point on a triangle whose object/world twice-areas and total object area are -// given. Uniform-in-object-area maps to world density (1/A_obj_total)·(A_obj_tri -// /A_world_tri), then to solid angle by dist²/|cosθ|. Exact under any affine -// instance transform. Shared by the sampler and the hit-side MIS pdf so the two -// can never drift — MIS unbiasedness depends on them being identical. +// given. Uniform-in-object-area maps to world density +// (1/A_obj_total)·(A_obj_tri /A_world_tri), then to solid angle by +// dist²/|cosθ|. Exact under any affine instance transform. Shared by the +// sampler and the hit-side MIS pdf so the two can never drift — MIS +// unbiasedness depends on them being identical. VISRTX_DEVICE float geometryLightSolidAnglePdf(float objTwiceArea, float worldTwiceArea, float totalObjArea, @@ -362,9 +365,9 @@ VISRTX_DEVICE float geometryLightSolidAnglePdf(float objTwiceArea, // entry point at a SYNTHETIC hit, so next-event radiance matches the path-hit // deposit exactly (MIS stays unbiased). `uvw` is the geometry's parametric // coordinate at the sample (see the per-geometry samplers). The synthetic hit -// points at the REAL surface instance (surfaceInstanceIndex) so instance-uniform -// attribute emission resolves identically to the deposit; a transform-only -// fallback covers the (unexpected) -1 case. +// points at the REAL surface instance (surfaceInstanceIndex) so +// instance-uniform attribute emission resolves identically to the deposit; a +// transform-only fallback covers the (unexpected) -1 case. VISRTX_DEVICE vec3 evalGeometryLightEmission(ScreenSample &ss, const LightGPUData &ld, const mat4 &xfm, @@ -392,6 +395,10 @@ VISRTX_DEVICE vec3 evalGeometryLightEmission(ScreenSample &ss, // toward the receiver so the EDF evaluates on the sampled side. A no-op for // the analytic samplers (outward normal, far side culled) and for // orientation-independent native-PBR emission. + // The geometric normal reused as shading normal, a synthesized tangent basis, + // and object id 0 make this hit faithful only for orientation-/tangent-/ + // object-id-independent emission — the diffuse case kFaithfulSet admits + // (material/EmissionPolicy.h). Enriching this hit is what grows that set. const vec3 ns = dot(nsWorld, outgoingDir) < 0.0f ? -nsWorld : nsWorld; hit.Ng = hit.Ns = ns; const mat3 basis = computeOrthonormalBasis(ns); @@ -408,15 +415,17 @@ VISRTX_DEVICE vec3 evalGeometryLightEmission(ScreenSample &ss, // path is the live one. const auto &world = ss.frameData->world; if (surfaceInstanceIndex >= 0 - && static_cast(surfaceInstanceIndex) < world.numSurfaceInstances) { + && static_cast(surfaceInstanceIndex) + < world.numSurfaceInstances) { hit.instance = &world.surfaceInstances[surfaceInstanceIndex]; return evaluateSurfaceEmission(*ss.frameData, md, hit, outgoingDir); } // Defensive fallback, unreachable for real geometry lights: a transform-only // instance. An out-of-range index means the host/device surface-instance - // layout drifted (the World-side assert is stripped under NDEBUG), so bounding - // it here keeps the drift from dereferencing out of bounds on device. + // layout drifted (the World-side assert is stripped under NDEBUG), so + // bounding it here keeps the drift from dereferencing out of bounds on + // device. InstanceSurfaceGPUData fallback{}; fallback.objectToWorld = glm::transpose(mat4x3(xfm)); fallback.worldToObject = glm::transpose(mat4x3(glm::affineInverse(xfm))); @@ -449,8 +458,9 @@ VISRTX_DEVICE LightSample sampleTriangleGeometryLight(const LightGPUData &ld, const vec3 e1o = tri.vertices[idx.y] - v0; const vec3 e2o = tri.vertices[idx.z] - v0; - // Uniform barycentric sample of the triangle; b0/b1/b2 weight v0/v1/v2 and are - // the hit's uvw for attribute/texcoord interpolation at the sampled point. + // Uniform barycentric sample of the triangle; b0/b1/b2 weight v0/v1/v2 and + // are the hit's uvw for attribute/texcoord interpolation at the sampled + // point. const float su = sqrtf(pcg_uniform(&ss.rs)); const float u2 = pcg_uniform(&ss.rs); const float b1 = su * (1.0f - u2); @@ -490,14 +500,15 @@ VISRTX_DEVICE LightSample sampleTriangleGeometryLight(const LightGPUData &ld, return ls; } -// Finish a SINGLE-sided (outward) Geometry Light area sample from an object-space -// surface point and its OUTWARD object normal: world direction/distance, the -// emitted radiance at the point, and the EXACT affine solid-angle pdf. The world -// area element and normal come from two transformed orthonormal object tangents -// (|cross(M t1, M t2)|), so it is exact under any affine instance transform — the -// same Jacobian the triangle path uses. Shared by the sphere/cylinder/cone -// samplers; the pdf must match geometryLightHitPdf on the deposit side for MIS to -// partition to 1. Radiance/pdf are left zero when the sample faces away. +// Finish a SINGLE-sided (outward) Geometry Light area sample from an +// object-space surface point and its OUTWARD object normal: world +// direction/distance, the emitted radiance at the point, and the EXACT affine +// solid-angle pdf. The world area element and normal come from two transformed +// orthonormal object tangents +// (|cross(M t1, M t2)|), so it is exact under any affine instance transform — +// the same Jacobian the triangle path uses. Shared by the sphere/cylinder/cone +// samplers; the pdf must match geometryLightHitPdf on the deposit side for MIS +// to partition to 1. Radiance/pdf are left zero when the sample faces away. VISRTX_DEVICE LightSample finishAreaLightSample(ScreenSample &ss, const LightGPUData &ld, const mat4 &xfm, @@ -516,7 +527,8 @@ VISRTX_DEVICE LightSample finishAreaLightSample(ScreenSample &ss, const mat3 basis = computeOrthonormalBasis(nObjOut); vec3 nWorld = cross(xfmVec(xfm, basis[0]), xfmVec(xfm, basis[1])); - const float worldAreaScale = length(nWorld); // world-area per unit object-area + const float worldAreaScale = + length(nWorld); // world-area per unit object-area if (worldAreaScale <= 0.0f) return ls; nWorld /= worldAreaScale; @@ -547,8 +559,9 @@ VISRTX_DEVICE LightSample finishAreaLightSample(ScreenSample &ss, } // Sample a point on a sphere-set Geometry Light. Picks a sphere by its -// object-space area (4πr²), samples that sphere's surface uniformly (Marsaglia). -// SINGLE-sided (outward): a closed sphere self-occludes its far hemisphere. +// object-space area (4πr²), samples that sphere's surface uniformly +// (Marsaglia). SINGLE-sided (outward): a closed sphere self-occludes its far +// hemisphere. VISRTX_DEVICE LightSample sampleSphereGeometryLight(const LightGPUData &ld, const SphereGeometryData &sph, const mat4 &xfm, @@ -574,10 +587,17 @@ VISRTX_DEVICE LightSample sampleSphereGeometryLight(const LightGPUData &ld, const float phi = kTwoPi * pcg_uniform(&ss.rs); const vec3 nObj = vec3(rho * cosf(phi), rho * sinf(phi), z); - // Sphere attributes are per-primitive (no interpolation); uvw = (0,0,1) matches - // the intersector's constant sphere parameter. - return finishAreaLightSample(ss, ld, xfm, origin, c + nObj * r, nObj, primID, - vec3(0.0f, 0.0f, 1.0f), surfaceInstanceIndex); + // Sphere attributes are per-primitive (no interpolation); uvw = (0,0,1) + // matches the intersector's constant sphere parameter. + return finishAreaLightSample(ss, + ld, + xfm, + origin, + c + nObj * r, + nObj, + primID, + vec3(0.0f, 0.0f, 1.0f), + surfaceInstanceIndex); } // Per-endpoint cap enablement, matching the intersector's resolveCapBits: a @@ -588,8 +608,10 @@ VISRTX_DEVICE void resolveEndpointCaps(const uint8_t *vertexCaps, bool &cap0, bool &cap1) { - cap0 = vertexCaps ? (vertexCaps[idx.x] != 0) : bool(defaultCapFlags & CAP_FIRST); - cap1 = vertexCaps ? (vertexCaps[idx.y] != 0) : bool(defaultCapFlags & CAP_SECOND); + cap0 = + vertexCaps ? (vertexCaps[idx.x] != 0) : bool(defaultCapFlags & CAP_FIRST); + cap1 = vertexCaps ? (vertexCaps[idx.y] != 0) + : bool(defaultCapFlags & CAP_SECOND); } // Uniform point on a disk of radius `rad` in the plane spanned by (e0,e1) @@ -602,10 +624,10 @@ VISRTX_DEVICE vec3 sampleDisk( return c + rr * (cosf(phi) * e0 + sinf(phi) * e1); } -// Sample a point on a cylinder-set Geometry Light: pick a cylinder by object area -// (lateral 2πrL + enabled caps πr²), then pick lateral vs a cap by sub-area and -// sample it uniformly. Outward object normal is radial on the wall, ±axis on a -// cap. Single-sided (outward). +// Sample a point on a cylinder-set Geometry Light: pick a cylinder by object +// area (lateral 2πrL + enabled caps πr²), then pick lateral vs a cap by +// sub-area and sample it uniformly. Outward object normal is radial on the +// wall, ±axis on a cap. Single-sided (outward). VISRTX_DEVICE LightSample sampleCylinderGeometryLight(const LightGPUData &ld, const CylinderGeometryData &cyl, const mat4 &xfm, @@ -620,7 +642,8 @@ VISRTX_DEVICE LightSample sampleCylinderGeometryLight(const LightGPUData &ld, uint32_t(inverseSampleCDF( cyl.primAreaCdf, int(cyl.numPrimitives), pcg_uniform(&ss.rs))), cyl.numPrimitives - 1); - const uvec2 idx = cyl.indices ? cyl.indices[primID] : uvec2(0, 1) + primID * 2; + const uvec2 idx = + cyl.indices ? cyl.indices[primID] : uvec2(0, 1) + primID * 2; const vec3 p0 = cyl.vertices[idx.x]; const vec3 p1 = cyl.vertices[idx.y]; const float r = fabsf(cyl.radii ? cyl.radii[primID] : cyl.radius); @@ -629,7 +652,8 @@ VISRTX_DEVICE LightSample sampleCylinderGeometryLight(const LightGPUData &ld, if (len <= 0.0f || r <= 0.0f) return {}; const vec3 axisN = axis / len; - const mat3 basis = computeOrthonormalBasis(axisN); // basis[0],basis[1] ⊥ axisN + const mat3 basis = + computeOrthonormalBasis(axisN); // basis[0],basis[1] ⊥ axisN bool cap0, cap1; resolveEndpointCaps(cyl.vertexCaps, cyl.defaultCapFlags, idx, cap0, cap1); @@ -657,14 +681,22 @@ VISRTX_DEVICE LightSample sampleCylinderGeometryLight(const LightGPUData &ld, } // uvw = (0, u, 1-u): u weights endpoint p1, (1-u) weights p0 (see the // intersector and readAttributeValue's cylinder branch). - return finishAreaLightSample(ss, ld, xfm, origin, pObj, nObj, primID, - vec3(0.0f, axialU, 1.0f - axialU), surfaceInstanceIndex); + return finishAreaLightSample(ss, + ld, + xfm, + origin, + pObj, + nObj, + primID, + vec3(0.0f, axialU, 1.0f - axialU), + surfaceInstanceIndex); } -// Sample a point on a cone-set Geometry Light: pick a cone by object area (frustum -// lateral π(r0+r1)·slant + enabled caps πr²), then pick lateral vs a cap. Lateral -// uses the radius-weighted axial CDF r(t)=√((1-u)r0²+u·r1²); the outward object -// normal is the tilted slant normal on the wall, ±axis on a cap. Single-sided. +// Sample a point on a cone-set Geometry Light: pick a cone by object area +// (frustum lateral π(r0+r1)·slant + enabled caps πr²), then pick lateral vs a +// cap. Lateral uses the radius-weighted axial CDF r(t)=√((1-u)r0²+u·r1²); the +// outward object normal is the tilted slant normal on the wall, ±axis on a cap. +// Single-sided. VISRTX_DEVICE LightSample sampleConeGeometryLight(const LightGPUData &ld, const ConeGeometryData &cone, const mat4 &xfm, @@ -722,8 +754,15 @@ VISRTX_DEVICE LightSample sampleConeGeometryLight(const LightGPUData &ld, nObj = axisN; axialT = 1.0f; } - return finishAreaLightSample(ss, ld, xfm, origin, pObj, nObj, primID, - vec3(0.0f, axialT, 1.0f - axialT), surfaceInstanceIndex); + return finishAreaLightSample(ss, + ld, + xfm, + origin, + pObj, + nObj, + primID, + vec3(0.0f, axialT, 1.0f - axialT), + surfaceInstanceIndex); } // Dispatch a Geometry Light sample by the backing geometry's type. The geometry @@ -761,8 +800,7 @@ VISRTX_DEVICE LightSample sampleHDRILight( auto thetaPhi = sphericalCoordsFromDirection(ld.hdri.xfm * dir); // Map spherical coordinates to UV texture coordinates // θ ∈ [0,π] → v ∈ [0,1], φ ∈ [0,2π] → u ∈ [0,1] - auto uv = glm::vec2(thetaPhi.y, thetaPhi.x) - / glm::vec2(kTwoPi, kPi); + auto uv = glm::vec2(thetaPhi.y, thetaPhi.x) / glm::vec2(kTwoPi, kPi); auto radiance = sampleHDRI(ld, uv); // pdf_ω = (L/totalL) · pdfWeight; the equirectangular sinθ jacobian is @@ -786,8 +824,8 @@ VISRTX_DEVICE LightSample sampleHDRILight( // Importance sampling using hierarchical (marginal/conditional) CDF approach // First sample row (y) using marginal CDF, then column (x) using conditional // CDF - auto y = inverseSampleCDF( - ld.hdri.marginalCDF, ld.hdri.size.y, pcg_uniform(&rs)); + auto y = + inverseSampleCDF(ld.hdri.marginalCDF, ld.hdri.size.y, pcg_uniform(&rs)); auto x = inverseSampleCDF(ld.hdri.conditionalCDF + y * ld.hdri.size.x, ld.hdri.size.x, pcg_uniform(&rs)); @@ -853,7 +891,8 @@ VISRTX_DEVICE LightSample sampleLight(ScreenSample &ss, case LightType::HDRI: return detail::sampleHDRILight(ld, xfm, ss.rs); case LightType::GEOMETRY: - return detail::sampleGeometryLight(ld, xfm, origin, ss, surfaceInstanceIndex); + return detail::sampleGeometryLight( + ld, xfm, origin, ss, surfaceInstanceIndex); default: break; } diff --git a/devices/rtx/device/material/EmissionPolicy.h b/devices/rtx/device/material/EmissionPolicy.h new file mode 100644 index 000000000..3fcdc6b41 --- /dev/null +++ b/devices/rtx/device/material/EmissionPolicy.h @@ -0,0 +1,38 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// The renderer-side registration policy (ADR 0007): the single source of truth +// for which described emission slots become next-event-sampled Geometry Lights. +// A slot registers iff it is non-null AND faithfully NEE-evaluable on the +// synthetic hit. +// +// `kFaithfulSet` is the EDF-kind half of "faithful". It is in lockstep with two +// GPU-side sites that encode the same diffuse-only assumption: +// - gpu/lightPickPower.h (double-sided Lambertian flux for GEOMETRY lights) +// - gpu/sampleLight.h (double-sided normal orientation at the synthetic +// hit, geometric normal reused as shading normal) +// Growing `kFaithfulSet` beyond {Diffuse} REQUIRES updating both — see the +// synthetic-hit-fidelity follow-up in ADR 0007. + +#include "libmdl/EmissionDescriptor.h" + +namespace visrtx { + +// EDF kinds the renderer can evaluate faithfully at the fidelity-limited +// next-event synthetic hit (geometric normal, synthesized tangent, object id 0, +// forced front). Today: diffuse only. +constexpr libmdl::EdfKind kFaithfulSet = libmdl::EdfKind::Diffuse; + +// Whether a described emission slot should register as a Geometry Light. +inline bool isRegisterable(const libmdl::SlotDescriptor &slot) +{ + return slot.verdict != libmdl::EmissionVerdict::ProvablyNull + && libmdl::isSubsetOf(slot.edfKinds, kFaithfulSet) + && slot.mode == libmdl::IntensityMode::RadiantExitance + && !slot.dependsOnGeometricState + && slot.sign == libmdl::EmissionSign::ProvablyNonnegative; +} + +} // namespace visrtx diff --git a/devices/rtx/device/material/MDL.cpp b/devices/rtx/device/material/MDL.cpp index 1d8cf7895..d4fb5e3ee 100644 --- a/devices/rtx/device/material/MDL.cpp +++ b/devices/rtx/device/material/MDL.cpp @@ -33,8 +33,12 @@ #include "gpu/gpu_objects.h" #include "gpu/sbt.h" +#include "material/EmissionPolicy.h" + #include "libmdl/ArgumentBlockDescriptor.h" #include "libmdl/ArgumentBlockInstance.h" +#include "libmdl/EmissionFold.h" +#include "libmdl/EmissionIR.h" #include "libmdl/helpers.h" #include "libmdl/source_name_utils.h" #include "material/Material.h" @@ -65,6 +69,58 @@ using namespace std::string_view_literals; namespace visrtx { +namespace { + +// Value source that folds the emission IR against this material instance's live +// arguments and bound samplers. Keyed by class-compilation argument name, which +// is how the argument block and the sampler descriptors are keyed. +struct MDLValueSource : libmdl::EmissionValueSource +{ + const libmdl::ArgumentBlockInstance *argBlock{nullptr}; + std::map samplersByName; + + bool color(const std::string &name, std::array &out) const override + { + if (!argBlock) + return false; + if (auto v = argBlock->getFloat3Value(name)) { + out = {(*v)[0], (*v)[1], (*v)[2]}; + return true; + } + return false; + } + bool boolean(const std::string &, bool &) const override + { + return false; + } + bool resourceByName( + const std::string &, libmdl::ResourceStats &) const override + { + // A module-default (body-literal) texture resolves under its URL, not a + // parameter name, so it is not reachable here — leaves the fold Unknown, + // which is safe (forward-only). Argument-bound textures resolve below. + return false; + } + bool resourceByParam( + const std::string &name, libmdl::ResourceStats &out) const override + { + auto it = samplersByName.find(name); + if (it == samplersByName.end()) { + // No sampler bound to this texture argument ⇒ the MDL lookup is invalid + // and folds to 0 (ADR 0007 A5). An unbound emissive texture emits nothing + // and is not a Geometry Light. + out = libmdl::ResourceStats{}; // valid == false + return true; + } + if (!it->second || !it->second->isValid()) + return false; // bound but not yet resolvable ⇒ Unknown (still registers) + out = it->second->emissionStats(); + return true; + } +}; + +} // namespace + MDL::MDL(DeviceGlobalState *d) : Material(d) {} MDL::~MDL() @@ -111,7 +167,8 @@ void MDL::finalize() syncParameters(); if (m_argumentBlockInstance.has_value()) { - if (const auto &argBlockData = m_argumentBlockInstance->getArgumentBlockData(); + if (const auto &argBlockData = + m_argumentBlockInstance->getArgumentBlockData(); !argBlockData.empty()) { m_argBlockBuffer.upload(data(argBlockData), size(argBlockData)); } else { @@ -121,14 +178,23 @@ void MDL::finalize() m_argBlockBuffer.reset(); } - // The emission hooks read the compile-time classification, keyed by the uuid - // syncSource just resolved. Resolve it BEFORE Material::finalize(), which - // uploads gpuData() — a stale flag there zeroes the hit-side next-event pdf - // and the deposit double-counts the emitter. The light-set refresh follows - // (not in commitParameters, where the flag would be stale on the commit that - // first introduces or removes emission). - m_emissionClassification = - deviceState()->mdl->materialRegistry.getEmissionClassification(m_uuid); + // Fold the compile-time emission IR (keyed by the uuid syncSource just + // resolved) against this instance's live arguments and samplers into the + // emission descriptor. Do it BEFORE Material::finalize(), which uploads + // gpuData() — a stale sampleability flag there zeroes the hit-side next-event + // pdf and the deposit double-counts. The light-set refresh follows (not in + // commitParameters, where the flag would be stale on the commit that first + // introduces or removes emission). + const libmdl::EmissionIR emissionIR = + deviceState()->mdl->materialRegistry.getEmissionIR(m_uuid); + MDLValueSource values; + values.argBlock = + m_argumentBlockInstance ? &*m_argumentBlockInstance : nullptr; + for (const auto &desc : m_samplers) { + if (desc.sampler) + values.samplersByName.emplace(desc.name, desc.sampler); + } + m_emissionDescriptor = libmdl::foldEmissionDescriptor(emissionIR, values); Material::finalize(); @@ -137,74 +203,23 @@ void MDL::finalize() bool MDL::emissionIsSampleable() const { - // Textured/procedural diffuse intensity has no host constant: sampleable, - // the device evaluating the true radiance at the sampled point, with the - // dynamic recipe (or the unit proxy) supplying the Pick Power. Conservative - // over-inclusion (e.g. an all-black texture) is unbiased — it merely wastes - // a pick slot. - const auto &radiance = m_emissionClassification.constantRadiance; - if (m_emissionClassification.isDiffuseEmission && !radiance) - return true; - return radiance - && std::max({(*radiance)[0], (*radiance)[1], (*radiance)[2]}) > 0.0f; + // The surface slot registers as a Geometry Light iff the folded descriptor is + // non-null and faithfully NEE-evaluable (ADR 0007). Textured/Unknown-verdict + // diffuse emission registers with the device evaluating the true radiance at + // the sampled point; an all-black texture folds to ProvablyNull and is + // excluded; signed, geometric-state-dependent, spot/measured, or power-mode + // emission stays forward-only (unbiased) rather than registering + // unfaithfully. + return isRegisterable(m_emissionDescriptor.surface); } vec3 MDL::emissionAverage() const { - const auto &radiance = m_emissionClassification.constantRadiance; - if (radiance) - return vec3((*radiance)[0], (*radiance)[1], (*radiance)[2]); - if (!m_emissionClassification.isDiffuseEmission) - return vec3(0.0f); - - // Dynamic recipe: the classification identified a single argument- or - // texture-driven intensity factor, so the mean radiance follows the LIVE - // argument value (or bound sampler mean) at light-build time — keeping the - // Pick Power true instead of the unit proxy. Under-picking a bright emitter - // is unbiased only in exact arithmetic: the firefly clamp and the last-depth - // MIS truncation both turn the resulting overweighted picks into visible - // dimming next to correctly-powered lights. - using DynamicSource = libmdl::Core::EmissionClassification::DynamicSource; - const auto &cls = m_emissionClassification; - const vec3 scale( - cls.dynamicScale[0], cls.dynamicScale[1], cls.dynamicScale[2]); - switch (cls.dynamicSource) { - case DynamicSource::Parameter: - if (m_argumentBlockInstance) { - if (auto v = - m_argumentBlockInstance->getFloat3Value(cls.dynamicArgumentName)) { - const vec3 mean = - glm::max(vec3((*v)[0], (*v)[1], (*v)[2]), vec3(0.0f)) * scale; - if (std::isfinite(mean.x) && std::isfinite(mean.y) - && std::isfinite(mean.z)) - return mean; - } - } - break; - case DynamicSource::Texture: { - // Known limitation: this resolves ANARI-bound samplers only. A texture - // parameter left at its MODULE DEFAULT registers in m_samplers under its - // URL (syncSource), not its parameter name, so it misses here and keeps - // the unit proxy — safe degrade, tracked as a follow-up. - auto it = std::find_if(cbegin(m_samplers), - cend(m_samplers), - [&](const auto &desc) { return desc.name == cls.dynamicArgumentName; }); - if (it != cend(m_samplers) && it->sampler && it->sampler->isValid()) { - const vec3 mean = - glm::max(vec3(it->sampler->averageValue()), vec3(0.0f)) * scale; - if (std::isfinite(mean.x) && std::isfinite(mean.y) - && std::isfinite(mean.z)) - return mean; - } - break; - } - default: - break; - } - // Unit-luminance proxy when no recipe (or its inputs) resolve: any nonzero - // value stays unbiased on both MIS estimator sides — it only steers - // importance (a power-weighted estimate is a follow-up). - return vec3(1.0f); + // The non-negative meanPositive magnitude proxy (radiance = intensity / PI), + // folded from the live arguments/samplers. Weights the Light Pick only; a + // unit proxy stands in when the intensity magnitude is not host-known. + const auto &m = m_emissionDescriptor.surface.magnitude; + return vec3(m[0], m[1], m[2]); } void MDL::syncSource() @@ -348,7 +363,8 @@ void MDL::syncParameters() const auto &name = param->first; if (name == "source"sv || name == "sourceType"sv || name == "materialName"sv) { - // Skip these control parameters, they are not part of the argument block + // Skip these control parameters, they are not part of the argument + // block continue; } @@ -367,8 +383,8 @@ void MDL::syncParameters() for (auto &&[name, type] : argumentBlockInstance.enumerateArguments()) { auto sourceParamAny = m_parameterMap.find(name) != m_parameterMap.end() - ? m_parameterMap[name] - : helium::AnariAny{}; + ? m_parameterMap[name] + : helium::AnariAny{}; if (sourceParamAny.valid() == 0) { // Parameter not set, reset to default value. @@ -376,7 +392,9 @@ void MDL::syncParameters() // Handle the texture case where we might have resources to cleanup if (type == libmdl::ArgumentBlockDescriptor::ArgumentType::Texture) { - if (auto it = find_if(begin(m_samplers), end(m_samplers), [name = name](auto &p) { return p.name == name; }); + if (auto it = find_if(begin(m_samplers), + end(m_samplers), + [name = name](auto &p) { return p.name == name; }); it != end(m_samplers)) { if (it->sampler) { if (it->isFromRegistry) { @@ -506,9 +524,8 @@ void MDL::syncParameters() // Check if this input if already bound and then release it auto it = std::find_if(begin(m_samplers), end(m_samplers), - [¶mName = name](const SamplerDesc &desc) { - return desc.name == paramName; - }); + [¶mName = name]( + const SamplerDesc &desc) { return desc.name == paramName; }); if (it != end(m_samplers)) { // Found, release if (it->sampler) { @@ -548,7 +565,8 @@ void MDL::syncParameters() void MDL::syncImplementationIndex() { - m_implementationIndex = deviceState()->mdl->materialRegistry.getMaterialImplementationIndex( + m_implementationIndex = + deviceState()->mdl->materialRegistry.getMaterialImplementationIndex( m_uuid); } @@ -560,13 +578,15 @@ MaterialGPUData MDL::gpuData() const retval.emissionIsSampleable = emissionIsSampleable(); retval.emissionAverage = emissionAverage(); - retval.callableBaseIndex = m_implementationIndex == mdl::MaterialRegistry::INVALID_IMPLEMENTATION_INDEX ? - ~0u : - uint32_t(SbtCallableEntryPoints::Last) + m_implementationIndex * uint32_t(SurfaceShaderEntryPoints::Count); + retval.callableBaseIndex = m_implementationIndex + == mdl::MaterialRegistry::INVALID_IMPLEMENTATION_INDEX + ? ~0u + : uint32_t(SbtCallableEntryPoints::Last) + + m_implementationIndex * uint32_t(SurfaceShaderEntryPoints::Count); if (m_argumentBlockInstance.has_value()) { retval.materialData.mdl.numSamplers = - std::min(std::size(retval.materialData.mdl.samplers), size(m_samplers)); + std::min(std::size(retval.materialData.mdl.samplers), size(m_samplers)); std::fill(std::begin(retval.materialData.mdl.samplers), std::end(retval.materialData.mdl.samplers), diff --git a/devices/rtx/device/material/MDL.h b/devices/rtx/device/material/MDL.h index dcdf3ab59..29ba263c5 100644 --- a/devices/rtx/device/material/MDL.h +++ b/devices/rtx/device/material/MDL.h @@ -38,6 +38,7 @@ #include "optix_visrtx.h" #include "libmdl/ArgumentBlockInstance.h" +#include "libmdl/EmissionDescriptor.h" #include "sampler/Sampler.h" #include @@ -53,15 +54,13 @@ struct MDL : public Material void commitParameters() override; void finalize() override; - // A raw `mdl` material whose compile-time classification found diffuse - // emission NOT provably zero is a sampleable Emissive Surface (Geometry - // Light): a nonzero body-literal intensity carries the emitted radiance - // (intensity over PI) as the Pick Power; a textured/procedural intensity — - // not host-knowable — uses a unit-luminance proxy, with the true radiance - // evaluated on the device at the sampled point. A folded zero is provably - // zero: no light. - // emissionIsConstant stays false: the EDF path is always taken and the - // average only weights the Light Pick. Per ADR 0006. + // A raw `mdl` material publishes an emission descriptor folded from its IR + // against its live arguments (ADR 0007). The renderer policy registers the + // surface slot as a Geometry Light iff it is non-null and faithfully + // NEE-evaluable (diffuse, radiant-exitance, non-negative, no geometric-state + // dependence). emissionAverage returns the non-negative meanPositive + // magnitude that weights the Light Pick. emissionIsConstant stays false: the + // EDF path is always taken and the average only weights the pick. bool emissionIsSampleable() const override; vec3 emissionAverage() const override; @@ -83,13 +82,15 @@ struct MDL : public Material std::string m_source; std::string m_sourceType; std::optional m_materialName; - struct SamplerDesc { - Sampler* sampler = nullptr; + struct SamplerDesc + { + Sampler *sampler = nullptr; std::string name; bool isFromRegistry = false; - bool operator==(const SamplerDesc &other) const { - return sampler == other.sampler && name == other.name && - isFromRegistry == other.isFromRegistry; + bool operator==(const SamplerDesc &other) const + { + return sampler == other.sampler && name == other.name + && isFromRegistry == other.isFromRegistry; } }; std::vector m_samplers; @@ -97,8 +98,9 @@ struct MDL : public Material libmdl::Uuid m_uuid{}; mdl::MaterialRegistry::ImplementationIndex m_implementationIndex{}; std::optional m_argumentBlockInstance; - // Read from the registry's compile-time cache at finalize (keyed by m_uuid). - libmdl::Core::EmissionClassification m_emissionClassification; + // Folded at finalize from the registry's compile-time IR (keyed by m_uuid) + // against this instance's live arguments and samplers. + libmdl::EmissionDescriptor m_emissionDescriptor; }; } // namespace visrtx diff --git a/devices/rtx/device/mdl/MaterialRegistry.cpp b/devices/rtx/device/mdl/MaterialRegistry.cpp index e148ac331..232b8794a 100644 --- a/devices/rtx/device/mdl/MaterialRegistry.cpp +++ b/devices/rtx/device/mdl/MaterialRegistry.cpp @@ -129,8 +129,8 @@ MaterialRegistry::compileAndCacheMaterial(const std::string &fullMaterialName, return std::nullopt; } - auto functionDef = make_handle(m_core->getFunctionDefinition( - module, materialName, transaction)); + auto functionDef = make_handle( + m_core->getFunctionDefinition(module, materialName, transaction)); if (!functionDef.is_valid_interface()) { m_core->logMessage(mi::base::MESSAGE_SEVERITY_ERROR, "Cannot find function {} definition in module {}", @@ -289,10 +289,11 @@ MaterialRegistry::compileAndCacheMaterial(const std::string &fullMaterialName, targetIt->ptxBlob = ptxBlob; targetIt->refCount = 1; } - // Classify emission now, while the compiled material is alive — it is not - // retained, and the classification is what lets an Emissive Surface be - // synthesized into a Geometry Light without recompiling (ADR 0006). - targetIt->emission = libmdl::Core::classifyEmission(compiledMaterial.get()); + // Extract the emission IR now, while the compiled material is alive — it is + // not retained. The material folds this IR against its live arguments at + // finalize to publish an emission descriptor (ADR 0007). + targetIt->emission = + libmdl::buildEmissionIR(compiledMaterial.get(), transaction); auto targetIndex = std::distance(std::begin(m_targetCodes), targetIt); @@ -334,10 +335,12 @@ MaterialRegistry::acquireMaterial( transaction->abort(); }); - auto module = - make_handle(m_core->loadModule(moduleName, transaction.get())); - auto material = compileAndCacheMaterial( - fullMaterialName, moduleName, materialName, module.get(), transaction.get()); + auto module = make_handle(m_core->loadModule(moduleName, transaction.get())); + auto material = compileAndCacheMaterial(fullMaterialName, + moduleName, + materialName, + module.get(), + transaction.get()); doCommit = material.has_value(); return material.value_or(AcquiredMaterial{}); } @@ -367,8 +370,11 @@ MaterialRegistry::acquireMaterialFromCode( auto module = make_handle( m_core->loadModuleFromString(moduleName, source, transaction.get())); - auto material = compileAndCacheMaterial( - fullMaterialName, moduleName, materialName, module.get(), transaction.get()); + auto material = compileAndCacheMaterial(fullMaterialName, + moduleName, + materialName, + module.get(), + transaction.get()); doCommit = material.has_value(); return material.value_or(AcquiredMaterial{}); } diff --git a/devices/rtx/device/mdl/MaterialRegistry.h b/devices/rtx/device/mdl/MaterialRegistry.h index cbc7a7267..93d978dc2 100644 --- a/devices/rtx/device/mdl/MaterialRegistry.h +++ b/devices/rtx/device/mdl/MaterialRegistry.h @@ -34,6 +34,7 @@ #include "libmdl/ArgumentBlockDescriptor.h" #include "libmdl/ArgumentBlockInstance.h" #include "libmdl/Core.h" +#include "libmdl/EmissionIR.h" #include "libmdl/TimeStamp.h" #include "libmdl/uuid.h" @@ -107,10 +108,10 @@ class MaterialRegistry } } - // Host-side emission classification of a compiled material (ADR 0006); - // default (non-emissive) when the uuid is unknown. - libmdl::Core::EmissionClassification getEmissionClassification( - const libmdl::Uuid &uuid) const + // Owned emission IR of a compiled material (ADR 0007), extracted while the + // compiled material was alive; empty when the uuid is unknown. The material + // folds it against its live arguments at finalize. + libmdl::EmissionIR getEmissionIR(const libmdl::Uuid &uuid) const { if (auto it = m_uuidToIndex.find(uuid); it != cend(m_uuidToIndex)) return m_targetCodes[it->second].emission; @@ -156,9 +157,9 @@ class MaterialRegistry // mi::base::Handle targetCode; std::vector ptxBlob; int refCount{}; - // Computed at compile time, while the compiled material is alive (it is + // Extracted at compile time, while the compiled material is alive (it is // not retained past compilation); evicted with the slot on release. - libmdl::Core::EmissionClassification emission; + libmdl::EmissionIR emission; }; // Per material PTX blobs. Stored in Sbt order. Sparse structure depending on diff --git a/devices/rtx/device/renderer/Interactive_ptx.cu b/devices/rtx/device/renderer/Interactive_ptx.cu index c98da28a2..173da0b95 100644 --- a/devices/rtx/device/renderer/Interactive_ptx.cu +++ b/devices/rtx/device/renderer/Interactive_ptx.cu @@ -218,6 +218,15 @@ struct InteractiveShadingPolicy auto cosineT = dot(bounceHit.Ns, sampleDir); auto color = materialEvaluateTint(bounceShadingState) * cosineT * rendererParams.ambientColor * rendererParams.ambientIntensity; + + // An emitter reached only by the reflection bounce: deposit its + // emission so it appears in reflections and lights via the forward + // path. Guarded to UNREGISTERED emitters — a registered (sampleable) + // one is already covered by the NEE loop above, so depositing here too + // would double- count it. Matches the ADR 0007 "miss = variance" goal + // for Interactive. + if (!bounceHit.material->emissionIsSampleable) + color += materialEvaluateEmission(bounceShadingState, -bounceRay.dir); contrib += color * nextRay.contributionWeight; } else { vec3 hdri; diff --git a/devices/rtx/device/renderer/Quality_ptx.cu b/devices/rtx/device/renderer/Quality_ptx.cu index caf91046b..75509671a 100644 --- a/devices/rtx/device/renderer/Quality_ptx.cu +++ b/devices/rtx/device/renderer/Quality_ptx.cu @@ -144,7 +144,8 @@ struct SurfaceLightSample // Power (relative flux) of the ambient term, treated as an infinite hemisphere // light so it competes in the same Pick Power currency as the light instances -// (irradiance × scene cross-section, matching lightPickPower's infinite lights). +// (irradiance × scene cross-section, matching lightPickPower's infinite +// lights). VISRTX_DEVICE float ambientPickPower(const FrameGPUData &frameData) { const auto &r = frameData.renderer; @@ -169,13 +170,18 @@ VISRTX_DEVICE size_t pickLightInstance(const WorldGPUData &world, float u) // Discrete probability that pickLightInstance selected `idx`, folded with the // ambient stratum so P(pick) sums to 1 across every pick candidate. The CDF is -// normalized by totalLightPower, so its per-slot delta is power_i/totalLightPower. +// normalized by totalLightPower, so its per-slot delta is +// power_i/totalLightPower. VISRTX_DEVICE float instancePickProbability( const WorldGPUData &world, size_t idx, float totalPower) { - const float lo = idx > 0 ? world.lightPickCdf[idx - 1] : 0.0f; - const float conditional = world.lightPickCdf[idx] - lo; - return conditional * world.totalLightPower / totalPower; + // Double delta: the CDF preserves masses below float epsilon, so the delta of + // adjacent entries must be differenced in double or a dim light's interval + // collapses to zero here even though it is selectable. + const double lo = idx > 0 ? world.lightPickCdf[idx - 1] : 0.0; + const double conditional = world.lightPickCdf[idx] - lo; + return float( + conditional * double(world.totalLightPower) / double(totalPower)); } // Aggregate probability that the Light Pick lands on the HDRI environment. Both @@ -190,10 +196,10 @@ VISRTX_DEVICE float envPickProbability(const FrameGPUData &frameData) return world.hdriPower / totalPower; // All-dark fallback: mirror sampleLights' uniform stratum count exactly // (ambient counted iff it carries Pick Power) so both MIS sides agree. - const size_t numStrata = world.numLightInstances + (ambientPower > 0.0f ? 1 : 0); - return numStrata > 0 - ? float(world.numHdriLightInstances) / float(numStrata) - : 0.0f; + const size_t numStrata = + world.numLightInstances + (ambientPower > 0.0f ? 1 : 0); + return numStrata > 0 ? float(world.numHdriLightInstances) / float(numStrata) + : 0.0f; } // NEE density a Geometry Light would report for a BSDF ray that hit it, for the @@ -206,21 +212,19 @@ VISRTX_DEVICE float envPickProbability(const FrameGPUData &frameData) // native textured emission; for MDL the dynamic-recipe live mean, or the unit // proxy when no recipe resolves), and the hit's own (constant) emission // otherwise. `emission` is the surface's evaluated radiance. -VISRTX_DEVICE float geometryLightHitPdf(const FrameGPUData &frameData, - const SurfaceHit &hit, - const vec3 &rayDir, - const vec3 &emission) +VISRTX_DEVICE float geometryLightHitPdf( + const FrameGPUData &frameData, const SurfaceHit &hit, const vec3 &rayDir) { // Only a sampleable-emissive area-samplable surface is a Geometry Light. The // type guard is load-bearing: GeometryGPUData is a union, so reading `.tri`/ - // `.sphere` on the wrong type (never NEE-sampled) would misread it and wrongly - // down-weight the deposit. + // `.sphere` on the wrong type (never NEE-sampled) would misread it and + // wrongly down-weight the deposit. if (!hit.material->emissionIsSampleable) return 0.0f; // objectToWorld is a row-stored mat3x4 (glm column i = OptiX row i of M), so // mat3(objectToWorld) is Mᵀ; transpose it back to M — the linear map the NEE - // sampler uses via xfmVec — or the area Jacobian is wrong under a non-symmetric - // instance transform (rotation + non-uniform scale). + // sampler uses via xfmVec — or the area Jacobian is wrong under a + // non-symmetric instance transform (rotation + non-uniform scale). const mat3 o2w = transpose(mat3(hit.instance->objectToWorld)); const float cosTheta = fabsf(dot(hit.Ng, rayDir)); if (cosTheta <= 0.0f) @@ -228,7 +232,8 @@ VISRTX_DEVICE float geometryLightHitPdf(const FrameGPUData &frameData, // Per-type: the exact solid-angle pdf the NEE sampler would report for this // hit, plus the object-space total area feeding the pick probability. Kept - // identical to the samplers in sampleLight.h so wNee and wBsdf partition to 1. + // identical to the samplers in sampleLight.h so wNee and wBsdf partition + // to 1. float solidAnglePdf = 0.0f; float totalArea = 0.0f; @@ -247,18 +252,19 @@ VISRTX_DEVICE float geometryLightHitPdf(const FrameGPUData &frameData, length(cross(e1o, e2o)), worldTwice, tri.totalArea, hit.t, cosTheta); totalArea = tri.totalArea; } else { - // Sphere/cylinder/cone samplers are SINGLE-sided (outward): finishAreaLightSample - // culls the far hemisphere (cosTheta <= 0). hit.Ng is ray-oriented so fabsf above - // cannot recover facing — an interior (back-face) hit is never NEE-sampled, so its - // NEE pdf must be 0. Otherwise the deposit sees pNee > 0 and down-weights via MIS - // while NEE contributes nothing, losing the interior fraction (dark shell inside). + // Sphere/cylinder/cone samplers are SINGLE-sided (outward): + // finishAreaLightSample culls the far hemisphere (cosTheta <= 0). hit.Ng is + // ray-oriented so fabsf above cannot recover facing — an interior + // (back-face) hit is never NEE-sampled, so its NEE pdf must be 0. Otherwise + // the deposit sees pNee > 0 and down-weights via MIS while NEE contributes + // nothing, losing the interior fraction (dark shell inside). if (!hit.isFrontFace) return 0.0f; - // The unit-tangent samplers depend only on the OUTWARD object normal at the point - // (worldAreaScale = |cross(M t1,M t2)|, invariant to the tangent basis). Recover it - // generically from the world normal — o2wᵀ·Ng ∝ nObj since Ng = normalize(M⁻ᵀ·nObj) - // — so no per-surface (lateral vs cap, slant) math is needed and it matches - // finishAreaLightSample exactly. + // The unit-tangent samplers depend only on the OUTWARD object normal at the + // point (worldAreaScale = |cross(M t1,M t2)|, invariant to the tangent + // basis). Recover it generically from the world normal — o2wᵀ·Ng ∝ nObj + // since Ng = normalize(M⁻ᵀ·nObj) — so no per-surface (lateral vs cap, + // slant) math is needed and it matches finishAreaLightSample exactly. uint32_t numPrimitives = 0; if (hit.geometry->type == GeometryType::SPHERE) { totalArea = hit.geometry->sphere.totalArea; @@ -291,15 +297,22 @@ VISRTX_DEVICE float geometryLightHitPdf(const FrameGPUData &frameData, LightGPUData ld{}; ld.type = LightType::GEOMETRY; ld.geometry.geometryIndex = -1; // unused by lightPickPower - // A constant emitter's CDF weight is its own radiance (== emissionAverage for - // the native path); a textured emitter's varies per point, so its CDF was - // built from the mean — use that here so pick probabilities match. - ld.geometry.radiance = hit.material->emissionIsConstant - ? emission - : hit.material->emissionAverage; + // Use the exact non-negative magnitude the Geometry Light's CDF pick power + // was built from — emissionAverage — so the hit-side pNee equals the + // selection pick probability. Reading the per-hit (possibly signed, + // per-point-varying) emission would disagree with the CDF and bias the + // deposit's MIS weight. + ld.geometry.radiance = hit.material->emissionAverage; ld.geometry.area = totalArea; + // Apply the SAME (raw > 0 && finite) clamp World::appendLight uses when it + // builds the CDF (World.cpp): a light whose recomputed Pick Power is + // non-finite or non-positive contributes 0 to the host CDF, so its hit-side + // pNee must be 0 too — else wEmission goes NaN/0 and disagrees with + // selection. + const float rawPick = + lightPickPower(ld, mat4(o2w), frameData.world.sceneRadius); const float pickProb = - lightPickPower(ld, mat4(o2w), frameData.world.sceneRadius) / totalPower; + (rawPick > 0.0f && isfinite(rawPick)) ? rawPick / totalPower : 0.0f; return solidAnglePdf * pickProb; } @@ -328,8 +341,8 @@ VISRTX_DEVICE PickedCandidate pickCandidate( const float totalPower = world.totalLightPower + ambientPower; - // Fallback when no candidate carries Pick Power (all dark): uniform pick keeps - // the estimator unbiased and avoids a divide-by-zero. + // Fallback when no candidate carries Pick Power (all dark): uniform pick + // keeps the estimator unbiased and avoids a divide-by-zero. if (!(totalPower > 0.0f)) { const size_t numStrata = world.numLightInstances + (hasAmbient ? 1 : 0); const size_t selected = @@ -563,8 +576,8 @@ VISRTX_GLOBAL void __raygen__() // Gate on a positive pdf, NOT a fixed epsilon: a dim light's pick // probability can make the joint pdf legitimately tiny, and dividing // by it stays unbiased. An epsilon floor would drop those samples and - // render the dim light black — the very bright+dim case the power pick - // targets. + // render the dim light black — the very bright+dim case the power + // pick targets. if (lightSample.pdf > 0.0f && lightSample.dist > 0.0f) { const vec3 directLight = volumeSample.albedo * lightSample.radiance * INV_4PI / lightSample.pdf; @@ -657,18 +670,20 @@ VISRTX_GLOBAL void __raygen__() } // Emission, direct lighting are scaled by opacity analytically rather - // than gated stochastically below. A Geometry Light reached by a finite- - // pdf bounce is also sampled by NEE, so MIS-weight the deposit against - // that; a delta/primary bounce (bsdfPdf == +inf) keeps weight 1 since - // NEE cannot reach it, as do non-sampled emissive surfaces (pNee == 0). + // than gated stochastically below. A Geometry Light reached by a + // finite- pdf bounce is also sampled by NEE, so MIS-weight the deposit + // against that; a delta/primary bounce (bsdfPdf == +inf) keeps weight 1 + // since NEE cannot reach it, as do non-sampled emissive surfaces (pNee + // == 0). float wEmission = 1.0f; if (!isinf(bsdfPdf)) { - const float pNee = geometryLightHitPdf( - frameData, surfaceHit, ray.dir, materialEmission); + const float pNee = + geometryLightHitPdf(frameData, surfaceHit, ray.dir); if (pNee > 0.0f) wEmission = bsdfPdf / (bsdfPdf + pNee); } - sample.color += wEmission * sampleContribution * opacity * materialEmission; + sample.color += + wEmission * sampleContribution * opacity * materialEmission; // Sample around the shading normal so the cosine-weighted hemisphere's // pdf matches the BRDF's NdotL (which uses Ns). Sampling around Ng // would bias the Lambertian estimator by cos_Ns/cos_Ng on smooth or @@ -692,10 +707,10 @@ VISRTX_GLOBAL void __raygen__() // Env MIS: only the HDRI environment can also be reached by the // BSDF escape, so only it gets a balance-heuristic weight. The // light density uses envPdf on BOTH sides (here and at the miss), - // not lightSample.pdf, so wNee and wBsdf use identical pdf functions - // and partition to 1 exactly — unbiased regardless of how closely - // envPdf tracks the NEE importance pdf (the NEE estimator still - // divides by its true lightSample.pdf, which carries the same + // not lightSample.pdf, so wNee and wBsdf use identical pdf + // functions and partition to 1 exactly — unbiased regardless of how + // closely envPdf tracks the NEE importance pdf (the NEE estimator + // still divides by its true lightSample.pdf, which carries the same // envPickProb, inside materialShadeSurface). // Other light types: p_bsdf = 0 => w_nee = 1 (behaviour unchanged). float wNee = 1.0f; @@ -709,7 +724,8 @@ VISRTX_GLOBAL void __raygen__() // lightSample.pdf is the exact NEE density (solid-angle × pick // probability); the BSDF continuation can also hit this Geometry // Light, - // so weight against it. Mirrors geometryLightHitPdf on the deposit. + // so weight against it. Mirrors geometryLightHitPdf on the + // deposit. const float pBsdf = materialEvalPdf(shadingState, -ray.dir, lightSample.dir); wNee = lightSample.pdf / (lightSample.pdf + pBsdf); @@ -772,8 +788,9 @@ VISRTX_GLOBAL void __raygen__() if (!surfaceHit.foundHit && !volumeSample.didScatter) { // Deposit the environment, MIS-weighted against NEE. pLight mirrors the - // NEE env density: the HDRI importance pdf (envPdf) folded with the same - // power-proportional env pick probability sampleLights applied. bsdfPdf + // NEE env density: the HDRI importance pdf (envPdf) folded with the + // same power-proportional env pick probability sampleLights applied. + // bsdfPdf // == +inf (delta / transmission / primary ray) => w_bsdf = 1. if (vec3 hdri; getBackgroundLight(frameData, ray.dir, hdri)) { const float pLight = envPdf(frameData, ray.dir) * envPickProb; diff --git a/devices/rtx/device/sampler/Image2D.cpp b/devices/rtx/device/sampler/Image2D.cpp index fdc14b5b1..0bc6e4cf8 100644 --- a/devices/rtx/device/sampler/Image2D.cpp +++ b/devices/rtx/device/sampler/Image2D.cpp @@ -31,13 +31,20 @@ #include "Image2D.h" +#include "TextureStats.h" #include "utility/AnariTypeHelpers.h" namespace visrtx { -static float srgbToLinear(float v) +static TexelFormat texelFormat(ANARIDataType t) { - return v <= 0.04045f ? v / 12.92f : powf((v + 0.055f) / 1.055f, 2.4f); + if (isFloat32(t)) + return TexelFormat::Float32; + if (isSrgb8(t)) + return TexelFormat::Srgb8; + if (isFixed8(t)) + return TexelFormat::Fixed8; + return TexelFormat::Unsupported; } Image2D::Image2D(DeviceGlobalState *d) : Sampler(d), m_image(this) {} @@ -90,9 +97,9 @@ void Image2D::finalize() m_texels = makeCudaTexelObject2D( cuArray, !isFp, "nearest", m_wrap1, m_wrap2, m_borderColor); - // The mean texel is NOT computed here: it is only needed by the emissive - // Pick-Power path, so averageValue() scans lazily on first query and memoizes - // against the image stamp (see m_averageValue). + // The reduction is NOT computed here: it is only needed by the emissive + // Pick-Power / classifier path, so it scans lazily on first query and + // memoizes against the image stamp (see textureReduction() / m_reduction). upload(); } @@ -104,81 +111,97 @@ bool Image2D::isValid() const vec4 Image2D::averageValue() const { - // Lazy + guarded: recompute only when the bound image's data actually changed. - // A fresh (0) stamp forces the first compute; a filter/wrap recommit leaves the - // image stamp untouched and returns the cache. Non-emissive samplers never - // reach here at all. + // The non-negative magnitude (meanPositive) — the same proxy the MDL + // classifier uses, so a signed texel never inflates or cancels an emitter's + // picked power. Native emission is radiance >= 0, so this equals the plain + // mean for every real emitter. Alpha is unused by emission. + const auto &m = textureReduction().meanPositive; + return vec4(m[0], m[1], m[2], 1.f); +} + +#if defined(USE_MDL) +libmdl::ResourceStats Image2D::emissionStats() const +{ + const TextureReduction &r = textureReduction(); + libmdl::ResourceStats s; + s.valid = r.valid; + if (!r.valid) + return s; // Unknown: not a real reduction + s.maxAbs = r.maxAbs; + s.meanPositive = r.meanPositive; + s.minValue = r.minValue; + s.transferPreservesZero = r.transferPreservesZero; + s.finite = r.finite; + return s; +} +#endif + +// Lazy + guarded: recompute only when the bound image's data actually changed. +// A fresh (0) stamp forces the first compute; a filter/wrap recommit leaves the +// image stamp untouched and returns the cache. Non-emissive samplers never +// query it at all. +const Image2D::TextureReduction &Image2D::textureReduction() const +{ const helium::TimeStamp stamp = m_image ? m_image->lastDataModified() : helium::TimeStamp{0}; - if (stamp != m_averageValueStamp) { - m_averageValue = computeAverageValueGPU(); - m_averageValueStamp = stamp; + if (stamp != m_reductionStamp) { + m_reduction = computeTextureReduction(); + m_reductionStamp = stamp; } - return m_averageValue; + return m_reduction; } -// Mean linear texel, used only to size a textured emitter's Pick Power -// (variance, never bias). Reads the retained host pixels; sRGB byte data is -// linearized to match the hardware sampler. Unsupported element types fall back -// to the fully-lit default so the emitter is still picked. Computed lazily and -// memoized; see averageValue() / m_averageValue. -vec4 Image2D::computeAverageValue() const +// One thrust pass over the resident device texels (Array::data(GPU) is a real +// H2D upload) yielding, per channel, the max absolute value (exact zero proof), +// the mean of the positive part (the non-negative magnitude that sizes a +// textured emitter's Pick Power — variance, never bias), and the min value +// (non-negative sign proof). The device functor linearizes sRGB byte data to +// match the hardware sampler. Unsupported element types or a missing device +// residency leave the reduction Unknown (magnitude stays unit so the emitter is +// still picked). +Image2D::TextureReduction Image2D::computeTextureReduction() const { + TextureReduction r; if (!m_image) - return Sampler::averageValue(); + return r; // valid=false, magnitude stays unit const ANARIDataType t = m_image->elementType(); const int nc = numANARIChannels(t); const size_t count = size_t(m_image->size().x) * m_image->size().y; - const void *host = m_image->data(AddressSpace::HOST); - if (nc == 0 || count == 0 || !host) - return Sampler::averageValue(); - - // sRGB 8-bit formats carry a linear alpha in the LAST channel (present for the - // RGBA/RA variants, i.e. even channel counts); only the color channels are - // gamma-encoded. - const bool srgb = isSrgb8(t); - const int colorChannels = (srgb && (nc == 2 || nc == 4)) ? nc - 1 : nc; - - glm::dvec4 sum(0.0); - if (isFloat32(t)) { - const auto *p = static_cast(host); - for (size_t i = 0; i < count; ++i) - for (int c = 0; c < nc; ++c) - sum[c] += double(p[i * nc + c]); - } else if (isFixed8(t) || srgb) { - const auto *p = static_cast(host); - for (size_t i = 0; i < count; ++i) - for (int c = 0; c < nc; ++c) { - const float v = p[i * nc + c] / 255.0f; - sum[c] += double((srgb && c < colorChannels) ? srgbToLinear(v) : v); - } - } else { - return Sampler::averageValue(); // uncommon type for emission; coarse fallback + const TexelFormat fmt = texelFormat(t); + if (nc == 0 || count == 0 || fmt == TexelFormat::Unsupported) + return r; + + // sRGB 8-bit formats carry a linear alpha in the LAST channel (present for + // the RGBA/RA variants, i.e. even channel counts); only the color channels + // are gamma-encoded. + const int colorChannels = + (fmt == TexelFormat::Srgb8 && (nc == 2 || nc == 4)) ? nc - 1 : nc; + + const void *dev = m_image->data(AddressSpace::GPU); + if (!dev) + return r; // no device residency ⇒ Unknown + + const TexelAccum a = reduceTexelsDevice(dev, fmt, nc, colorChannels, count); + + // Broadcast source channels to rgb: a 1/2-channel texture drives all three + // color channels from channel 0, so a grayscale emissive texture still yields + // a sensible magnitude color. + auto channelForRGB = [&](int rgb) { return nc >= 3 ? rgb : 0; }; + for (int rgb = 0; rgb < 3; ++rgb) { + const int c = channelForRGB(rgb); + r.meanPositive[rgb] = float(a.posSum[c] / double(count)); + r.maxAbs[rgb] = a.maxAbs[c]; + r.minValue[rgb] = a.minValue[c]; } - glm::dvec4 avg = sum / double(count); - // Broadcast 1/2-channel textures to RGB so a grayscale emissive texture still - // yields a sensible average color. - if (nc == 1) - return vec4(float(avg.x), float(avg.x), float(avg.x), 1.f); - if (nc == 2) - return vec4(float(avg.x), float(avg.x), float(avg.x), float(avg.y)); - if (nc == 3) - return vec4(float(avg.x), float(avg.y), float(avg.z), 1.f); - return vec4(avg); -} - -// TODO(perf): reduce over the resident texels on the device instead of the host -// scan above — the image is already uploaded as a cudaArray for sampling -// (m_texels), so a device reduction (cf. the thrust::reduce-over-image in -// light/sampling/CDF.cu) avoids the host readback entirely for large emissive -// textures. The kernel must reproduce computeAverageValue()'s per-channel -// sRGB->linear (color channels only) and the 1/2-channel broadcast, then read -// back a single vec4. Stubbed: delegates to the host scan for now. -vec4 Image2D::computeAverageValueGPU() const -{ - return computeAverageValue(); + r.valid = true; + r.finite = a.finite; + // The sRGB and linear transfers satisfy T(0)=0; the only way a stored-zero + // texel samples nonzero is a nonzero border color under a border wrap mode. + r.transferPreservesZero = m_borderColor.x == 0.0f && m_borderColor.y == 0.0f + && m_borderColor.z == 0.0f; + return r; } int Image2D::numChannels() const diff --git a/devices/rtx/device/sampler/Image2D.h b/devices/rtx/device/sampler/Image2D.h index 78c60a86c..7037aa3f1 100644 --- a/devices/rtx/device/sampler/Image2D.h +++ b/devices/rtx/device/sampler/Image2D.h @@ -35,6 +35,8 @@ #include "array/Array2D.h" #include "utility/CudaImageTexture.h" +#include + namespace visrtx { struct Image2D : public Sampler @@ -48,17 +50,36 @@ struct Image2D : public Sampler int numChannels() const override; vec4 averageValue() const override; +#if defined(USE_MDL) + libmdl::ResourceStats emissionStats() const override; +#endif cudaTextureObject_t textureObject() const; private: SamplerGPUData gpuData() const override; - // Mean linear texel for emissive Pick Power. computeAverageValue() is the host - // scan; computeAverageValueGPU() is the intended device-side reduction over the - // already-resident texels (stubbed — currently delegates to the host scan). - vec4 computeAverageValue() const; - vec4 computeAverageValueGPU() const; + // Single-pass texel reduction feeding both the emissive Pick Power and the + // MDL emission classifier from one scan (replacing the former separate + // averageValue and emissionStats scans). Pick Power reads meanPositive — the + // same non-negative magnitude proxy the classifier uses — so a signed texel + // never inflates or cancels an emitter's picked power. + struct TextureReduction + { + // Per channel; unit default so an un-reduced emitter is still picked. + std::array maxAbs{{1.f, 1.f, 1.f}}; // maxAbs==0 ⇒ exact zero + std::array meanPositive{{1.f, 1.f, 1.f}}; // magnitude / Pick Power + std::array minValue{{-1.f, -1.f, -1.f}}; // minValue>=0 ⇒ nonneg + bool transferPreservesZero{false}; // T(0)==0 unless a nonzero border color + bool finite{true}; + bool valid{false}; // false ⇒ unbound/unsupported ⇒ classifier Unknown + }; + + // Lazy, memoized against the image data stamp: computed on the first query + // and reused until the bound image's texels actually change. Non-emissive + // samplers never query it and so never scan. + const TextureReduction &textureReduction() const; + TextureReduction computeTextureReduction() const; void cleanupImageCudaArray(); void cleanupImageTextureObjects(); @@ -71,13 +92,11 @@ struct Image2D : public Sampler cudaTextureObject_t m_texture{}; cudaTextureObject_t m_texels{}; - // Mean linear texel, consumed only by the emissive Pick-Power path. Computed - // lazily on the first averageValue() query and memoized against the image's - // lastDataModified stamp: non-emissive samplers (base color, normal, roughness, - // ...) never query it and so never scan, and a no-op recommit (filter/wrap - // change, scene churn) does not rescan. mutable: filled from the const query. - mutable vec4 m_averageValue{1.f}; - mutable helium::TimeStamp m_averageValueStamp{0}; + // Memoized reduction, filled lazily from the const query and guarded on the + // image's lastDataModified stamp: a no-op recommit (filter/wrap change, scene + // churn) does not rescan. mutable: filled from the const query. + mutable TextureReduction m_reduction; + mutable helium::TimeStamp m_reductionStamp{0}; }; } // namespace visrtx diff --git a/devices/rtx/device/sampler/Sampler.cpp b/devices/rtx/device/sampler/Sampler.cpp index e6af6c719..1e7affe20 100644 --- a/devices/rtx/device/sampler/Sampler.cpp +++ b/devices/rtx/device/sampler/Sampler.cpp @@ -70,6 +70,24 @@ vec4 Sampler::averageValue() const return vec4(1.f); } +#if defined(USE_MDL) +libmdl::ResourceStats Sampler::emissionStats() const +{ + // A sampler that cannot reduce its texels proves nothing: not zero (maxAbs + // nonzero), not non-negative (minValue negative), unit magnitude proxy. Its + // emission stays register-ineligible via the policy's sign gate — + // forward-only, unbiased — until a real reduction (Image2D) is available. + libmdl::ResourceStats s; + s.valid = true; + s.maxAbs = {1.f, 1.f, 1.f}; + s.meanPositive = {1.f, 1.f, 1.f}; + s.minValue = {-1.f, -1.f, -1.f}; + s.transferPreservesZero = false; + s.finite = true; + return s; +} +#endif + void Sampler::commitParameters() { m_inAttribute = getParamString("inAttribute", "attribute0"); diff --git a/devices/rtx/device/sampler/Sampler.h b/devices/rtx/device/sampler/Sampler.h index 087439f1d..4cb0c95ee 100644 --- a/devices/rtx/device/sampler/Sampler.h +++ b/devices/rtx/device/sampler/Sampler.h @@ -32,6 +32,9 @@ #pragma once #include "RegisteredObject.h" +#if defined(USE_MDL) +#include "libmdl/ResourceStats.h" +#endif namespace visrtx { @@ -48,6 +51,17 @@ struct Sampler : public RegisteredObject // un-averaged sampler is still picked; Image2D overrides with the mean texel. virtual vec4 averageValue() const; +#if defined(USE_MDL) + // Per-channel texel reduction consumed by the MDL emission classifier's value + // source (maxAbs for the zero proof, meanPositive for the magnitude proxy, + // minValue for the sign proof). The default is Unknown (valid but unproven): + // a sampler that cannot reduce its texels neither proves zero nor proves a + // non-negative sign, so its emission stays register-eligible only under the + // policy's Unknown path. Image2D overrides with a real scan. MDL-only: the + // classifier is the sole consumer and lives behind USE_MDL. + virtual libmdl::ResourceStats emissionStats() const; +#endif + static Sampler *createInstance( std::string_view subtype, DeviceGlobalState *d); diff --git a/devices/rtx/device/sampler/TextureStats.cu b/devices/rtx/device/sampler/TextureStats.cu new file mode 100644 index 000000000..ebef16c0c --- /dev/null +++ b/devices/rtx/device/sampler/TextureStats.cu @@ -0,0 +1,118 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#include "TextureStats.h" + +#include +#include +#include + +#include +#include +#include + +namespace visrtx { + +namespace { + +__device__ inline float srgbToLinear(float v) +{ + return v <= 0.04045f ? v / 12.92f : powf((v + 0.055f) / 1.055f, 2.4f); +} + +// Device-friendly POD mirror of TexelAccum (plain arrays so thrust can hold it +// in registers during the fold). Identity element: zero sums, zero maxAbs, +// +inf minValue, finite. +struct Accum4 +{ + double posSum[4]; + float maxAbs[4]; + float minValue[4]; + int finite; +}; + +__host__ __device__ inline Accum4 identityAccum() +{ + Accum4 a; + for (int c = 0; c < 4; ++c) { + a.posSum[c] = 0.0; + a.maxAbs[c] = 0.0f; + a.minValue[c] = FLT_MAX; + } + a.finite = 1; + return a; +} + +// Decode one texel's nc channels to linear and fold them into an Accum4. Only +// the touched channels [0,nc) are set; the rest keep the identity so unused +// channels never pollute the reduction. +struct DecodeTexel +{ + const void *data; + TexelFormat fmt; + int nc; + int colorChannels; + + __device__ Accum4 operator()(std::size_t i) const + { + Accum4 a = identityAccum(); + for (int c = 0; c < nc; ++c) { + float v; + if (fmt == TexelFormat::Float32) { + v = static_cast(data)[i * nc + c]; + } else { + const float u = static_cast(data)[i * nc + c] / 255.0f; + v = (fmt == TexelFormat::Srgb8 && c < colorChannels) ? srgbToLinear(u) + : u; + } + a.posSum[c] = fmaxf(v, 0.0f); + a.maxAbs[c] = fabsf(v); + a.minValue[c] = v; + if (!isfinite(v)) + a.finite = 0; + } + return a; + } +}; + +struct MergeAccum +{ + __host__ __device__ Accum4 operator()(const Accum4 &x, const Accum4 &y) const + { + Accum4 r; + for (int c = 0; c < 4; ++c) { + r.posSum[c] = x.posSum[c] + y.posSum[c]; + r.maxAbs[c] = fmaxf(x.maxAbs[c], y.maxAbs[c]); + r.minValue[c] = fminf(x.minValue[c], y.minValue[c]); + } + r.finite = x.finite & y.finite; + return r; + } +}; + +} // namespace + +TexelAccum reduceTexelsDevice(const void *deviceData, + TexelFormat fmt, + int nc, + int colorChannels, + std::size_t count) +{ + const Accum4 folded = thrust::transform_reduce(thrust::device, + thrust::counting_iterator(0), + thrust::counting_iterator(count), + DecodeTexel{deviceData, fmt, nc, colorChannels}, + identityAccum(), + MergeAccum{}); + + TexelAccum out; + for (int c = 0; c < 4; ++c) { + out.posSum[c] = folded.posSum[c]; + out.maxAbs[c] = folded.maxAbs[c]; + out.minValue[c] = folded.minValue[c]; + } + out.finite = folded.finite != 0; + return out; +} + +} // namespace visrtx diff --git a/devices/rtx/device/sampler/TextureStats.h b/devices/rtx/device/sampler/TextureStats.h new file mode 100644 index 000000000..d3efd5632 --- /dev/null +++ b/devices/rtx/device/sampler/TextureStats.h @@ -0,0 +1,45 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// Device-side texel reduction for Image2D: one thrust pass over the resident +// texels yields the raw per-source-channel accumulators the emissive Pick Power +// and the emission classifier (maxAbs / meanPositive / minValue) are built +// from. SDK-free and MDL-agnostic — averageValue() feeds the non-MDL PBR path. + +#include +#include +#include + +namespace visrtx { + +enum class TexelFormat +{ + Unsupported, + Float32, // raw float channels + Fixed8, // uint8/255, linear + Srgb8 // uint8/255, sRGB->linear on color channels +}; + +// Raw per-source-channel reduction (up to 4 channels), decoded to linear. The +// caller applies the magnitude/broadcast and sign policy. +struct TexelAccum +{ + std::array posSum{}; // Σ max(v,0) → meanPositive + std::array maxAbs{}; // max |v| → exact zero proof + std::array minValue{ + {FLT_MAX, FLT_MAX, FLT_MAX, FLT_MAX}}; // min v → sign proof + bool finite{true}; +}; + +// Reduce `count` texels of `nc` channels resident at `deviceData` (a device +// pointer from Array::data(AddressSpace::GPU)). sRGB color channels (source +// index < colorChannels) are linearized to match the hardware sampler. +TexelAccum reduceTexelsDevice(const void *deviceData, + TexelFormat fmt, + int nc, + int colorChannels, + std::size_t count); + +} // namespace visrtx diff --git a/devices/rtx/device/world/World.cpp b/devices/rtx/device/world/World.cpp index 87ff8449a..47a49817c 100644 --- a/devices/rtx/device/world/World.cpp +++ b/devices/rtx/device/world/World.cpp @@ -681,14 +681,23 @@ void World::buildInstanceLightGPUData() // surface instance and emit silently wrong radiance — catch it here. assert(surfaceInstanceCursor == m_instanceSurfaceGPUData.size()); - // Turn the per-instance Pick Powers into a normalized cumulative CDF in place. - // A zero total (every light dark) leaves the CDF unused: the renderer falls - // back to a uniform pick. + // Turn the per-instance Pick Powers into a normalized cumulative CDF in + // place. Accumulate and normalize in double, dividing by the DOUBLE + // cumulative total (not the float m_totalLightPower): normalizing by the + // float total can push the last entry above 1.0 when float lost a dim light's + // mass, leaving trailing dim lights unselectable while their hit-side pNee + // stays positive — bias. The double total makes the last entry exactly 1.0 + // and preserves those masses. A zero total (every light dark) leaves the CDF + // unused: uniform pick. if (m_totalLightPower > 0.0f) { - float cumulative = 0.0f; + double total = 0.0; + for (size_t i = 0; i < totalLights; ++i) + total += pickCdf[i]; + const double invTotal = total > 0.0 ? 1.0 / total : 0.0; + double cumulative = 0.0; for (size_t i = 0; i < totalLights; ++i) { cumulative += pickCdf[i]; - pickCdf[i] = cumulative / m_totalLightPower; + pickCdf[i] = cumulative * invTotal; } } diff --git a/devices/rtx/device/world/World.h b/devices/rtx/device/world/World.h index 8b5d19fe5..8a9194870 100644 --- a/devices/rtx/device/world/World.h +++ b/devices/rtx/device/world/World.h @@ -115,7 +115,7 @@ struct World : public Object HostDeviceArray m_instanceHdriLightGPUData; // Power-proportional Light Pick, rebuilt with the light instances. - HostDeviceArray m_lightPickCdf; + HostDeviceArray m_lightPickCdf; float m_totalLightPower{0.f}; float m_hdriPower{0.f}; float m_sceneRadius{0.f}; diff --git a/devices/rtx/docs/adr/0005-geometry-lights-share-the-light-sampling-path.md b/devices/rtx/docs/adr/0005-geometry-lights-share-the-light-sampling-path.md index 6d6e15020..056b3a392 100644 --- a/devices/rtx/docs/adr/0005-geometry-lights-share-the-light-sampling-path.md +++ b/devices/rtx/docs/adr/0005-geometry-lights-share-the-light-sampling-path.md @@ -28,10 +28,11 @@ equivalent. The pick probability is `pickPower / totalPower`, the same formula the pick CDF uses, so the two sides agree (and MIS stays unbiased regardless of float drift, since the balance-heuristic weights sum to 1). This keeps the - gate to one per-material flag (`emissionIsConstant`) plus the geometry's - area, both of which each object keeps current in its own GPU slot, so - neither can go stale — no reverse-lookup table, no extra `SurfaceHit` - field. + gate to one per-material flag (`emissionIsSampleable`; the constant-radiance + deposit fast path this ADR's Stage-1 emitters take reads a second flag, + `emissionIsConstant`) plus the geometry's area, both of which each object + keeps current in its own GPU slot, so neither can go stale — no + reverse-lookup table, no extra `SurfaceHit` field. - `{lightIndex, xfm}` cannot represent emission driven by instance-uniform attributes. Once non-constant emission becomes sampleable, the light instance must also reference its surface instance so next-event radiance diff --git a/devices/rtx/docs/adr/0007-emission-descriptor-registration-policy.md b/devices/rtx/docs/adr/0007-emission-descriptor-registration-policy.md new file mode 100644 index 000000000..342f95dd0 --- /dev/null +++ b/devices/rtx/docs/adr/0007-emission-descriptor-registration-policy.md @@ -0,0 +1,310 @@ +# Emission descriptor and faithfulSet registration policy + +An MDL material's emission is analyzed into an **immutable descriptor** that +*describes* — it never decides registration. A thin, renderer-side **policy** +decides which described slots become next-event-sampled Geometry Lights. This +splits a previously entangled classifier into three concerns: + +1. **Classifier** (MDL-pure, in `libmdl`): walks the compiled material's emission + DAG, emits an owned IR, and folds a descriptor `{verdict, edfKinds, magnitude, + mode}` per slot plus the argument/resource dependencies the emission reads. It + carries no renderer knowledge and retains no MDL-SDK expression pointers. +2. **Contract** (the seam): the descriptor shape plus a consumer-exported + `faithfulSet` — the EDF kinds the renderer can evaluate faithfully on its + synthetic next-event hit. +3. **Policy** (renderer-side): `register(slot)` iff the slot is consumed **and** + `verdict ≠ ProvablyNull` **and** `edfKinds ⊆ faithfulSet` **and** the intensity + mode is faithfully handled. + +This supersedes the classification rule of +[ADR-0006](0006-sampleable-mdl-emission.md): the "emission not provably zero → +Emissive Surface" premise and the diffuse-EDF fidelity scope are preserved, but +the *decision* of whether to register moves out of the classifier into the +policy, and the classifier's output becomes a complete, honest descriptor rather +than a sampleability verdict. ADR-0006's synthetic-hit fidelity limits +(§EDF fidelity scope) and Pick-Power semantics remain in force. + +## Emission IR (lowering the MDL expression DAG) + +The classifier does **not** fold the MDL SDK's `IExpression` DAG directly. It +first lowers the compiled material's emission expressions (surface/backface × +`{EDF, intensity, mode}` plus `thin_walled`) into an owned **emission IR** — a +one-time projection walked by `IExpression::get_kind()`: + +| MDL SDK expression | IR node | +|---|---| +| `EK_TEMPORARY` (`IExpression_temporary`) | dereferenced and memoized, so a CSE-shared temporary collapses to a single node | +| `EK_CONSTANT` (`IExpression_constant` → `IValue`) | `Constant` — the value is copied out | +| `EK_PARAMETER` (`IExpression_parameter`) | `Parameter` | +| `EK_DIRECT_CALL` (`IExpression_direct_call`) | `Call` / `Texture`, tagged with `IFunction_definition::get_semantic()` | + +Each node stores a `Semantic` enum, operand indices, copied constant values, and +resource names — and **no `IExpression` handle**. Three properties motivate the +lowering rather than folding `IExpression` in place: + +- **Lifetime**: the device does not retain the compiled material, so any held + `IExpression*` would dangle. The IR owns everything the fold reads and outlives + the SDK objects. +- **Stable keying**: nodes key on `get_semantic()`, never DB names, so no user + module can masquerade as a `df::`/`tex::` intrinsic. +- **SDK-runtime-free fold**: `foldEmissionDescriptor` interprets the IR with no + MDL-SDK *runtime* calls or objects — link-level SDK-free. It still includes SDK + headers for the `Semantic` type (`IFunction_definition::Semantics`) the IR keys + on. What crosses into the renderer is the descriptor, which is fully SDK-free + (no `mi::` types), letting the device consume it without the SDK. + +Temporary memoization also makes the `−` exact-identity test decidable: shared +subexpressions become the same node, so the lattice's subtraction rule +(below) can prove `ProvablyZero` by IR-node ref-compare. + +## Error model + +The default renderer is a Quality path tracer whose forward estimator deposits +emission on any BSDF closest-hit at MIS weight 1 when the surface is **not** +registered as a light. This asymmetry drives the whole design: + +- **Miss (false negative) = variance, not bias**, in unbounded Quality: an + unregistered emitter still deposits via the forward path, unbiased, just noisier + (no next-event contribution). Miss is *bias* (dark) only where no forward + estimator survives — the Interactive/Matte/finite-depth modes below. (Even the + default Quality path is finite — `maxRayDepth=5` — so a miss leaves a small + final-vertex darkening; strict "variance, converges" holds only at unbounded + depth. The renderer work item that closes the finite-depth gap removes this.) +- **Over-register (false positive) = BIAS**, not free perf: the next-event + synthetic hit is fidelity-limited (geometric normal, synthesized tangent, + object id 0, forced front — ADR-0006:102). Registering an EDF the renderer + cannot evaluate faithfully makes next-event repay a *wrong* integrand while the + correct forward deposit is MIS-downweighted — a systematic error the sample + count never removes. + +Therefore the policy registers **only** what is both non-null and faithfully +evaluable. An unfaithful or unknown EDF kind is *described, never registered* — +its light still arrives via the forward path, unbiased. + +### Render-mode matrix (why a miss matters where) + +| Mode | A missed (unregistered) emitter | +|---|---| +| Quality, unbounded depth | variance (converges) | +| Quality, `maxRayDepth=1` / final path vertex | dark | +| Matte receivers | dark | +| Interactive | dark receivers (no hit-side geometry-light MIS) | +| Fast | emission invisible regardless | + +The dark cases are pre-existing forward-estimator gaps (today's classifier +already rejects every non-diffuse EDF, so those emitters are already +forward-only). Collapsing this matrix to "variance everywhere" is tracked as an +independent renderer work item, not a prerequisite of this policy. + +## Lattice and op table + +Analysis is a three-valued abstract interpretation over the **current immutable +snapshot** (argument block + resource table); the reactive re-fold owns future +writes. `ProvablyZero` means identically zero over all uv/state/time at the +current snapshot; any doubt joins to `Unknown`. + +Scalar/color lattice `{ProvablyZero, ProvablyNonZero, Unknown}`; EDF lattice +`{ProvablyNull, ProvablyEmissive, Unknown}`. + +- `·` is zero-absorbing; `+` is zero iff both operands are zero. +- `−`: `ProvablyZero` **only** on exact IR-node identity (CSE temporaries make + ref-compare decidable; distinct-but-equal nodes fall to `Unknown`). Otherwise + `Unknown`, never `ProvablyNonZero` (`1 − w` with `w` unknown can be zero). +- `?:`: fold if the condition folds to a constant at the current snapshot, else + **join = least-upper-bound with `Unknown` as top**: `Zero ⊔ Zero = Zero`; + `Zero ⊔ NonZero = Unknown`; any `Unknown ⇒ Unknown`. EDF lattice likewise: + `Null ⊔ Null = Null`, `Null ⊔ Emissive = Unknown`, else `Unknown`. +- Mixes are analyzed conservatively: the fold unions the EDF kinds of every + reachable component regardless of weight (a zero-weight component still + contributes its kind) and never proves a mix `ProvablyNull`. Weight-based + pruning — dropping a `ProvablyZero`-weighted component (cross-channel for + color weights) — is a sound refinement, not yet implemented. +- `/`, `pow`, `exp`, and any **unmodeled node ⇒ `Unknown`**. +- Color zero test is **cross-channel**: `ProvablyZero` iff the max over *all* + channels is zero — never a luminance/scalar reduction (`(+1,−1,0)` is not + provably zero). +- Proofs assume finite values; a non-finite texel flags `Unknown`. + +### Texture reductions + +Per canvas, one pass at load / content-change yields +`{maxAbs, meanPositive, minValue}` per channel (memoized on the image's +content-version stamp): + +- `maxAbs == 0 ⇒ ProvablyZero` — exact, no epsilon, over all texels and canvases. +- The gate is in **sampler-output space**: a transfer `T` with `T(0) ≠ 0` + (LUT / ICC / nonzero border) breaks the stored-texel bound ⇒ `Unknown`. The + standard MDL transfers (`tex::gamma_*`) and `wrap_clip` satisfy `T(0) = 0`. +- The reduction also yields a per-channel **`minValue`**: `minValue ≥ 0` over all + texels proves the texture's `sign` contribution is `ProvablyNonnegative`; a + negative texel makes it `Unknown`. +- The magnitude proxy is `meanPositive` — the per-channel mean of `max(texel, 0)` + — never mean-absolute (see the negative-emission decision). It never gates + zero; it is non-negative by construction so a CDF can represent it. This one + magnitude sizes the emissive Pick Power for **every** emissive material type — + native PBR and MDL alike read it, not a separate signed mean. The two coincide + for any registerable emitter (whose texels are all ≥ 0), and `meanPositive` + stays CDF-valid even for the signed emitters that never register. + +## Descriptor and contract + +``` +SlotDesc = { verdict: Null|Emissive|Unknown, + edfKinds: set, + magnitude: meanPositiveRadianceProxy (per channel, >= 0), + mode: radiant_exitance|power, + dependsOnGeometricState: bool, // intensity/EDF reads normal/tangent/objectId/position + sign: ProvablyNonnegative|Unknown } + +EmissionDescriptor = { + surface: SlotDesc, backface: SlotDesc, // each folded from its own sub-expressions +} +``` + +The emission's argument/resource dependencies are **not** on the descriptor — +they live on the owned IR (`EmissionIR`), computed by `collectDeps`: + +``` +emissionDeps: set // structural — every branch of every ?: +resourceDeps: set +``` + +The descriptor is **complete and honest**: it describes the backface slot and +unfaithful EDF kinds even though today's consumer ignores them. The classifier +needs no change when the renderer's fidelity grows — only `faithfulSet` grows. + +Two per-slot flags make the faithfulness gate *sufficient*, not just necessary +(an EDF-kind check alone is not enough — see Considered options): + +- **`dependsOnGeometricState`** — set when the emission (EDF **or** intensity) + reads a geometric-state quantity the synthetic next-event hit fabricates: + shading normal (`state::normal`, bump), tangent (`state::texture_tangent_*`), + object/instance id, or position. Such emission evaluates a *different* + integrand at the synthetic hit than at the real forward hit, so registering it + biases even when the EDF kind is faithful. The fold detects it structurally + from the IR (any reachable `state::` intrinsic in the diffuse-fidelity set). +- **`sign`** — `ProvablyNonnegative` when no reachable emission value can be + negative (constants with all channels ≥ 0; textures whose reduction proves a + non-negative minimum; otherwise `Unknown`). Signed emission renders correctly + via the forward path (the device does not clamp), but it cannot be *registered* + faithfully: an all-negative next-event contribution is dropped by the shadow + ray's positive-contribution epsilon gate while the forward deposit is + MIS-downweighted, under-applying the light. So only `ProvablyNonnegative` + emission is registerable; the rest stays forward-only, unbiased. + +`faithfulSet` is a single consumer-exported constant. Today it is `{diffuse}` +(ADR-0006: emission must be invariant to normal/tangent/object-id). It lives in +one header (`devices/rtx/device/material/EmissionPolicy.h`) and is cross- +referenced from the two GPU-side sites that encode the same diffuse assumption +(`lightPickPower.h` double-sided Lambertian flux, `sampleLight.h` double-sided +normal orientation) so the assumption has a single source of truth. + +### Registration policy + +``` +register(slot) iff consumed(slot) // call-site: the caller passes only consumed slots (surface today) + and slot.verdict ≠ ProvablyNull // ┐ isRegisterable(slot) — the five + and slot.edfKinds ⊆ faithfulSet // │ faithfulness conjuncts, in + and slot.mode is faithfully handled // │ EmissionPolicy.h (radiant_exitance today) + and not slot.dependsOnGeometricState // │ synthetic hit fabricates it + and slot.sign == ProvablyNonnegative // ┘ else forward-only, unbiased +``` + +`consumed(slot)` is not a term inside `isRegisterable`; it is expressed by which +slots the caller folds and tests (surface only today — backface is described but +not yet consumed). + +`Unknown` *verdict* with faithful kinds (and the two faithfulness flags +satisfied) ⇒ register: the worst case is a spurious zero-radiance diffuse light, +genuinely perf-only and unbiased, because the magnitude proxy is non-negative and +the emission is state-invariant. Any unfaithful/unknown *kind*, a +geometric-state dependence, a possibly-negative value, or a power mode ⇒ do not +register — the forward path keeps that light unbiased; registering would bias. + +## Considered options + +- **Descriptor vs verdict.** The classifier emits a descriptor and a separate + policy decides (chosen). Rejected: the classifier deciding sampleability + directly (ADR-0006's model) — it entangles renderer fidelity into MDL-pure + analysis, so every fidelity gain (a new faithful EDF, backface consumption) + requires editing the classifier. A describing classifier plus a thin policy + restores renderer-agnosticism: `faithfulSet` grows, classifier unchanged. +- **`faithfulSet` export.** A static `constexpr` set in one renderer header, + cross-referenced from the GPU lockstep sites (chosen). Rejected: a queried + runtime capability — the renderer owner is the same team, and a compile-time + constant makes the three lockstep sites (classifier gate, Pick-Power flux, + synthetic-normal orientation) impossible to desynchronize silently. +- **Negative-emission policy / magnitude proxy (#6).** The device does **not** + clamp negative emission (`MDLShader_ptx.cu`; the host mirrors this deliberately, + as ADR-0006 established), so signed emission must keep rendering correctly. + Chosen: **signed emission renders via the forward path; only + `ProvablyNonnegative` emission is registered, with a non-negative + `meanPositive` magnitude proxy.** This is the only coherent choice: a + mean-**absolute** proxy (an earlier draft) gives a negative emitter a *positive* + Pick Power, so next-event selects it, but its all-negative shadow-ray + contribution is then dropped by the positive-contribution epsilon gate while + the forward deposit is MIS-downweighted — the light is under-applied, a bias + the sample count never removes. Making the magnitude non-negative *and* gating + registration on `sign == ProvablyNonnegative` keeps negative emission on the + unbiased forward path (a purely-negative emitter folds to magnitude 0 → zero + Pick Power → never NEE-selected → forward deposit at weight 1). Rejected: + clamping negative emission to zero **on the device** — it would let any signed + emitter register but silently changes rendered output; a device clamp is a + separate change (renderer ticket) that, if made, would let `sign` treat clamped + emission as non-negative. The host magnitude and the device hit-side pNee must + read the **same** non-negative proxy (consumer obligation below), or a CDF + built from `meanPositive` disagrees with a hit-side pNee derived from signed + radiance — MIS bias. Today's constant path feeds signed radiance to the hit-side + pdf, which this ADR flags as a renderer work item. +- **Emission `.mode`.** The descriptor records `radiant_exitance` vs `power`; + only `radiant_exitance` is faithfully handled today, so `power` is + described-but-not-registered (forward-path only), identical to current behavior + but now principled rather than a hard classifier rejection. Power-mode + area-normalization is a follow-up. +- **Reactive re-fold and atomic publish.** The descriptor is re-folded whenever + the emission's argument/resource dependencies change, which happens on the + device's existing commit path (all argument writes flow through + `syncParameters` inside `finalize`); the existing light-set refresh + (`refreshEmissionLightSet` → `lastLightSetChange`) then republishes light data, + CDF, and sampleability flags together in one host-serialized pre-launch rebuild. + No new double-buffer machinery: "atomic publish" reduces to the invariant *every + descriptor change affecting verdict or magnitude bumps `lastLightSetChange` + before launch*, enforced by test. + +## Consequences + +- The classifier becomes MDL-pure and independently unit-testable without a GPU: + `libmdl` is a standalone static library, so the static pass, IR, and fold are + exercised device-free against inline `.mdl` snippets. +- ADR-0006's `EmissionClassification` (constant radiance + textured-diffuse flag + + single-factor dynamic recipe) and its string-prefix DAG walk are replaced by + the owned IR + descriptor fold, keyed on inlined `df::` **semantics** + (`IFunction_definition::get_semantic`) rather than DB-name prefixes. +- `PhysicallyBasedMDL` continues to read its post-translate keys (it does not + introspect the DAG) but now publishes the same descriptor type, so both raw and + wrapper materials feed one policy. Its live-sampler-mean Pick Power is preserved. +- `Unknown`-verdict diffuse emission now **registers** (worst case a spurious, + unbiased zero-radiance light) where the old classifier's non-fold fell back to a + unit proxy; spot/measured/power emission is **described but not registered**, + identical rendered result to today (forward-only) but now via an explicit + fidelity gate rather than an EDF-kind rejection. +- Volume emission (`SLOT_VOLUME_EMISSION_INTENSITY`) is out of scope; the surface + and backface slots are the analysis subject. + +## Follow-ups + +1. Synthetic-hit fidelity (real normal/tangent/object-id, angular EDF eval) → + grows `faithfulSet` to spot/directional/measured. +2. Consume the descriptor's `backface` slot (side-aware hit/PDF/Pick-Power). +3. Collapse the render-mode matrix: forward-estimator deposits in + Interactive/Matte/finite-depth so a miss is variance in every mode. +4. Float-CDF dim-light bias: preserve every positive mass or derive hit-side + `pNee` from the quantized CDF delta. +5. `intensity_power` mode area-normalization; register power-mode emitters. +6. Hit-side pNee must read the **same** non-negative `meanPositive` magnitude the + Pick-Power CDF is built from; today the constant path feeds signed radiance to + the hit-side geometry-light pdf, which would disagree with a `meanPositive` CDF. + (A device negative-emission clamp would resolve this and would let `sign` treat + clamped emission as non-negative.) + +These are filed as independent renderer tickets; none blocks the classifier. diff --git a/devices/rtx/libmdl/CMakeLists.txt b/devices/rtx/libmdl/CMakeLists.txt index dcdb33831..539d450fb 100644 --- a/devices/rtx/libmdl/CMakeLists.txt +++ b/devices/rtx/libmdl/CMakeLists.txt @@ -10,6 +10,8 @@ project_add_library(STATIC ArgumentBlockDescriptor.cpp ArgumentBlockInstance.cpp Core.cpp + EmissionFold.cpp + EmissionIR.cpp helpers.cpp ptx.cpp TimeStamp.cpp diff --git a/devices/rtx/libmdl/Core.cpp b/devices/rtx/libmdl/Core.cpp index e58c1fb6d..1eed179f0 100644 --- a/devices/rtx/libmdl/Core.cpp +++ b/devices/rtx/libmdl/Core.cpp @@ -384,11 +384,10 @@ const mi::neuraylib::IModule *Core::loadModuleByCanonicalName( // resolver pick a complete copy on the search path. auto pathsCount = mdlConfiguration->get_mdl_paths_length(); for (auto i = decltype(pathsCount)(0); i < pathsCount; ++i) { - auto rootName = - std::filesystem::path( - make_handle(mdlConfiguration->get_mdl_path(i))->get_c_str()) - .filename() - .string(); + auto rootName = std::filesystem::path( + make_handle(mdlConfiguration->get_mdl_path(i))->get_c_str()) + .filename() + .string(); if (rootName.empty()) continue; @@ -490,293 +489,6 @@ mi::neuraylib::ICompiled_material *Core::getCompiledMaterial( return compiledMaterial; } -namespace { - -// Resolve `let`-block indirection: a compiled sub-expression may reference a -// temporary slot instead of holding the value; without this, literal-bodied -// emitters would silently classify as non-constant. -mi::base::Handle derefTemporaries( - const mi::neuraylib::ICompiled_material *compiledMaterial, - mi::base::Handle expr) -{ - using namespace mi::neuraylib; - using mi::base::make_handle; - while (expr && expr->get_kind() == IExpression::EK_TEMPORARY) { - auto temporary = - make_handle(expr->get_interface()); - expr = make_handle(compiledMaterial->get_temporary(temporary->get_index())); - } - return expr; -} - -// Multiply a constant color/float factor into `scale` componentwise. -bool foldColorFactor( - const mi::neuraylib::IValue *value, std::array &scale) -{ - using namespace mi::neuraylib; - using mi::base::make_handle; - if (!value) - return false; - if (value->get_kind() == IValue::VK_COLOR) { - auto color = make_handle(value->get_interface()); - for (int i = 0; i < 3; ++i) { - auto channel = make_handle(color->get_value(i)); - if (!channel) - return false; - auto f = make_handle(channel->get_interface()); - if (!f) - return false; - scale[i] *= f->get_value(); - } - return true; - } - if (value->get_kind() == IValue::VK_FLOAT) { - const float f = - make_handle(value->get_interface())->get_value(); - for (auto &c : scale) - c *= f; - return true; - } - return false; -} - -struct DynamicIntensity -{ - std::array scale{1.f, 1.f, 1.f}; - Core::EmissionClassification::DynamicSource source = - Core::EmissionClassification::DynamicSource::None; - std::string argumentName; -}; - -// Multiplicative walk of a non-folding intensity expression: accumulate -// constant factors into `scale` and identify at most ONE dynamic factor — a -// color/float parameter or the `tex` of a tex::lookup_color. Anything outside -// that shape (sums, other calls, two dynamic factors, a multiply chain deeper -// than any sane authoring) fails the walk and the material keeps the -// unit-proxy Pick Power. -constexpr int kMaxIntensityWalkDepth = 16; - -bool walkIntensityFactors( - const mi::neuraylib::ICompiled_material *compiledMaterial, - mi::base::Handle expr, - DynamicIntensity &out, - int depth = 0) -{ - using namespace mi::neuraylib; - using mi::base::make_handle; - using DynamicSource = Core::EmissionClassification::DynamicSource; - - if (depth > kMaxIntensityWalkDepth) - return false; - expr = derefTemporaries(compiledMaterial, expr); - if (!expr) - return false; - - switch (expr->get_kind()) { - case IExpression::EK_CONSTANT: { - auto constant = - make_handle(expr->get_interface()); - return foldColorFactor(make_handle(constant->get_value()).get(), out.scale); - } - case IExpression::EK_PARAMETER: { - if (out.source != DynamicSource::None) - return false; - auto param = - make_handle(expr->get_interface()); - const char *name = compiledMaterial->get_parameter_name(param->get_index()); - if (!name) - return false; - out.source = DynamicSource::Parameter; - out.argumentName = name; - return true; - } - case IExpression::EK_DIRECT_CALL: { - auto call = - make_handle(expr->get_interface()); - const char *definition = call ? call->get_definition() : nullptr; - if (!definition) - return false; - const std::string_view def(definition); - auto args = make_handle(call->get_arguments()); - if (!args) - return false; - // Exact DB-name prefixes, same masquerade guard as the EDF check above. - if (def.rfind("mdl::operator*(", 0) == 0) { - if (args->get_size() != 2) - return false; - return walkIntensityFactors(compiledMaterial, - make_handle(args->get_expression(mi::Size(0))), - out, - depth + 1) - && walkIntensityFactors(compiledMaterial, - make_handle(args->get_expression(mi::Size(1))), - out, - depth + 1); - } - if (def.rfind("mdl::color(float)", 0) == 0) { - // Single-float color constructor only: color(r,g,b) factors would - // wrongly multiply as three scalars. - if (args->get_size() != 1) - return false; - return walkIntensityFactors(compiledMaterial, - make_handle(args->get_expression(mi::Size(0))), - out, - depth + 1); - } - // The lookup's coord/crop/wrap arguments are deliberately ignored: the - // recipe's mean is TEXTURE-domain (the bound sampler's full-image mean), - // an approximation of the surface mean — variance-only, never bias. - if (def.rfind("mdl::tex::lookup_color(", 0) == 0) { - if (out.source != DynamicSource::None) - return false; - auto tex = derefTemporaries( - compiledMaterial, make_handle(args->get_expression("tex"))); - if (!tex || tex->get_kind() != IExpression::EK_PARAMETER) - return false; - auto param = - make_handle(tex->get_interface()); - const char *name = - compiledMaterial->get_parameter_name(param->get_index()); - if (!name) - return false; - out.source = DynamicSource::Texture; - out.argumentName = name; - return true; - } - return false; - } - default: - return false; - } -} - -} // namespace - -Core::EmissionClassification Core::classifyEmission( - const mi::neuraylib::ICompiled_material *compiledMaterial) -{ - using namespace mi::neuraylib; - using mi::base::make_handle; - - EmissionClassification result; - - // Author-declared emission is a direct call to df::diffuse_edf; the default - // edf() compiles to a constant invalid-df. Only the diffuse EDF has uniform - // radiance over the hemisphere, matching the double-sided Geometry Light - // sampler and the synthetic next-event hit (ADR 0006's fidelity scope). - auto edf = derefTemporaries(compiledMaterial, - make_handle(compiledMaterial->lookup_sub_expression( - "surface.emission.emission"))); - if (!edf || edf->get_kind() != IExpression::EK_DIRECT_CALL) - return result; - { - auto call = - make_handle(edf->get_interface()); - const char *definition = call ? call->get_definition() : nullptr; - // Exact prefix match on the DB name of the elemental EDF, so no user - // module (::somepdf::, ::pkg::df::) can masquerade as it. - constexpr std::string_view DIFFUSE_EDF_PREFIX = "mdl::df::diffuse_edf("; - if (!definition - || std::string_view(definition).rfind(DIFFUSE_EDF_PREFIX, 0) != 0) - return result; - } - - // Only the (default) radiant-exitance intensity mode is handled; power mode - // needs area normalization the host cannot do here. - constexpr mi::Sint32 INTENSITY_RADIANT_EXITANCE = 0; // ::df::intensity_mode - auto mode = derefTemporaries(compiledMaterial, - make_handle( - compiledMaterial->lookup_sub_expression("surface.emission.mode"))); - if (mode) { - if (mode->get_kind() != IExpression::EK_CONSTANT) - return result; - auto constant = - make_handle(mode->get_interface()); - auto value = make_handle(constant->get_value()); - if (value->get_kind() != IValue::VK_ENUM) - return result; - if (make_handle(value->get_interface())->get_value() - != INTENSITY_RADIANT_EXITANCE) - return result; - } - - result.isDiffuseEmission = true; - - // Body-literal intensity folds to a constant; anything argument-, texture- or - // state-driven stays symbolic under class compilation and has no host value. - // A symbolic single-factor shape still yields a dynamic recipe below, so the - // Pick Power can track the live argument instead of the unit proxy. - constexpr float INV_PI = 0.31830988618379067154f; - auto intensity = derefTemporaries(compiledMaterial, - make_handle(compiledMaterial->lookup_sub_expression( - "surface.emission.intensity"))); - if (!intensity) - return result; - if (intensity->get_kind() != IExpression::EK_CONSTANT) { - DynamicIntensity dyn; - if (walkIntensityFactors(compiledMaterial, intensity, dyn) - && dyn.source != EmissionClassification::DynamicSource::None) { - // Radiance domain (= intensity scale / PI). Non-finite OR negative - // folded factors keep the proxy: clamping a negative scale to zero - // would zero the Pick Power while the device could still emit - // (negative scale x negative argument), making the light NEE-dead — - // exactly the truncation dimming this recipe exists to avoid. - bool usable = true; - for (float &c : dyn.scale) { - if (!std::isfinite(c) || c < 0.0f) { - usable = false; - break; - } - c *= INV_PI; - } - if (usable) { - result.dynamicSource = dyn.source; - result.dynamicArgumentName = std::move(dyn.argumentName); - result.dynamicScale = dyn.scale; - } - } - return result; - } - - auto constant = - make_handle(intensity->get_interface()); - auto value = make_handle(constant->get_value()); - std::array rgb{}; - if (value->get_kind() == IValue::VK_COLOR) { - auto color = make_handle(value->get_interface()); - for (int i = 0; i < 3; ++i) { - auto channel = make_handle(color->get_value(i)); - auto f = make_handle(channel->get_interface()); - if (!f) - return result; // not a plain float channel: treat as not host-known - rgb[i] = f->get_value(); - } - } else if (value->get_kind() == IValue::VK_FLOAT) { - const float f = - make_handle(value->get_interface())->get_value(); - rgb = {f, f, f}; - } else { - return result; - } - - // Emitted radiance = intensity / PI: the device emission callable returns - // edf * intensity and a diffuse EDF's value is 1/PI. Storing the unfolded - // intensity would overweight this emitter PI x in the light-pick CDF. - // A non-finite channel disqualifies the material entirely — clearing the - // diffuse flag too, or the textured branch would make it sampleable and - // next-event estimation would spray the NaN/Inf to every receiver the pick - // selects (one poisoned pick per sample). Negatives are clamped. - for (float &c : rgb) { - if (!std::isfinite(c)) { - result.isDiffuseEmission = false; - return result; - } - c = std::max(c, 0.0f) * INV_PI; - } - result.constantRadiance = rgb; - return result; -} - mi::neuraylib::ICompiled_material *Core::getDistilledToDiffuse( const mi::neuraylib::ICompiled_material *compiledMaterial) { @@ -823,8 +535,8 @@ const mi::neuraylib::ITarget_code *Core::getPtxTargetCode( ptxBackend->set_option("opt_level", "2"); ptxBackend->set_option("enable_exceptions", "off"); - // Generate init, surface scattering, surface emission (emission/intensity/mode), - // volume scattering and cutout opacity. + // Generate init, surface scattering, surface emission + // (emission/intensity/mode), volume scattering and cutout opacity. static mi::neuraylib::Target_function_description materialFunctions[] = { {"init", "mdlInit"}, {"thin_walled", "mdlThinWalled"}, diff --git a/devices/rtx/libmdl/Core.h b/devices/rtx/libmdl/Core.h index 06e5cec22..ddf295def 100644 --- a/devices/rtx/libmdl/Core.h +++ b/devices/rtx/libmdl/Core.h @@ -41,43 +41,6 @@ namespace visrtx::libmdl { class Core { public: - // Host-side emission classification of a compiled material, per ADR 0006. - // Computed once at compile time (the compiled material is not retained), so - // an Emissive Surface can be synthesized into a Geometry Light without - // recompiling or resolving class-compiled arguments. - struct EmissionClassification - { - // Emitted radiance (= folded intensity / PI, a diffuse EDF's value being - // 1/PI) when `surface.emission.intensity` folds to a body-literal constant; - // nullopt when it does not (texture / procedural / parameter-driven — not - // host-knowable under class compilation). - std::optional> constantRadiance; - // A diffuse radiant-exitance emission EDF is present AND eligible — a - // non-finite folded constant clears it, disqualifying the material. - // Invariant: constantRadiance.has_value() implies isDiffuseEmission. - bool isDiffuseEmission{false}; - - // Dynamic mean-radiance recipe: when the intensity does not fold but is a - // SINGLE class-compilation argument (color/float parameter, or the `tex` - // of a tex::lookup_color) times folded constants, the host can still - // compute a live mean radiance at light-build time: - // mean = * dynamicScale - // dynamicScale is radiance-domain (the folded constants already carry the - // diffuse EDF's 1/PI). Without a recipe the Pick Power falls back to the - // unit proxy — unbiased, but under-picking a bright emitter turns the - // firefly clamp and the last-depth MIS truncation into visible dimming - // next to correctly-powered lights. - enum class DynamicSource - { - None, - Parameter, - Texture, - }; - DynamicSource dynamicSource{DynamicSource::None}; - std::string dynamicArgumentName; - std::array dynamicScale{}; - }; - // The main neuray interface can only be acquired once. Possibly get it // as a parameter instead of allocating it internally. // Note that we allow overriding the logger only if we own the @@ -100,7 +63,8 @@ class Core // Load an MDL module from in-memory source into `transaction` and return it. // Returns null on failure (diagnostics are logged). Mirrors loadModule. - const mi::neuraylib::IModule *loadModuleFromString(std::string_view moduleName, + const mi::neuraylib::IModule *loadModuleFromString( + std::string_view moduleName, std::string_view moduleSource, mi::neuraylib::ITransaction *transaction); @@ -170,9 +134,6 @@ class Core const mi::neuraylib::IFunction_definition *, bool classCompilation = true); - static EmissionClassification classifyEmission( - const mi::neuraylib::ICompiled_material *compiledMaterial); - mi::neuraylib::ICompiled_material *getDistilledToDiffuse( const mi::neuraylib::ICompiled_material *compiledMaterial); diff --git a/devices/rtx/libmdl/EmissionDescriptor.h b/devices/rtx/libmdl/EmissionDescriptor.h new file mode 100644 index 000000000..7cb497e01 --- /dev/null +++ b/devices/rtx/libmdl/EmissionDescriptor.h @@ -0,0 +1,91 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// The emission descriptor — the classifier's output and the seam between +// MDL-pure analysis and the renderer. It is SDK-free (no mi:: types) so the +// device/renderer consume it without the MDL SDK. It DESCRIBES; a thin +// renderer-side policy decides registration. See ADR 0007. + +#include +#include + +namespace visrtx::libmdl { + +// Three-valued emission verdict for one material side. +enum class EmissionVerdict : std::uint8_t +{ + ProvablyNull, // identically zero at the current snapshot — never an emitter + ProvablyEmissive, // provably nonzero somewhere + Unknown, // could be either (register-eligible if faithful; worst case perf) +}; + +// EDF kinds present in a slot, as a bitmask. `Unknown` marks an unmodeled EDF +// leaf — described, never registered. +enum class EdfKind : std::uint8_t +{ + None = 0, + Diffuse = 1 << 0, + Spot = 1 << 1, + Directional = 1 << 2, + Measured = 1 << 3, + Unknown = 1 << 4, +}; + +inline EdfKind operator|(EdfKind a, EdfKind b) +{ + return EdfKind(std::uint8_t(a) | std::uint8_t(b)); +} +inline EdfKind &operator|=(EdfKind &a, EdfKind b) +{ + a = a | b; + return a; +} +inline bool hasKind(EdfKind set, EdfKind k) +{ + return (std::uint8_t(set) & std::uint8_t(k)) != 0; +} +// True iff every kind in `set` is contained in `allowed` (the ⊆ test), and the +// set is non-empty. +inline bool isSubsetOf(EdfKind set, EdfKind allowed) +{ + return set != EdfKind::None + && (std::uint8_t(set) & ~std::uint8_t(allowed)) == 0; +} + +enum class IntensityMode : std::uint8_t +{ + RadiantExitance, // radiance = intensity / PI (the faithful, default mode) + Power, // total power over area — not faithfully handled yet +}; + +enum class EmissionSign : std::uint8_t +{ + ProvablyNonnegative, // no reachable emission value can be negative + Unknown, // possibly negative — register would drop the NEE term, so + // forward-only +}; + +struct SlotDescriptor +{ + EmissionVerdict verdict{EmissionVerdict::ProvablyNull}; + EdfKind edfKinds{EdfKind::None}; + // Non-negative mean-radiance proxy (meanPositive · folded scale), per + // channel; feeds Pick Power only, never gates zero. Zero is a valid value. + std::array magnitude{}; + IntensityMode mode{IntensityMode::RadiantExitance}; + // The emission (EDF or intensity) reads a geometric-state quantity the + // synthetic next-event hit fabricates (normal/tangent/object-id/position); + // registering it would evaluate a different integrand than the forward hit. + bool dependsOnGeometricState{false}; + EmissionSign sign{EmissionSign::Unknown}; +}; + +struct EmissionDescriptor +{ + SlotDescriptor surface; + SlotDescriptor backface; +}; + +} // namespace visrtx::libmdl diff --git a/devices/rtx/libmdl/EmissionFold.cpp b/devices/rtx/libmdl/EmissionFold.cpp new file mode 100644 index 000000000..3df633ba1 --- /dev/null +++ b/devices/rtx/libmdl/EmissionFold.cpp @@ -0,0 +1,727 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#include "EmissionFold.h" + +#include +#include + +namespace visrtx::libmdl { + +namespace { + +constexpr float INV_PI = 0.31830988618379067154f; + +// Geometric-state quantities the synthetic next-event hit fabricates, so +// emission that reads them is not faithfully NEE-evaluable (ADR 0007). NOT +// texture_coordinate: uv-driven (textured) emission is the shipped sampleable +// case, evaluated via the sampler mean; nor the globally-uniform state +// (animation_time, wavelengths, scene units) which the synthetic hit +// reproduces. +bool isFabricatedState(Semantic s) +{ + switch (s) { + case Semantic::DS_INTRINSIC_STATE_POSITION: + case Semantic::DS_INTRINSIC_STATE_NORMAL: + case Semantic::DS_INTRINSIC_STATE_GEOMETRY_NORMAL: + case Semantic::DS_INTRINSIC_STATE_MOTION: + case Semantic::DS_INTRINSIC_STATE_TEXTURE_TANGENT_U: + case Semantic::DS_INTRINSIC_STATE_TEXTURE_TANGENT_V: + case Semantic::DS_INTRINSIC_STATE_TANGENT_SPACE: + case Semantic::DS_INTRINSIC_STATE_GEOMETRY_TANGENT_U: + case Semantic::DS_INTRINSIC_STATE_GEOMETRY_TANGENT_V: + case Semantic::DS_INTRINSIC_STATE_DIRECTION: + case Semantic::DS_INTRINSIC_STATE_TRANSFORM: + case Semantic::DS_INTRINSIC_STATE_TRANSFORM_POINT: + case Semantic::DS_INTRINSIC_STATE_TRANSFORM_VECTOR: + case Semantic::DS_INTRINSIC_STATE_TRANSFORM_NORMAL: + case Semantic::DS_INTRINSIC_STATE_TRANSFORM_SCALE: + case Semantic::DS_INTRINSIC_STATE_ROUNDED_CORNER_NORMAL: + case Semantic::DS_INTRINSIC_STATE_OBJECT_ID: + return true; + default: + return false; + } +} + +EdfKind edfKindOf(Semantic s) +{ + switch (s) { + case Semantic::DS_INTRINSIC_DF_DIFFUSE_EDF: + return EdfKind::Diffuse; + case Semantic::DS_INTRINSIC_DF_SPOT_EDF: + return EdfKind::Spot; + case Semantic::DS_INTRINSIC_DF_MEASURED_EDF: + return EdfKind::Measured; + default: + return EdfKind::None; + } +} + +bool isMixSemantic(Semantic s) +{ + switch (s) { + case Semantic::DS_INTRINSIC_DF_NORMALIZED_MIX: + case Semantic::DS_INTRINSIC_DF_CLAMPED_MIX: + case Semantic::DS_INTRINSIC_DF_UNBOUNDED_MIX: + case Semantic::DS_INTRINSIC_DF_COLOR_NORMALIZED_MIX: + case Semantic::DS_INTRINSIC_DF_COLOR_CLAMPED_MIX: + case Semantic::DS_INTRINSIC_DF_COLOR_UNBOUNDED_MIX: + return true; + default: + return false; + } +} + +enum class Tri : std::uint8_t +{ + False, + True, + Unknown +}; + +// Abstract value of a scalar/color sub-expression. +struct Scalar +{ + Tri zero{Tri::Unknown}; // is the value identically zero? + bool magnitudeKnown{false}; + std::array magnitude{}; // meanPositive per channel (>= 0), if known + EmissionSign sign{EmissionSign::Unknown}; + bool dependsOnState{false}; + bool finite{true}; +}; + +// Abstract value of an EDF sub-expression. +struct Edf +{ + Tri null{Tri::Unknown}; // is the EDF the null (non-emitting) EDF? + EdfKind kinds{EdfKind::None}; + bool dependsOnState{false}; +}; + +std::array positivePart(const std::array &v) +{ + return {std::max(v[0], 0.0f), std::max(v[1], 0.0f), std::max(v[2], 0.0f)}; +} + +bool allZero(const std::array &v) +{ + return v[0] == 0.0f && v[1] == 0.0f && v[2] == 0.0f; +} + +bool anyNonzero(const std::array &v) +{ + return v[0] != 0.0f || v[1] != 0.0f || v[2] != 0.0f; +} + +bool allNonnegative(const std::array &v) +{ + return v[0] >= 0.0f && v[1] >= 0.0f && v[2] >= 0.0f; +} + +bool allFinite(const std::array &v) +{ + return std::isfinite(v[0]) && std::isfinite(v[1]) && std::isfinite(v[2]); +} + +class Fold +{ + public: + Fold(const EmissionIR &ir, const EmissionValueSource &values) + : m_ir(ir), m_values(values) + {} + + SlotDescriptor foldSlot(const EmissionSlotIR &slot) + { + SlotDescriptor desc; + if (slot.edfRoot < 0) { + desc.verdict = EmissionVerdict::ProvablyNull; + return desc; + } + + const Edf edf = evalEdf(slot.edfRoot, 0); + Scalar intensity; + if (slot.intensityRoot >= 0) + intensity = evalScalar(slot.intensityRoot, 0); + else + intensity.zero = Tri::True; // no intensity ⇒ no emission + + desc.edfKinds = edf.kinds; + desc.mode = foldMode(slot.modeRoot); + desc.dependsOnGeometricState = + edf.dependsOnState || intensity.dependsOnState; + + // verdict = nullness of (EDF · intensity): null iff EDF null OR intensity + // 0. + const Tri emissionNull = + (edf.null == Tri::True || intensity.zero == Tri::True) + ? Tri::True + : ((edf.null == Tri::False && intensity.zero == Tri::False) + ? Tri::False + : Tri::Unknown); + if (!intensity.finite) + desc.verdict = EmissionVerdict::Unknown; + else if (emissionNull == Tri::True) + desc.verdict = EmissionVerdict::ProvablyNull; + else if (emissionNull == Tri::False) + desc.verdict = EmissionVerdict::ProvablyEmissive; + else + desc.verdict = EmissionVerdict::Unknown; + + // magnitude = meanPositive(intensity) / PI (diffuse EDF value 1/PI). A unit + // proxy stands in when the intensity magnitude is not host-known — + // unbiased, it only weights the Light Pick. + if (intensity.magnitudeKnown) { + for (int i = 0; i < 3; ++i) + desc.magnitude[i] = intensity.magnitude[i] * INV_PI; + } else { + desc.magnitude = {1.0f, 1.0f, 1.0f}; + } + + // sign gates registration; the EDF itself is nonnegative. A non-finite + // intensity fails the sign gate (so it never registers an inf/NaN Pick + // Power) — its light still arrives via the forward path. + desc.sign = intensity.finite ? intensity.sign : EmissionSign::Unknown; + return desc; + } + + private: + const EmissionNode &node(int i) const + { + return m_ir.nodes[std::size_t(i)]; + } + + IntensityMode foldMode(int modeRoot) const + { + if (modeRoot < 0) + return IntensityMode::RadiantExitance; // default + const auto &n = node(modeRoot); + if (n.kind == EmissionNodeKind::Constant + && n.constantKind == ConstantKind::Enum) + return n.intValue == 0 ? IntensityMode::RadiantExitance + : IntensityMode::Power; + // Non-constant mode: cannot prove radiant-exitance ⇒ treat as power (not + // registered) rather than silently assume the faithful mode. + return IntensityMode::Power; + } + + Edf evalEdf(int index, int depth) + { + Edf r; + if (index < 0 || depth > 64) + return r; // Unknown + const auto &n = node(index); + + if (n.kind == EmissionNodeKind::Constant) { + if (n.constantKind == ConstantKind::InvalidDf) { + r.null = Tri::True; + return r; + } + return r; // a non-df constant in EDF position: Unknown + } + if (n.kind != EmissionNodeKind::Call) { + r.kinds = EdfKind::Unknown; + return r; + } + + const EdfKind leaf = edfKindOf(n.semantic); + if (leaf != EdfKind::None) { + r.kinds = leaf; + r.null = Tri::False; // a present emissive leaf + // Scan the leaf's own arguments (tint/roughness/exponent/normal/...) for + // fabricated-state reads: a state-dependent EDF parameter is unfaithful + // at the synthetic hit even when the leaf KIND is faithful — e.g. + // diffuse_edf(tint: color(state::object_id())) varies radiance by the + // object id the synthetic hit fakes to 0. The mix/tint/directional paths + // already scan their operands; the leaf branch must too. + for (int op : n.operands) + r.dependsOnState = + r.dependsOnState || evalScalar(op, depth + 1).dependsOnState; + return r; + } + + switch (n.semantic) { + case Semantic::DS_INTRINSIC_DF_TINT: { + // tint(color tint, edf base): null iff tint zero OR base null. + Edf base; + Scalar tint; + for (int op : n.operands) { + if (isEdfNode(op)) + base = evalEdf(op, depth + 1); + else + tint = evalScalar(op, depth + 1); + } + r.kinds = base.kinds; + r.dependsOnState = base.dependsOnState || tint.dependsOnState; + if (tint.zero == Tri::True || base.null == Tri::True) + r.null = Tri::True; + else if (tint.zero == Tri::False && base.null == Tri::False) + r.null = Tri::False; + else + r.null = Tri::Unknown; + return r; + } + case Semantic::DS_INTRINSIC_DF_DIRECTIONAL_FACTOR: { + // Directional emission: kind = directional + base kinds; null iff base + // null or both endpoint tints cross-channel zero. Endpoints are the two + // color operands; base is the edf operand. + Edf base; + bool sawEdf = false; + Tri anyTintNonzero = Tri::False; + bool allTintZeroKnown = true; + for (int op : n.operands) { + if (isEdfNode(op)) { + base = evalEdf(op, depth + 1); + sawEdf = true; + r.dependsOnState = r.dependsOnState || base.dependsOnState; + } else { + Scalar tint = evalScalar(op, depth + 1); + r.dependsOnState = r.dependsOnState || tint.dependsOnState; + if (tint.zero != Tri::True) + allTintZeroKnown = false; + if (tint.zero == Tri::False) + anyTintNonzero = Tri::True; + } + } + r.kinds = EdfKind::Directional; + if (sawEdf) + r.kinds |= base.kinds; + if (base.null == Tri::True || allTintZeroKnown) + r.null = Tri::True; + else if (base.null == Tri::False && anyTintNonzero == Tri::True) + r.null = Tri::False; + else + r.null = Tri::Unknown; + return r; + } + default: + break; + } + + if (isMixSemantic(n.semantic)) { + // A mix's single operand is a df_component[] array, whose elements are + // component constructors wrapping {weight, edf}. The EDF leaves nest + // several levels down through array/struct constructors, NOT as direct + // operands — so deep-scan for reachable EDF-semantic nodes to union their + // kinds. Never prove Null here (the array shape makes an all-null proof + // unreliable); an Unknown verdict still registers a faithful all-diffuse + // mix, and mis-registering a null mix only wastes a pick slot (unbiased). + collectEdfKindsDeep(index, r, depth); + if (r.kinds == EdfKind::None) + r.kinds = EdfKind::Unknown; + r.null = Tri::Unknown; + return r; + } + + // Unmodeled df:: call. + r.kinds = EdfKind::Unknown; + return r; + } + + // Deep-scan for EDF leaves reachable through array/struct constructors (the + // shape a df::*_mix's component array takes), unioning their kinds and any + // fabricated-state reads. Idempotent over the DAG; depth-bounded. + void collectEdfKindsDeep(int index, Edf &acc, int depth) + { + if (index < 0 || depth > 128) + return; + const auto &n = node(index); + if (n.kind == EmissionNodeKind::Call) { + if (isFabricatedState(n.semantic)) + acc.dependsOnState = true; + const EdfKind leaf = edfKindOf(n.semantic); + acc.kinds |= leaf; + if (n.semantic == Semantic::DS_INTRINSIC_DF_DIRECTIONAL_FACTOR) + acc.kinds |= EdfKind::Directional; + // A df:: intrinsic in the EDF tree that is neither a modeled leaf nor a + // modeled combinator is an EMISSIVE distribution we don't understand — + // poison to Unknown so a modeled diffuse sibling can't mask it and let + // the mix register as pure diffuse. (Scalar operands like weights are not + // in the DF range, so they don't poison.) + if (leaf == EdfKind::None && isUnmodeledDfIntrinsic(n.semantic)) + acc.kinds |= EdfKind::Unknown; + for (int op : n.operands) + collectEdfKindsDeep(op, acc, depth + 1); + } else if (n.kind == EmissionNodeKind::Texture) { + for (int op : n.operands) + collectEdfKindsDeep(op, acc, depth + 1); + } else if (n.kind == EmissionNodeKind::Opaque) { + acc.kinds |= EdfKind::Unknown; // an unresolved node under the mix + } + } + + // A DF-intrinsic semantic (df::* function) that this fold does not model as a + // leaf or a known combinator (tint/directional_factor/mix families). + static bool isUnmodeledDfIntrinsic(Semantic s) + { + if (s < Semantic::DS_INTRINSIC_DF_FIRST + || s > Semantic::DS_INTRINSIC_DF_LAST) + return false; + if (edfKindOf(s) != EdfKind::None || isMixSemantic(s)) + return false; + switch (s) { + case Semantic::DS_INTRINSIC_DF_TINT: + case Semantic::DS_INTRINSIC_DF_DIRECTIONAL_FACTOR: + return false; // modeled combinators + default: + return true; + } + } + + bool isEdfNode(int index) const + { + if (index < 0) + return false; + const auto &n = node(index); + if (n.kind == EmissionNodeKind::Constant) + return n.constantKind == ConstantKind::InvalidDf; + if (n.kind == EmissionNodeKind::Call) + return edfKindOf(n.semantic) != EdfKind::None + || n.semantic == Semantic::DS_INTRINSIC_DF_TINT + || n.semantic == Semantic::DS_INTRINSIC_DF_DIRECTIONAL_FACTOR + || isMixSemantic(n.semantic); + return false; + } + + Scalar evalScalar(int index, int depth) + { + Scalar r; + if (index < 0 || depth > 64) + return r; // Unknown + const auto &n = node(index); + + switch (n.kind) { + case EmissionNodeKind::Constant: + return scalarFromConstant(n); + case EmissionNodeKind::Parameter: + return scalarFromParameter(n); + case EmissionNodeKind::Texture: + return scalarFromTexture(n); + case EmissionNodeKind::Call: + return scalarFromCall(n, depth); + default: + return r; // Opaque ⇒ Unknown + } + } + + Scalar scalarFromConstant(const EmissionNode &n) + { + Scalar r; + if (n.constantKind == ConstantKind::Color + || n.constantKind == ConstantKind::Float) { + r.finite = allFinite(n.value); + r.zero = allZero(n.value) ? Tri::True + : anyNonzero(n.value) ? Tri::False + : Tri::Unknown; + r.magnitudeKnown = true; + r.magnitude = positivePart(n.value); + r.sign = allNonnegative(n.value) ? EmissionSign::ProvablyNonnegative + : EmissionSign::Unknown; + return r; + } + // Bool/int/enum in a radiance context: leave Unknown (should not occur). + return r; + } + + Scalar scalarFromParameter(const EmissionNode &n) + { + Scalar r; + std::array value; + if (m_values.color(n.parameterName, value)) { + r.finite = allFinite(value); + r.zero = allZero(value) ? Tri::True + : anyNonzero(value) ? Tri::False + : Tri::Unknown; + r.magnitudeKnown = true; + r.magnitude = positivePart(value); + r.sign = allNonnegative(value) ? EmissionSign::ProvablyNonnegative + : EmissionSign::Unknown; + } + // Unknown value ⇒ all Unknown (default), sign Unknown, magnitude unknown. + return r; + } + + Scalar scalarFromTexture(const EmissionNode &n) + { + Scalar r; + // A coord/wrap/crop argument that reads a fabricated geometric-state + // quantity makes the lookup unfaithful at the synthetic hit. + for (int op : n.operands) { + Scalar s = evalScalar(op, 1); + r.dependsOnState = r.dependsOnState || s.dependsOnState; + r.finite = r.finite && s.finite; + } + + ResourceStats stats; + const bool known = n.parameterIndex >= 0 + ? m_values.resourceByParam(n.parameterName, stats) + : m_values.resourceByName(n.resourceName, stats); + if (!known) + return r; // Unknown (state-dependence preserved) + + if (!stats.valid) { + // Unbound/invalid texture ⇒ lookup folds to 0. + r.zero = Tri::True; + r.magnitudeKnown = true; + r.magnitude = {0.0f, 0.0f, 0.0f}; + r.sign = EmissionSign::ProvablyNonnegative; + return r; + } + r.finite = stats.finite; + // Zero only provable in sampler-output space (T(0)==0). + if (stats.transferPreservesZero && allZero(stats.maxAbs)) + r.zero = Tri::True; + else + r.zero = Tri::Unknown; + r.magnitudeKnown = true; + r.magnitude = stats.meanPositive; + r.sign = allNonnegative(stats.minValue) ? EmissionSign::ProvablyNonnegative + : EmissionSign::Unknown; + return r; + } + + Scalar scalarFromCall(const EmissionNode &n, int depth) + { + // State-dependence propagates regardless of whether the op is modeled. + bool stateFromThis = isFabricatedState(n.semantic); + + switch (n.semantic) { + case Semantic::DS_MULTIPLY: + return combineMul(n, depth, stateFromThis); + case Semantic::DS_PLUS: + return combineAdd(n, depth, stateFromThis); + case Semantic::DS_MINUS: + return combineSub(n, depth, stateFromThis); + case Semantic::DS_TERNARY: + return combineTernary(n, depth, stateFromThis); + case Semantic::DS_CONV_CONSTRUCTOR: + case Semantic::DS_ELEM_CONSTRUCTOR: + case Semantic::DS_CONV_OPERATOR: + case Semantic::DS_COPY_CONSTRUCTOR: + return combineConstructor(n, depth, stateFromThis); + default: + break; + } + + // Unmodeled call: Unknown value, but still scan operands for state and + // finiteness so the flags stay sound. + Scalar r; + r.dependsOnState = stateFromThis; + for (int op : n.operands) { + Scalar s = evalScalar(op, depth + 1); + r.dependsOnState = r.dependsOnState || s.dependsOnState; + } + return r; + } + + Scalar combineMul(const EmissionNode &n, int depth, bool stateFromThis) + { + Scalar r; + r.zero = Tri::True; // identity for the running product's zero test + r.magnitudeKnown = true; + r.magnitude = {1.0f, 1.0f, 1.0f}; + r.sign = EmissionSign::ProvablyNonnegative; + r.dependsOnState = stateFromThis; + bool any = false; + Tri zeroAcc = + Tri::False; // zero-absorbing: product zero iff any factor zero + for (int op : n.operands) { + Scalar s = evalScalar(op, depth + 1); + any = true; + r.dependsOnState = r.dependsOnState || s.dependsOnState; + r.finite = r.finite && s.finite; + if (s.zero == Tri::True) + zeroAcc = Tri::True; + else if (s.zero == Tri::Unknown && zeroAcc != Tri::True) + zeroAcc = Tri::Unknown; + r.magnitudeKnown = r.magnitudeKnown && s.magnitudeKnown; + if (s.magnitudeKnown) + for (int i = 0; i < 3; ++i) + r.magnitude[i] *= s.magnitude[i]; + if (s.sign != EmissionSign::ProvablyNonnegative) + r.sign = EmissionSign::Unknown; + } + if (!any) + return Scalar{}; + r.zero = zeroAcc; + if (!r.magnitudeKnown) + r.magnitude = {}; + return r; + } + + Scalar combineAdd(const EmissionNode &n, int depth, bool stateFromThis) + { + Scalar r; + r.magnitudeKnown = true; + r.magnitude = {0.0f, 0.0f, 0.0f}; + r.sign = EmissionSign::ProvablyNonnegative; + r.dependsOnState = stateFromThis; + bool allZeroKnown = true; + bool anyNonzero = false; + bool anyUnknownZero = false; + for (int op : n.operands) { + Scalar s = evalScalar(op, depth + 1); + r.dependsOnState = r.dependsOnState || s.dependsOnState; + r.finite = r.finite && s.finite; + if (s.zero != Tri::True) + allZeroKnown = false; + if (s.zero == Tri::False) + anyNonzero = true; + if (s.zero == Tri::Unknown) + anyUnknownZero = true; + r.magnitudeKnown = r.magnitudeKnown && s.magnitudeKnown; + if (s.magnitudeKnown) + for (int i = 0; i < 3; ++i) + r.magnitude[i] += s.magnitude[i]; + if (s.sign != EmissionSign::ProvablyNonnegative) + r.sign = EmissionSign::Unknown; + } + // sum zero iff all zero; nonzero needs nonnegativity (a negative could + // cancel a positive), so only claim Nonzero when the sum is provably so. + if (allZeroKnown) + r.zero = Tri::True; + else if (anyNonzero && !anyUnknownZero + && r.sign == EmissionSign::ProvablyNonnegative) + r.zero = Tri::False; + else + r.zero = Tri::Unknown; + if (!r.magnitudeKnown) + r.magnitude = {}; + return r; + } + + Scalar combineSub(const EmissionNode &n, int depth, bool stateFromThis) + { + Scalar r; + r.dependsOnState = stateFromThis; + // ProvablyZero only on exact node identity (a - a). Otherwise Unknown — + // never ProvablyNonzero (1 - w with w unknown can be zero). + if (n.operands.size() == 2 && n.operands[0] == n.operands[1]) { + r.zero = Tri::True; + r.magnitudeKnown = true; + r.magnitude = {0.0f, 0.0f, 0.0f}; + r.sign = EmissionSign::ProvablyNonnegative; + } + for (int op : n.operands) { + Scalar s = evalScalar(op, depth + 1); + r.dependsOnState = r.dependsOnState || s.dependsOnState; + r.finite = r.finite && s.finite; + } + return r; + } + + Scalar combineTernary(const EmissionNode &n, int depth, bool stateFromThis) + { + if (n.operands.size() != 3) + return Scalar{}; + Scalar cond = evalScalar(n.operands[0], depth + 1); + const bool condState = cond.dependsOnState; + + // Fold if the condition is a known constant bool. + const auto &condNode = node(n.operands[0]); + if (condNode.kind == EmissionNodeKind::Constant + && condNode.constantKind == ConstantKind::Bool) { + const int arm = condNode.boolValue ? 1 : 2; + Scalar r = evalScalar(n.operands[arm], depth + 1); + r.dependsOnState = r.dependsOnState || stateFromThis || condState; + return r; + } + + // Otherwise join the arms (LUB with Unknown as top). + Scalar a = evalScalar(n.operands[1], depth + 1); + Scalar b = evalScalar(n.operands[2], depth + 1); + Scalar r; + r.dependsOnState = + stateFromThis || condState || a.dependsOnState || b.dependsOnState; + r.finite = a.finite && b.finite; + r.zero = joinTri(a.zero, b.zero); + r.sign = (a.sign == EmissionSign::ProvablyNonnegative + && b.sign == EmissionSign::ProvablyNonnegative) + ? EmissionSign::ProvablyNonnegative + : EmissionSign::Unknown; + r.magnitudeKnown = a.magnitudeKnown && b.magnitudeKnown; + if (r.magnitudeKnown) + for (int i = 0; i < 3; ++i) + r.magnitude[i] = std::max(a.magnitude[i], b.magnitude[i]); + return r; + } + + Scalar combineConstructor( + const EmissionNode &n, int depth, bool stateFromThis) + { + // color(float) / color(r,g,b) / conversions: fold operands into a value + // when all are known constants; otherwise propagate zero/sign/state + // conservatively. + Scalar r; + r.dependsOnState = stateFromThis; + if (n.operands.empty()) + return r; + + // Single-operand broadcast (color(float), conversions). + if (n.operands.size() == 1) + return applyState(evalScalar(n.operands[0], depth + 1), stateFromThis); + + // Multi-channel constructor: combine per-channel scalars. + r.magnitudeKnown = true; + r.magnitude = {0.0f, 0.0f, 0.0f}; + r.sign = EmissionSign::ProvablyNonnegative; + Tri zeroAcc = Tri::True; + for (std::size_t i = 0; i < n.operands.size() && i < 3; ++i) { + Scalar s = evalScalar(n.operands[int(i)], depth + 1); + r.dependsOnState = r.dependsOnState || s.dependsOnState; + r.finite = r.finite && s.finite; + zeroAcc = zeroAcc == Tri::True ? s.zero : joinNonzero(zeroAcc, s.zero); + r.magnitudeKnown = r.magnitudeKnown && s.magnitudeKnown; + if (s.magnitudeKnown) + r.magnitude[i] = s.magnitude[0]; + if (s.sign != EmissionSign::ProvablyNonnegative) + r.sign = EmissionSign::Unknown; + } + r.zero = zeroAcc; + if (!r.magnitudeKnown) + r.magnitude = {}; + return r; + } + + static Scalar applyState(Scalar s, bool stateFromThis) + { + s.dependsOnState = s.dependsOnState || stateFromThis; + return s; + } + + static Tri joinTri(Tri a, Tri b) + { + if (a == b) + return a; + return Tri::Unknown; + } + + // For a channel accumulation where "all zero ⇒ zero, any nonzero ⇒ nonzero". + static Tri joinNonzero(Tri acc, Tri ch) + { + if (acc == Tri::False || ch == Tri::False) + return Tri::False; // a nonzero channel makes the color nonzero + if (acc == Tri::True && ch == Tri::True) + return Tri::True; + return Tri::Unknown; + } + + const EmissionIR &m_ir; + const EmissionValueSource &m_values; +}; + +} // namespace + +EmissionDescriptor foldEmissionDescriptor( + const EmissionIR &ir, const EmissionValueSource &values) +{ + EmissionDescriptor desc; + if (ir.empty()) + return desc; + Fold fold(ir, values); + desc.surface = fold.foldSlot(ir.surface); + desc.backface = fold.foldSlot(ir.backface); + return desc; +} + +} // namespace visrtx::libmdl diff --git a/devices/rtx/libmdl/EmissionFold.h b/devices/rtx/libmdl/EmissionFold.h new file mode 100644 index 000000000..198d3d2ec --- /dev/null +++ b/devices/rtx/libmdl/EmissionFold.h @@ -0,0 +1,69 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// Folds an emission IR against a value source into an EmissionDescriptor, using +// the three-valued abstract interpretation of ADR 0007. Pure and SDK-free: it +// walks the IR (which holds no SDK pointers) plus an abstract value source, so +// it runs device-free in tests and against live argument bytes on the device. + +#include "EmissionDescriptor.h" +#include "EmissionIR.h" +#include "ResourceStats.h" + +#include +#include + +namespace visrtx::libmdl { + +// Supplies current parameter values and resource stats to the fold, keyed by +// the class-compilation argument NAME (the device's argument block and samplers +// are name-keyed). The device backs it with live argument bytes + sampler +// reductions; tests back it with fakes. +class EmissionValueSource +{ + public: + virtual ~EmissionValueSource() = default; + + // Current value of a parameter, if known. A float scalar is broadcast to rgb. + // Returning false makes the parameter symbolic (Unknown), never zero. + virtual bool color( + const std::string ¶meterName, std::array &out) const = 0; + virtual bool boolean(const std::string ¶meterName, bool &out) const = 0; + + // Stats for a body-literal bound texture (by DB name) or an argument-bound + // texture (by parameter name). Returning false ⇒ Unknown. + virtual bool resourceByName( + const std::string &name, ResourceStats &out) const = 0; + virtual bool resourceByParam( + const std::string ¶meterName, ResourceStats &out) const = 0; +}; + +// A value source that knows nothing — every parameter/resource is Unknown. Used +// for the topology-only path and as a base for partial fakes. +class NullValueSource : public EmissionValueSource +{ + public: + bool color(const std::string &, std::array &) const override + { + return false; + } + bool boolean(const std::string &, bool &) const override + { + return false; + } + bool resourceByName(const std::string &, ResourceStats &) const override + { + return false; + } + bool resourceByParam(const std::string &, ResourceStats &) const override + { + return false; + } +}; + +EmissionDescriptor foldEmissionDescriptor( + const EmissionIR &ir, const EmissionValueSource &values); + +} // namespace visrtx::libmdl diff --git a/devices/rtx/libmdl/EmissionIR.cpp b/devices/rtx/libmdl/EmissionIR.cpp new file mode 100644 index 000000000..a670b5623 --- /dev/null +++ b/devices/rtx/libmdl/EmissionIR.cpp @@ -0,0 +1,309 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#include "EmissionIR.h" + +#include +#include +#include + +#include +#include +#include + +namespace visrtx::libmdl { + +namespace { + +using namespace mi::neuraylib; +using mi::base::make_handle; + +// Bound the walk so a pathological DAG cannot blow the stack; deeper than this +// no real emission graph goes, and anything truncated folds to Opaque/Unknown. +constexpr int kMaxWalkDepth = 64; + +// tex::lookup_color / _float3 / _float — the texture-sampling intrinsics whose +// first argument is the texture. Keyed on semantics below. +bool isTextureLookup(Semantic s) +{ + switch (s) { + case Semantic::DS_INTRINSIC_TEX_LOOKUP_FLOAT: + case Semantic::DS_INTRINSIC_TEX_LOOKUP_FLOAT2: + case Semantic::DS_INTRINSIC_TEX_LOOKUP_FLOAT3: + case Semantic::DS_INTRINSIC_TEX_LOOKUP_FLOAT4: + case Semantic::DS_INTRINSIC_TEX_LOOKUP_COLOR: + return true; + default: + return false; + } +} + +class Builder +{ + public: + Builder(const ICompiled_material *material, ITransaction *transaction) + : m_material(material), m_transaction(transaction) + {} + + EmissionIR build() + { + EmissionIR ir; + m_ir = &ir; + + ir.surface = buildSlot("surface"); + ir.backface = buildSlot("backface"); + ir.thinWalledRoot = buildPath("thin_walled"); + + collectDeps(ir); + m_ir = nullptr; + return ir; + } + + private: + EmissionSlotIR buildSlot(const char *side) + { + EmissionSlotIR slot; + slot.edfRoot = buildPath(std::string(side) + ".emission.emission"); + slot.intensityRoot = buildPath(std::string(side) + ".emission.intensity"); + slot.modeRoot = buildPath(std::string(side) + ".emission.mode"); + return slot; + } + + int buildPath(const std::string &path) + { + auto expr = make_handle(m_material->lookup_sub_expression(path.c_str())); + if (!expr) + return -1; + return buildExpr(expr, 0); + } + + // Resolve `let`-block temporaries and memoize on the temporary index so a + // subexpression shared through CSE maps to ONE node. + mi::base::Handle deref( + mi::base::Handle expr, int *sharedTemporary) + { + *sharedTemporary = -1; + while (expr && expr->get_kind() == IExpression::EK_TEMPORARY) { + auto tmp = + make_handle(expr->get_interface()); + *sharedTemporary = int(tmp->get_index()); + expr = make_handle(m_material->get_temporary(tmp->get_index())); + } + return expr; + } + + int buildExpr(mi::base::Handle expr, int depth) + { + int sharedTemporary = -1; + expr = deref(expr, &sharedTemporary); + if (!expr || depth > kMaxWalkDepth) + return addNode(EmissionNode{}); // Opaque + + if (sharedTemporary >= 0) { + auto it = m_temporaryNodes.find(sharedTemporary); + if (it != m_temporaryNodes.end()) + return it->second; + } + + const int nodeIndex = buildExprUncached(expr, depth); + if (sharedTemporary >= 0) + m_temporaryNodes.emplace(sharedTemporary, nodeIndex); + return nodeIndex; + } + + int buildExprUncached(mi::base::Handle expr, int depth) + { + switch (expr->get_kind()) { + case IExpression::EK_CONSTANT: + return buildConstant( + make_handle(expr->get_interface())); + case IExpression::EK_PARAMETER: + return buildParameter( + make_handle(expr->get_interface())); + case IExpression::EK_DIRECT_CALL: + return buildCall( + make_handle(expr->get_interface()), + depth); + default: + return addNode(EmissionNode{}); // Opaque + } + } + + int buildConstant(mi::base::Handle constant) + { + EmissionNode node; + node.kind = EmissionNodeKind::Constant; + auto value = make_handle(constant->get_value()); + if (!value) + return addNode(EmissionNode{}); + + switch (value->get_kind()) { + case IValue::VK_COLOR: { + auto color = make_handle(value->get_interface()); + node.constantKind = ConstantKind::Color; + for (int i = 0; i < 3; ++i) { + auto ch = make_handle(color->get_value(i)); + auto f = make_handle(ch->get_interface()); + if (!f) + return addNode(EmissionNode{}); // non-literal channel: Opaque + node.value[i] = f->get_value(); + } + break; + } + case IValue::VK_FLOAT: { + const float f = + make_handle(value->get_interface())->get_value(); + node.constantKind = ConstantKind::Float; + node.value = {f, f, f}; + break; + } + case IValue::VK_BOOL: + node.constantKind = ConstantKind::Bool; + node.boolValue = + make_handle(value->get_interface())->get_value(); + break; + case IValue::VK_INT: + node.constantKind = ConstantKind::Int; + node.intValue = + make_handle(value->get_interface())->get_value(); + break; + case IValue::VK_ENUM: + node.constantKind = ConstantKind::Enum; + node.intValue = + make_handle(value->get_interface())->get_value(); + break; + case IValue::VK_INVALID_DF: + node.constantKind = ConstantKind::InvalidDf; + break; + default: + return addNode(EmissionNode{}); // unmodeled value kind: Opaque + } + return addNode(std::move(node)); + } + + int buildParameter(mi::base::Handle param) + { + EmissionNode node; + node.kind = EmissionNodeKind::Parameter; + node.parameterIndex = int(param->get_index()); + const char *name = m_material->get_parameter_name(param->get_index()); + node.parameterName = name ? name : ""; + return addNode(std::move(node)); + } + + Semantic semanticOf(const IExpression_direct_call *call) + { + const char *definition = call->get_definition(); + if (!definition) + return Semantic::DS_UNKNOWN; + auto fn = + make_handle(m_transaction->access(definition)); + if (!fn) + return Semantic::DS_UNKNOWN; + return fn->get_semantic(); + } + + int buildCall(mi::base::Handle call, int depth) + { + const Semantic semantic = semanticOf(call.get()); + auto args = make_handle(call->get_arguments()); + if (!args) + return addNode(EmissionNode{}); + + if (isTextureLookup(semantic)) + return buildTexture(semantic, args, depth); + + EmissionNode node; + node.kind = EmissionNodeKind::Call; + node.semantic = semantic; + for (mi::Size i = 0; i < args->get_size(); ++i) { + node.operands.push_back( + buildExpr(make_handle(args->get_expression(i)), depth + 1)); + } + return addNode(std::move(node)); + } + + int buildTexture(Semantic semantic, + mi::base::Handle args, + int depth) + { + EmissionNode node; + node.kind = EmissionNodeKind::Texture; + node.semantic = semantic; + + int sharedTemporary = -1; + auto tex = + deref(make_handle(args->get_expression("tex")), &sharedTemporary); + if (tex && tex->get_kind() == IExpression::EK_PARAMETER) { + auto param = + make_handle(tex->get_interface()); + node.parameterIndex = int(param->get_index()); + const char *name = m_material->get_parameter_name(param->get_index()); + node.parameterName = name ? name : ""; + } else if (tex && tex->get_kind() == IExpression::EK_CONSTANT) { + auto constant = + make_handle(tex->get_interface()); + auto value = make_handle(constant->get_value()); + if (value && value->get_kind() == IValue::VK_TEXTURE) { + auto texValue = + make_handle(value->get_interface()); + const char *name = texValue->get_value(); + node.resourceName = name ? name : ""; + } + } + + // Build the non-`tex` arguments (coord, wrap, crop) as operands so the fold + // scans them for geometric-state reads: a coord driven by state::position + // or state::normal is unfaithful at the synthetic hit even though the + // sampler mean covers a texture_coordinate-driven lookup. + for (mi::Size i = 0; i < args->get_size(); ++i) { + const char *argName = args->get_name(i); + if (argName && std::string_view(argName) == "tex") + continue; + node.operands.push_back( + buildExpr(make_handle(args->get_expression(i)), depth + 1)); + } + return addNode(std::move(node)); + } + + int addNode(EmissionNode node) + { + m_ir->nodes.push_back(std::move(node)); + return int(m_ir->nodes.size()) - 1; + } + + void collectDeps(EmissionIR &ir) + { + for (const auto &node : ir.nodes) { + if (node.parameterIndex >= 0) + ir.emissionDeps.push_back(node.parameterIndex); + if (node.kind == EmissionNodeKind::Texture && !node.resourceName.empty()) + ir.resourceDeps.push_back(node.resourceName); + } + std::sort(ir.emissionDeps.begin(), ir.emissionDeps.end()); + ir.emissionDeps.erase( + std::unique(ir.emissionDeps.begin(), ir.emissionDeps.end()), + ir.emissionDeps.end()); + std::sort(ir.resourceDeps.begin(), ir.resourceDeps.end()); + ir.resourceDeps.erase( + std::unique(ir.resourceDeps.begin(), ir.resourceDeps.end()), + ir.resourceDeps.end()); + } + + const ICompiled_material *m_material; + ITransaction *m_transaction; + EmissionIR *m_ir{nullptr}; + std::unordered_map m_temporaryNodes; +}; + +} // namespace + +EmissionIR buildEmissionIR( + const ICompiled_material *compiledMaterial, ITransaction *transaction) +{ + if (!compiledMaterial || !transaction) + return {}; + return Builder(compiledMaterial, transaction).build(); +} + +} // namespace visrtx::libmdl diff --git a/devices/rtx/libmdl/EmissionIR.h b/devices/rtx/libmdl/EmissionIR.h new file mode 100644 index 000000000..067ec4893 --- /dev/null +++ b/devices/rtx/libmdl/EmissionIR.h @@ -0,0 +1,118 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// Owned, immutable intermediate representation of an MDL material's emission +// sub-DAGs (surface/backface × {emission EDF, intensity, mode}, plus +// thin_walled), extracted once from a compiled material while it is alive in +// the registry. The IR retains ZERO MDL-SDK expression pointers — the device +// does not keep the compiled material, so any retained IExpression* would +// dangle. The descriptor fold (EmissionFold) walks this IR plus a value source; +// it never touches the SDK. See ADR 0007. + +#include +#include +#include + +#include +#include +#include +#include + +namespace visrtx::libmdl { + +// Raw MDL semantic of a modeled call node. Stored as the SDK enum (a plain +// integer, not a pointer) so the fold keys on semantics, never DB names — no +// user module can masquerade as an intrinsic. DS_UNKNOWN marks an unmodeled +// call, which the fold treats as lattice-top (Unknown). +using Semantic = mi::neuraylib::IFunction_definition::Semantics; + +enum class EmissionNodeKind : std::uint8_t +{ + Constant, // a folded literal (color/float/bool/int/enum) + Parameter, // a class-compilation argument slot (symbolic under class compile) + Texture, // a tex::lookup_* whose texture is a parameter or bound resource + Call, // a modeled call (df::, operator, math::, constructor) over operands + Opaque, // an unmodeled node — the fold joins it to Unknown +}; + +enum class ConstantKind : std::uint8_t +{ + None, + Color, // rgb in value[0..2] + Float, // scalar broadcast into value[0..2] + Bool, // in boolValue + Int, // in intValue + Enum, // enum ordinal in intValue + InvalidDf, // the null EDF (default edf()) — an EDF root folds to ProvablyNull +}; + +// A single IR node. Operand references are indices into EmissionIR::nodes, so +// common subexpressions (MDL `let` temporaries) resolve to ONE shared node — +// giving the fold's `−` rule a decidable exact-identity test. +struct EmissionNode +{ + EmissionNodeKind kind{EmissionNodeKind::Opaque}; + + // Constant payload (kind == Constant). + ConstantKind constantKind{ConstantKind::None}; + std::array value{}; + bool boolValue{false}; + int intValue{0}; + + // Parameter / Texture payload: the class-compilation argument this node + // reads, or -1 for a body-literal/bound texture with no argument slot. + int parameterIndex{-1}; + std::string parameterName; + + // Texture payload: the DB name of a body-literal bound texture resource + // (empty when the texture is argument-driven — then parameterIndex is set). + // The device maps this name to a target-code texture slot; the IR stays + // device-free. + std::string resourceName; + + // Call payload (kind == Call): the MDL semantic and operand node indices. + Semantic semantic{Semantic::DS_UNKNOWN}; + std::vector operands; +}; + +// The three emission roots of one material side. Each is a node index, or -1 if +// absent (e.g. a material with no emission has edfRoot referencing an +// invalid-df constant, never -1; -1 means the sub-expression was missing). +struct EmissionSlotIR +{ + int edfRoot{-1}; // surface.emission.emission (the EDF) + int intensityRoot{-1}; // surface.emission.intensity + int modeRoot{-1}; // surface.emission.mode (an intensity_mode enum) +}; + +struct EmissionIR +{ + std::vector nodes; + EmissionSlotIR surface; + EmissionSlotIR backface; + int thinWalledRoot{-1}; + + // Structural dependency sets: every class-compilation argument slot and + // resource slot reachable from the emission roots, collected across ALL + // branches of every ternary (dead arms included) so a later condition flip + // never reads a slot the collection missed. Topology is frozen per compiled + // material, so a superset is safe. + std::vector emissionDeps; // parameter indices + std::vector resourceDeps; // bound-texture resource DB names + + bool empty() const + { + return nodes.empty(); + } +}; + +// Build the emission IR from a compiled material. `transaction` must be open +// (it is used only to resolve direct-call definitions to their semantics; no +// SDK pointer is retained past return). Never returns SDK handles. +EmissionIR buildEmissionIR( + const mi::neuraylib::ICompiled_material *compiledMaterial, + mi::neuraylib::ITransaction *transaction); + +} // namespace visrtx::libmdl diff --git a/devices/rtx/libmdl/ResourceStats.h b/devices/rtx/libmdl/ResourceStats.h new file mode 100644 index 000000000..4ea013ad9 --- /dev/null +++ b/devices/rtx/libmdl/ResourceStats.h @@ -0,0 +1,26 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +// Per-channel texel reduction of a bound texture, in sampler-output space. The +// emission classifier's value source hands these to the fold. SDK-free so both +// libmdl and the device sampler produce/consume it without the MDL SDK. + +#include + +namespace visrtx::libmdl { + +struct ResourceStats +{ + bool valid{false}; // false ⇒ unbound/invalid ⇒ a lookup folds to 0 + std::array + maxAbs{}; // maxAbs==0 (with transferPreservesZero) ⇒ zero + std::array meanPositive{}; // mean of max(texel,0) ⇒ magnitude proxy + std::array minValue{}; // minValue>=0 ⇒ sign ProvablyNonnegative + bool transferPreservesZero{ + true}; // T(0)==0; else the stored-zero bound breaks + bool finite{true}; +}; + +} // namespace visrtx::libmdl