From 5b5411ae6c7e03f4837cf0b73d9ee4a96f425912 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 15:46:22 +0000 Subject: [PATCH 01/12] docs(rtx): ADR 0007 emission descriptor/faithfulSet policy + ADR 0005 gate fix Add ADR 0007 (emission descriptor and faithfulSet registration policy). Correct ADR 0005: the Geometry Light gate is emissionIsSampleable, not emissionIsConstant (the latter is only the Stage-1 constant-radiance deposit fast path); the text predated the Stage-2 generalization. --- ...ry-lights-share-the-light-sampling-path.md | 9 +- ...emission-descriptor-registration-policy.md | 310 ++++++++++++++++++ 2 files changed, 315 insertions(+), 4 deletions(-) create mode 100644 devices/rtx/docs/adr/0007-emission-descriptor-registration-policy.md 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. From 9a00314e9006ae4ccdd00d5b9e458f063dfe1e04 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 15:54:02 +0000 Subject: [PATCH 02/12] test(rtx): device-free MDL emission classifier test target Add rtx_test_MdlEmissionClassifier, a host-only unit test linking libmdl (no CUDA/OptiX/GPU; the MDL SDK is dlopen'd at runtime). It compiles inline .mdl snippets through the same Core flow the device uses and asserts on classifyEmission, characterizing current behavior as a baseline before the descriptor refactor: constant-literal fold, parameter/texture dynamic recipes, no-emission, and power-mode rejection. SKIPs (return 77) when the MDL SDK is not loadable. --- devices/rtx/apps/tests/unit/CMakeLists.txt | 17 ++ .../tests/unit/test_MdlEmissionClassifier.cpp | 216 ++++++++++++++++++ 2 files changed, 233 insertions(+) create mode 100644 devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp 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..d505c1013 --- /dev/null +++ b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp @@ -0,0 +1,216 @@ +/* + * 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. libmdl is a +// standalone static library (no CUDA/OptiX/GPU), so classifyEmission is +// 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) rather than +// failing. +// +// These lock the CURRENT classifyEmission behavior as a baseline before the +// descriptor refactor (ADR 0007). The descriptor-fold vectors are added to this +// same target as the refactor lands. + +#include "libmdl/Core.h" +#include "libmdl/source_name_utils.h" + +#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(...)` and classify its emission. The +// compiled material must outlive the returned classification's use, so it is +// kept alive for the duration of the caller's asserts via `keepAlive`. +Core::EmissionClassification classify(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; + return {}; + } + return Core::classifyEmission(keepAlive.get()); +} + +// 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"; + +using DynamicSource = Core::EmissionClassification::DynamicSource; + +void run(Core &core, mi::neuraylib::ITransaction *txn) +{ + mi::base::Handle keepAlive; + + { + auto c = classify(core, txn, kConstLiteral, keepAlive); + CHECK(c.isDiffuseEmission); + CHECK(c.constantRadiance.has_value()); + if (c.constantRadiance) { + CHECK(approxEqual((*c.constantRadiance)[0], 2.0f)); + CHECK(approxEqual((*c.constantRadiance)[1], 2.0f)); + CHECK(approxEqual((*c.constantRadiance)[2], 2.0f)); + } + CHECK(c.dynamicSource == DynamicSource::None); + } + + { + auto c = classify(core, txn, kParamDriven, keepAlive); + CHECK(c.isDiffuseEmission); + CHECK(!c.constantRadiance.has_value()); + CHECK(c.dynamicSource == DynamicSource::Parameter); + CHECK(c.dynamicArgumentName == "value"); + CHECK(approxEqual(c.dynamicScale[0], 1.0f)); + CHECK(approxEqual(c.dynamicScale[1], 1.0f)); + CHECK(approxEqual(c.dynamicScale[2], 1.0f)); + } + + { + auto c = classify(core, txn, kTextured, keepAlive); + CHECK(c.isDiffuseEmission); + CHECK(!c.constantRadiance.has_value()); + CHECK(c.dynamicSource == DynamicSource::Texture); + CHECK(c.dynamicArgumentName == "tex"); + } + + { + auto c = classify(core, txn, kNoEmission, keepAlive); + CHECK(!c.isDiffuseEmission); + CHECK(!c.constantRadiance.has_value()); + CHECK(c.dynamicSource == DynamicSource::None); + } + + { + auto c = classify(core, txn, kPowerMode, keepAlive); + CHECK(!c.isDiffuseEmission); + CHECK(!c.constantRadiance.has_value()); + } +} + +} // namespace + +int main() +{ + using mi::base::make_handle; + + try { + Core core; + auto scope = core.createScope("MdlEmissionClassifierTestScope"); + auto txn = make_handle(core.createTransaction(scope)); + run(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; +} From a8dd05c981ce16990a2220c961865915adea6e8e Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:02:27 +0000 Subject: [PATCH 03/12] feat(libmdl): emission static pass builds an owned IR MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit buildEmissionIR walks a compiled material's emission sub-DAGs (surface/backface × {EDF, intensity, mode} plus thin_walled) into an owned, immutable IR that retains zero MDL-SDK pointers — the device does not keep the compiled material, so any retained IExpression* would dangle. Nodes are keyed on MDL semantics (IFunction_definition::get_semantic), never DB names, so no user module can masquerade as an intrinsic. Let-block temporaries are resolved and memoized so a CSE-shared subexpression maps to a single node, giving the later descriptor fold a decidable exact-identity test. Ternary arms are all walked (dead arms included) for a structural dependency superset. Device-free unit tests cover the diffuse-EDF root, argument/texture deps, the null-EDF (invalid-df) case, and shared-temporary identity. --- .../tests/unit/test_MdlEmissionClassifier.cpp | 96 ++++++ devices/rtx/libmdl/CMakeLists.txt | 1 + devices/rtx/libmdl/EmissionIR.cpp | 296 ++++++++++++++++++ devices/rtx/libmdl/EmissionIR.h | 118 +++++++ 4 files changed, 511 insertions(+) create mode 100644 devices/rtx/libmdl/EmissionIR.cpp create mode 100644 devices/rtx/libmdl/EmissionIR.h diff --git a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp index d505c1013..7bf641faa 100644 --- a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp +++ b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp @@ -16,6 +16,7 @@ // same target as the refactor lands. #include "libmdl/Core.h" +#include "libmdl/EmissionIR.h" #include "libmdl/source_name_utils.h" #include @@ -140,6 +141,100 @@ export material emissive() = material( using DynamicSource = Core::EmissionClassification::DynamicSource; +// 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) { + (void)classify(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]); + } +} + void run(Core &core, mi::neuraylib::ITransaction *txn) { mi::base::Handle keepAlive; @@ -200,6 +295,7 @@ int main() auto scope = core.createScope("MdlEmissionClassifierTestScope"); auto txn = make_handle(core.createTransaction(scope)); run(core, txn.get()); + runIR(core, txn.get()); txn->commit(); core.removeScope(scope); } catch (const std::exception &e) { diff --git a/devices/rtx/libmdl/CMakeLists.txt b/devices/rtx/libmdl/CMakeLists.txt index dcdb33831..971066fdf 100644 --- a/devices/rtx/libmdl/CMakeLists.txt +++ b/devices/rtx/libmdl/CMakeLists.txt @@ -10,6 +10,7 @@ project_add_library(STATIC ArgumentBlockDescriptor.cpp ArgumentBlockInstance.cpp Core.cpp + EmissionIR.cpp helpers.cpp ptx.cpp TimeStamp.cpp diff --git a/devices/rtx/libmdl/EmissionIR.cpp b/devices/rtx/libmdl/EmissionIR.cpp new file mode 100644 index 000000000..11d5f34e0 --- /dev/null +++ b/devices/rtx/libmdl/EmissionIR.cpp @@ -0,0 +1,296 @@ +// Copyright (c) 2019-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#include "EmissionIR.h" + +#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 : ""; + } + } + 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 From 0f3aae91e1eab32f763bd971b9d6a9ba884537cc Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:20:36 +0000 Subject: [PATCH 04/12] feat(libmdl): fold emission IR into an SDK-free descriptor MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit foldEmissionDescriptor runs the three-valued abstract interpretation of ADR 0007 over the emission IR plus an abstract value source, producing a per-slot descriptor {verdict, edfKinds, magnitude, mode, dependsOnGeometricState, sign}. The descriptor is SDK-free so the renderer consumes it without the MDL SDK. Lattice per ADR 0007: zero-absorbing product, sum zero iff all zero (nonzero only when provably nonnegative), subtraction ProvablyZero only on exact node identity, ternary folds a constant condition else joins arms (LUB, Unknown top), cross-channel color zero. EDF taxonomy keys on df:: semantics (diffuse/spot/ measured leaves, tint, directional_factor, mix families); a null (invalid-df) EDF folds to ProvablyNull. The two faithfulness flags make the registration gate sufficient, not just necessary: sign is ProvablyNonnegative only when every reachable value is proven >= 0 (so signed emission stays forward-only, unbiased), and dependsOnGeometricState is set when the emission reaches a fabricated state quantity (normal/tangent/object-id/position) — the IR now captures texture coord/wrap args as operands so a state-driven coordinate is caught, while texture_coordinate stays faithful. Magnitude is the non-negative meanPositive proxy. Device-free unit vectors cover constant/parameter/textured emission, unbound and all-black textures, negative emission (sign Unknown), power mode, and geometric- state dependence via both intensity and texture coordinate. --- .../tests/unit/test_MdlEmissionClassifier.cpp | 215 ++++++ devices/rtx/libmdl/CMakeLists.txt | 1 + devices/rtx/libmdl/EmissionDescriptor.h | 91 +++ devices/rtx/libmdl/EmissionFold.cpp | 671 ++++++++++++++++++ devices/rtx/libmdl/EmissionFold.h | 78 ++ devices/rtx/libmdl/EmissionIR.cpp | 13 + 6 files changed, 1069 insertions(+) create mode 100644 devices/rtx/libmdl/EmissionDescriptor.h create mode 100644 devices/rtx/libmdl/EmissionFold.cpp create mode 100644 devices/rtx/libmdl/EmissionFold.h diff --git a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp index 7bf641faa..997d2edbd 100644 --- a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp +++ b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp @@ -16,6 +16,8 @@ // same target as the refactor lands. #include "libmdl/Core.h" +#include "libmdl/EmissionDescriptor.h" +#include "libmdl/EmissionFold.h" #include "libmdl/EmissionIR.h" #include "libmdl/source_name_utils.h" @@ -24,6 +26,8 @@ #include #include #include +#include +#include #include using visrtx::libmdl::Core; @@ -235,6 +239,216 @@ void runIR(Core &core, mi::neuraylib::ITransaction *txn) } } +// 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; + +int paramIndexByName(const EmissionIR &ir, const char *name) +{ + for (const auto &n : ir.nodes) + if (n.parameterName == name) + return n.parameterIndex; + return -1; +} + +// Map-backed value source for the fold vectors. +struct MapValueSource : EmissionValueSource +{ + std::map> colors; + std::map resByParam; + + bool color(int i, std::array &o) const override + { + auto it = colors.find(i); + if (it == colors.end()) + return false; + o = it->second; + return true; + } + bool boolean(int, bool &) const override + { + return false; + } + bool resourceByName(const std::string &, ResourceStats &) const override + { + return false; + } + bool resourceByParam(int i, ResourceStats &o) const override + { + auto it = resByParam.find(i); + if (it == resByParam.end()) + return false; + o = it->second; + return true; + } +}; + +void runFold(Core &core, mi::neuraylib::ITransaction *txn) +{ + mi::base::Handle keepAlive; + NullValueSource none; + + auto irOf = [&](std::string_view src) { + (void)classify(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[paramIndexByName(ir, "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[paramIndexByName(ir, "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[paramIndexByName(ir, "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[paramIndexByName(ir, "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[paramIndexByName(ir, "tex")] = s; + auto d = foldEmissionDescriptor(ir, vs).surface; + CHECK(d.dependsOnGeometricState); + } +} + void run(Core &core, mi::neuraylib::ITransaction *txn) { mi::base::Handle keepAlive; @@ -296,6 +510,7 @@ int main() auto txn = make_handle(core.createTransaction(scope)); run(core, txn.get()); runIR(core, txn.get()); + runFold(core, txn.get()); txn->commit(); core.removeScope(scope); } catch (const std::exception &e) { diff --git a/devices/rtx/libmdl/CMakeLists.txt b/devices/rtx/libmdl/CMakeLists.txt index 971066fdf..539d450fb 100644 --- a/devices/rtx/libmdl/CMakeLists.txt +++ b/devices/rtx/libmdl/CMakeLists.txt @@ -10,6 +10,7 @@ project_add_library(STATIC ArgumentBlockDescriptor.cpp ArgumentBlockInstance.cpp Core.cpp + EmissionFold.cpp EmissionIR.cpp helpers.cpp ptx.cpp 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..b19570fc1 --- /dev/null +++ b/devices/rtx/libmdl/EmissionFold.cpp @@ -0,0 +1,671 @@ +// 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. + desc.sign = intensity.sign; + 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 + 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)) { + // Union component kinds; null iff every EDF operand is null. A precise + // per-weight analysis needs the df_component array shape; the + // conservative union keeps kinds honest (drives the fidelity gate) + // without proving Emissive, which only affects variance. + bool allNull = true; + for (int op : n.operands) { + if (!isEdfNode(op)) + continue; + Edf c = evalEdf(op, depth + 1); + r.kinds |= c.kinds; + r.dependsOnState = r.dependsOnState || c.dependsOnState; + if (c.null != Tri::True) + allNull = false; + } + if (r.kinds == EdfKind::None) + r.kinds = EdfKind::Unknown; + r.null = allNull ? Tri::True : Tri::Unknown; + return r; + } + + // Unmodeled df:: call. + r.kinds = EdfKind::Unknown; + return r; + } + + 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.parameterIndex, 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; + } + + ResourceStats stats; + const bool known = n.parameterIndex >= 0 + ? m_values.resourceByParam(n.parameterIndex, 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..89dfca2a7 --- /dev/null +++ b/devices/rtx/libmdl/EmissionFold.h @@ -0,0 +1,78 @@ +// 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 +#include + +namespace visrtx::libmdl { + +// Per-channel texel reduction of a bound texture, in sampler-output space. +struct ResourceStats +{ + bool valid{false}; // false ⇒ unbound/invalid ⇒ lookup folds to 0 + std::array maxAbs{}; // maxAbs==0 ⇒ ProvablyZero (exact) + 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 zero bound breaks + bool finite{true}; +}; + +// Supplies current parameter values and resource stats to the fold. The device +// backs it with live argument bytes + the resource table; tests back it with +// fakes or the compiled material's own arguments. All indices are IR parameter +// indices; resource names are IR resource DB names. +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(int parameterIndex, std::array &out) const = 0; + virtual bool boolean(int parameterIndex, bool &out) const = 0; + + // Stats for a body-literal bound texture (by DB name) or an argument-bound + // texture (by parameter index; name empty). Returning false ⇒ Unknown. + virtual bool resourceByName( + const std::string &name, ResourceStats &out) const = 0; + virtual bool resourceByParam( + int parameterIndex, 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(int, std::array &) const override + { + return false; + } + bool boolean(int, bool &) const override + { + return false; + } + bool resourceByName(const std::string &, ResourceStats &) const override + { + return false; + } + bool resourceByParam(int, 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 index 11d5f34e0..a670b5623 100644 --- a/devices/rtx/libmdl/EmissionIR.cpp +++ b/devices/rtx/libmdl/EmissionIR.cpp @@ -8,6 +8,7 @@ #include #include +#include #include namespace visrtx::libmdl { @@ -250,6 +251,18 @@ class Builder 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)); } From 2277a3965f5214462f565224eceda664cc2575c7 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:28:58 +0000 Subject: [PATCH 05/12] feat(rtx): unify emissive Pick Power on the non-negative magnitude MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Image2D folds the emissive Pick Power and the emission-classifier stats into ONE memoized host scan (computeTextureReduction, guarded on the image data stamp): per channel, maxAbs (the exact zero proof), meanPositive = mean of max(texel,0) (the non-negative magnitude proxy), and minValue (the non-negative sign proof). This replaces the two separate scans (averageValue's mean and a standalone emissionStats scan) and their two stamps with a single reduction; averageValue() and emissionStats() both read it. The scan keeps the sRGB->linear on color channels and the 1/2-channel broadcast. transferPreservesZero is false only under a nonzero border color, the one case a stored-zero texel can sample nonzero. averageValue() now returns meanPositive — the SAME non-negative magnitude proxy the MDL classifier uses — instead of the plain signed mean, so all three emissive material types (PBR, PhysicallyBasedMDL, MDL) size Pick Power from one projection. Native emission is radiance >= 0, so meanPositive == mean for every real emitter (the parity tests are unchanged); the two only diverge on a signed-texel emissive texture, where clamping the negative part out of the picked power is the correct, CDF-representable magnitude. The dead signed-mean accumulator is dropped from the reduction. The emission classifier is the sole consumer of ResourceStats and lives behind USE_MDL, so Sampler::emissionStats(), Image2D::emissionStats(), and the libmdl/ResourceStats.h include are all gated on USE_MDL — the base sampler no longer reaches into libmdl in an MDL-off build. The Sampler base default proves nothing (not zero, not non-negative), so a sampler that cannot reduce its texels keeps its emission forward-only via the policy's sign gate rather than registering an unproven light. The classifier stats are not yet consumed — the descriptor value source wires them in the next commit. Emissive/MDL API tests (incl. the three-way parity test) and the classifier unit test stay green. --- devices/rtx/device/sampler/Image2D.cpp | 142 +++++++++++++++++-------- devices/rtx/device/sampler/Image2D.h | 43 +++++--- devices/rtx/device/sampler/Sampler.cpp | 18 ++++ devices/rtx/device/sampler/Sampler.h | 14 +++ devices/rtx/libmdl/EmissionFold.h | 12 +-- devices/rtx/libmdl/ResourceStats.h | 26 +++++ 6 files changed, 185 insertions(+), 70 deletions(-) create mode 100644 devices/rtx/libmdl/ResourceStats.h diff --git a/devices/rtx/device/sampler/Image2D.cpp b/devices/rtx/device/sampler/Image2D.cpp index fdc14b5b1..f3dba8919 100644 --- a/devices/rtx/device/sampler/Image2D.cpp +++ b/devices/rtx/device/sampler/Image2D.cpp @@ -33,6 +33,10 @@ #include "utility/AnariTypeHelpers.h" +#include +#include +#include + namespace visrtx { static float srgbToLinear(float v) @@ -90,9 +94,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 +108,125 @@ 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 host scan 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). Reads the retained host pixels; sRGB byte data is +// linearized to match the hardware sampler. Unsupported element types leave the +// reduction Unknown (magnitude stays unit so the emitter is still picked). +// +// TODO(perf): reduce over the resident device texels instead of this host scan +// — the image already uploads a linear device buffer (data(AddressSpace:: +// DEVICE)), so a thrust reduction (cf. light/sampling/CDF.cu) avoids the host +// readback for large emissive textures. The device functor must reproduce this +// per-channel sRGB->linear (color channels only) decode before reducing. +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(); + 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. + // 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); + // Per-source-channel accumulators (up to 4). + std::array posSum{}; + std::array maxAbs{}; + std::array minVal{{FLT_MAX, FLT_MAX, FLT_MAX, FLT_MAX}}; + bool finite = true; + + auto accumulate = [&](int c, float v) { + if (!std::isfinite(v)) + finite = false; + posSum[c] += double(std::max(v, 0.0f)); + maxAbs[c] = std::max(maxAbs[c], std::fabs(v)); + minVal[c] = std::min(minVal[c], v); + }; + 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]); + accumulate(c, 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); + accumulate(c, (srgb && c < colorChannels) ? srgbToLinear(v) : v); } } else { - return Sampler::averageValue(); // uncommon type for emission; coarse fallback + return r; // uncommon type for emission ⇒ Unknown, magnitude stays unit } - 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); -} + // 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(posSum[c] / double(count)); + r.maxAbs[rgb] = maxAbs[c]; + r.minValue[rgb] = minVal[c]; + } -// 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 = 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/libmdl/EmissionFold.h b/devices/rtx/libmdl/EmissionFold.h index 89dfca2a7..270a13f7f 100644 --- a/devices/rtx/libmdl/EmissionFold.h +++ b/devices/rtx/libmdl/EmissionFold.h @@ -10,23 +10,13 @@ #include "EmissionDescriptor.h" #include "EmissionIR.h" +#include "ResourceStats.h" #include #include namespace visrtx::libmdl { -// Per-channel texel reduction of a bound texture, in sampler-output space. -struct ResourceStats -{ - bool valid{false}; // false ⇒ unbound/invalid ⇒ lookup folds to 0 - std::array maxAbs{}; // maxAbs==0 ⇒ ProvablyZero (exact) - 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 zero bound breaks - bool finite{true}; -}; - // Supplies current parameter values and resource stats to the fold. The device // backs it with live argument bytes + the resource table; tests back it with // fakes or the compiled material's own arguments. All indices are IR parameter 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 From 16f6b38662998a7e4d052487c72e403560d77a16 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:51:33 +0000 Subject: [PATCH 06/12] feat(rtx): MDL materials publish an emission descriptor (ADR 0007) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replace the raw-mdl compile-time EmissionClassification with the owned emission IR + descriptor fold. At finalize, MDL folds the registry's cached IR against this instance's live arguments and bound-sampler reductions into an EmissionDescriptor, then refreshes the light set. emissionIsSampleable applies the registration policy (non-null, diffuse, radiant-exitance, non-negative sign, no geometric-state dependence); emissionAverage returns the non-negative meanPositive magnitude. Deletes classifyEmission, walkIntensityFactors, the dynamic-recipe evaluation, and the EmissionClassification struct — the IR fold subsumes them and handles mixes, tints, directional factors, and geometric-state detection the string- prefix walk could not. Also hardens the fold against two defects: a df::*_mix's EDF leaves nest inside its df_component[] array, so the fold now deep-scans for them instead of proving a false ProvablyNull; and a non-finite intensity now fails the sign gate rather than registering an infinite Pick Power. (The texture-coordinate geometric-state scan is already present.) Behavior change: an UNBOUND emissive texture is no longer a Geometry Light — its MDL lookup is invalid and folds to 0, so it emits nothing (ADR 0007 A5), superseding ADR-0006's zero-radiance over-inclusion. Bound textured emitters still register with correct Pick Power (TestEmissiveMdlLight radiance parity). PhysicallyBasedMDL is unchanged (its post-translate capture already yields the right sampleability/average). All 60 tests pass. --- .../apps/tests/api/TestEmissiveMdlLight.cpp | 10 +- .../tests/unit/test_MdlEmissionClassifier.cpp | 151 ++++----- devices/rtx/device/material/MDL.cpp | 199 +++++++----- devices/rtx/device/material/MDL.h | 34 +- devices/rtx/device/mdl/MaterialRegistry.cpp | 30 +- devices/rtx/device/mdl/MaterialRegistry.h | 13 +- devices/rtx/libmdl/Core.cpp | 300 +----------------- devices/rtx/libmdl/Core.h | 43 +-- devices/rtx/libmdl/EmissionFold.cpp | 94 ++++-- devices/rtx/libmdl/EmissionFold.h | 23 +- 10 files changed, 320 insertions(+), 577 deletions(-) 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/test_MdlEmissionClassifier.cpp b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp index 997d2edbd..b6a64b1cc 100644 --- a/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp +++ b/devices/rtx/apps/tests/unit/test_MdlEmissionClassifier.cpp @@ -3,17 +3,18 @@ * SPDX-License-Identifier: BSD-3-Clause */ -// Device-free unit tests for the MDL emission classifier. libmdl is a -// standalone static library (no CUDA/OptiX/GPU), so classifyEmission is -// 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) rather than -// failing. +// 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). // -// These lock the CURRENT classifyEmission behavior as a baseline before the -// descriptor refactor (ADR 0007). The descriptor-fold vectors are added to this -// same target as the refactor lands. +// 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" @@ -26,6 +27,7 @@ #include #include #include +#include #include #include #include @@ -52,10 +54,9 @@ 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(...)` and classify its emission. The -// compiled material must outlive the returned classification's use, so it is -// kept alive for the duration of the caller's asserts via `keepAlive`. -Core::EmissionClassification classify(Core &core, +// 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) @@ -66,22 +67,20 @@ Core::EmissionClassification classify(Core &core, if (!module.is_valid_interface()) { std::printf("FAIL could not load inline module\n"); ++g_failures; - return {}; + 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 {}; + return; } keepAlive = make_handle(core.getCompiledMaterial(fnDef.get())); if (!keepAlive.is_valid_interface()) { std::printf("FAIL could not compile material\n"); ++g_failures; - return {}; } - return Core::classifyEmission(keepAlive.get()); } // intensity = color(2.0) * math::PI folds to a body-literal constant: emitted @@ -143,8 +142,6 @@ export material emissive() = material( mode: intensity_power))); )mdl"; -using DynamicSource = Core::EmissionClassification::DynamicSource; - // 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. @@ -185,7 +182,7 @@ void runIR(Core &core, mi::neuraylib::ITransaction *txn) mi::base::Handle keepAlive; auto buildIR = [&](std::string_view src) { - (void)classify(core, txn, src, keepAlive); // reuse the compile flow + compileMaterial(core, txn, src, keepAlive); // reuse the compile flow return buildEmissionIR(keepAlive.get(), txn); }; @@ -289,29 +286,21 @@ using visrtx::libmdl::IntensityMode; using visrtx::libmdl::NullValueSource; using visrtx::libmdl::ResourceStats; -int paramIndexByName(const EmissionIR &ir, const char *name) -{ - for (const auto &n : ir.nodes) - if (n.parameterName == name) - return n.parameterIndex; - return -1; -} - -// Map-backed value source for the fold vectors. +// Map-backed value source for the fold vectors, keyed by parameter name. struct MapValueSource : EmissionValueSource { - std::map> colors; - std::map resByParam; + std::map> colors; + std::map resByParam; - bool color(int i, std::array &o) const override + bool color(const std::string &name, std::array &o) const override { - auto it = colors.find(i); + auto it = colors.find(name); if (it == colors.end()) return false; o = it->second; return true; } - bool boolean(int, bool &) const override + bool boolean(const std::string &, bool &) const override { return false; } @@ -319,9 +308,9 @@ struct MapValueSource : EmissionValueSource { return false; } - bool resourceByParam(int i, ResourceStats &o) const override + bool resourceByParam(const std::string &name, ResourceStats &o) const override { - auto it = resByParam.find(i); + auto it = resByParam.find(name); if (it == resByParam.end()) return false; o = it->second; @@ -329,13 +318,29 @@ struct MapValueSource : EmissionValueSource } }; +// 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) { - (void)classify(core, txn, src, keepAlive); + compileMaterial(core, txn, src, keepAlive); return buildEmissionIR(keepAlive.get(), txn); }; @@ -353,7 +358,7 @@ void runFold(Core &core, mi::neuraylib::ITransaction *txn) { // parameter with a known live value ⇒ magnitude tracks it auto ir = irOf(kParamDriven); MapValueSource vs; - vs.colors[paramIndexByName(ir, "value")] = {3.0f, 3.0f, 3.0f}; + 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); @@ -389,7 +394,7 @@ void runFold(Core &core, mi::neuraylib::ITransaction *txn) s.meanPositive = {4.0f, 4.0f, 4.0f}; s.minValue = {0.0f, 0.0f, 0.0f}; s.transferPreservesZero = true; - vs.resByParam[paramIndexByName(ir, "tex")] = s; + vs.resByParam[std::string("tex")] = s; auto d = foldEmissionDescriptor(ir, vs).surface; CHECK(d.verdict == EmissionVerdict::Unknown); CHECK(d.sign == EmissionSign::ProvablyNonnegative); @@ -406,7 +411,7 @@ void runFold(Core &core, mi::neuraylib::ITransaction *txn) s.meanPositive = {0.0f, 0.0f, 0.0f}; s.minValue = {0.0f, 0.0f, 0.0f}; s.transferPreservesZero = true; - vs.resByParam[paramIndexByName(ir, "tex")] = s; + vs.resByParam[std::string("tex")] = s; auto d = foldEmissionDescriptor(ir, vs).surface; CHECK(d.verdict == EmissionVerdict::ProvablyNull); } @@ -416,7 +421,7 @@ void runFold(Core &core, mi::neuraylib::ITransaction *txn) MapValueSource vs; ResourceStats s; s.valid = false; - vs.resByParam[paramIndexByName(ir, "tex")] = s; + vs.resByParam[std::string("tex")] = s; auto d = foldEmissionDescriptor(ir, vs).surface; CHECK(d.verdict == EmissionVerdict::ProvablyNull); } @@ -443,58 +448,27 @@ void runFold(Core &core, mi::neuraylib::ITransaction *txn) 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[paramIndexByName(ir, "tex")] = s; + vs.resByParam[std::string("tex")] = s; auto d = foldEmissionDescriptor(ir, vs).surface; CHECK(d.dependsOnGeometricState); } -} - -void run(Core &core, mi::neuraylib::ITransaction *txn) -{ - mi::base::Handle keepAlive; - - { - auto c = classify(core, txn, kConstLiteral, keepAlive); - CHECK(c.isDiffuseEmission); - CHECK(c.constantRadiance.has_value()); - if (c.constantRadiance) { - CHECK(approxEqual((*c.constantRadiance)[0], 2.0f)); - CHECK(approxEqual((*c.constantRadiance)[1], 2.0f)); - CHECK(approxEqual((*c.constantRadiance)[2], 2.0f)); - } - CHECK(c.dynamicSource == DynamicSource::None); - } - - { - auto c = classify(core, txn, kParamDriven, keepAlive); - CHECK(c.isDiffuseEmission); - CHECK(!c.constantRadiance.has_value()); - CHECK(c.dynamicSource == DynamicSource::Parameter); - CHECK(c.dynamicArgumentName == "value"); - CHECK(approxEqual(c.dynamicScale[0], 1.0f)); - CHECK(approxEqual(c.dynamicScale[1], 1.0f)); - CHECK(approxEqual(c.dynamicScale[2], 1.0f)); - } - - { - auto c = classify(core, txn, kTextured, keepAlive); - CHECK(c.isDiffuseEmission); - CHECK(!c.constantRadiance.has_value()); - CHECK(c.dynamicSource == DynamicSource::Texture); - CHECK(c.dynamicArgumentName == "tex"); - } - { - auto c = classify(core, txn, kNoEmission, keepAlive); - CHECK(!c.isDiffuseEmission); - CHECK(!c.constantRadiance.has_value()); - CHECK(c.dynamicSource == DynamicSource::None); + { // 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)); } - { - auto c = classify(core, txn, kPowerMode, keepAlive); - CHECK(!c.isDiffuseEmission); - CHECK(!c.constantRadiance.has_value()); + { // 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); } } @@ -508,7 +482,6 @@ int main() Core core; auto scope = core.createScope("MdlEmissionClassifierTestScope"); auto txn = make_handle(core.createTransaction(scope)); - run(core, txn.get()); runIR(core, txn.get()); runFold(core, txn.get()); txn->commit(); diff --git a/devices/rtx/device/material/MDL.cpp b/devices/rtx/device/material/MDL.cpp index 1d8cf7895..f8c5e825f 100644 --- a/devices/rtx/device/material/MDL.cpp +++ b/devices/rtx/device/material/MDL.cpp @@ -35,6 +35,8 @@ #include "gpu/sbt.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 +67,69 @@ 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; + } +}; + +// The renderer's registration policy (ADR 0007). C7 extracts this to a shared +// EmissionPolicy header cross-referenced from the GPU Pick-Power/normal sites. +bool isRegisterable(const libmdl::SlotDescriptor &s) +{ + return s.verdict != libmdl::EmissionVerdict::ProvablyNull + && libmdl::isSubsetOf(s.edfKinds, libmdl::EdfKind::Diffuse) + && s.mode == libmdl::IntensityMode::RadiantExitance + && !s.dependsOnGeometricState + && s.sign == libmdl::EmissionSign::ProvablyNonnegative; +} + +} // namespace + MDL::MDL(DeviceGlobalState *d) : Material(d) {} MDL::~MDL() @@ -111,7 +176,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 +187,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 +212,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 +372,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 +392,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 +401,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 +533,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 +574,8 @@ void MDL::syncParameters() void MDL::syncImplementationIndex() { - m_implementationIndex = deviceState()->mdl->materialRegistry.getMaterialImplementationIndex( + m_implementationIndex = + deviceState()->mdl->materialRegistry.getMaterialImplementationIndex( m_uuid); } @@ -560,13 +587,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/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/EmissionFold.cpp b/devices/rtx/libmdl/EmissionFold.cpp index b19570fc1..3df633ba1 100644 --- a/devices/rtx/libmdl/EmissionFold.cpp +++ b/devices/rtx/libmdl/EmissionFold.cpp @@ -178,8 +178,10 @@ class Fold desc.magnitude = {1.0f, 1.0f, 1.0f}; } - // sign gates registration; the EDF itself is nonnegative. - desc.sign = intensity.sign; + // 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; } @@ -226,6 +228,15 @@ class Fold 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; } @@ -288,23 +299,17 @@ class Fold } if (isMixSemantic(n.semantic)) { - // Union component kinds; null iff every EDF operand is null. A precise - // per-weight analysis needs the df_component array shape; the - // conservative union keeps kinds honest (drives the fidelity gate) - // without proving Emissive, which only affects variance. - bool allNull = true; - for (int op : n.operands) { - if (!isEdfNode(op)) - continue; - Edf c = evalEdf(op, depth + 1); - r.kinds |= c.kinds; - r.dependsOnState = r.dependsOnState || c.dependsOnState; - if (c.null != Tri::True) - allNull = false; - } + // 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 = allNull ? Tri::True : Tri::Unknown; + r.null = Tri::Unknown; return r; } @@ -313,6 +318,56 @@ class Fold 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) @@ -372,7 +427,7 @@ class Fold { Scalar r; std::array value; - if (m_values.color(n.parameterIndex, value)) { + if (m_values.color(n.parameterName, value)) { r.finite = allFinite(value); r.zero = allZero(value) ? Tri::True : anyNonzero(value) ? Tri::False @@ -394,11 +449,12 @@ class Fold 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.parameterIndex, stats) + ? m_values.resourceByParam(n.parameterName, stats) : m_values.resourceByName(n.resourceName, stats); if (!known) return r; // Unknown (state-dependence preserved) diff --git a/devices/rtx/libmdl/EmissionFold.h b/devices/rtx/libmdl/EmissionFold.h index 270a13f7f..198d3d2ec 100644 --- a/devices/rtx/libmdl/EmissionFold.h +++ b/devices/rtx/libmdl/EmissionFold.h @@ -17,10 +17,10 @@ namespace visrtx::libmdl { -// Supplies current parameter values and resource stats to the fold. The device -// backs it with live argument bytes + the resource table; tests back it with -// fakes or the compiled material's own arguments. All indices are IR parameter -// indices; resource names are IR resource DB names. +// 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: @@ -28,15 +28,16 @@ class EmissionValueSource // 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(int parameterIndex, std::array &out) const = 0; - virtual bool boolean(int parameterIndex, bool &out) const = 0; + 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 index; name empty). Returning false ⇒ Unknown. + // texture (by parameter name). Returning false ⇒ Unknown. virtual bool resourceByName( const std::string &name, ResourceStats &out) const = 0; virtual bool resourceByParam( - int parameterIndex, ResourceStats &out) const = 0; + const std::string ¶meterName, ResourceStats &out) const = 0; }; // A value source that knows nothing — every parameter/resource is Unknown. Used @@ -44,11 +45,11 @@ class EmissionValueSource class NullValueSource : public EmissionValueSource { public: - bool color(int, std::array &) const override + bool color(const std::string &, std::array &) const override { return false; } - bool boolean(int, bool &) const override + bool boolean(const std::string &, bool &) const override { return false; } @@ -56,7 +57,7 @@ class NullValueSource : public EmissionValueSource { return false; } - bool resourceByParam(int, ResourceStats &) const override + bool resourceByParam(const std::string &, ResourceStats &) const override { return false; } From a1e27a3e3042b36299bce21f53790ea162942c9e Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:54:50 +0000 Subject: [PATCH 07/12] feat(rtx): extract the faithfulSet registration policy MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Move the inline registration gate to material/EmissionPolicy.h as the single source of truth: kFaithfulSet = {Diffuse} and isRegisterable(SlotDescriptor) (non-null, kinds ⊆ faithfulSet, radiant-exitance, non-negative sign, no geometric-state dependence). MDL consumes it instead of a private helper. Cross-reference the two GPU-side sites that encode the same diffuse-only assumption — lightPickPower.h (double-sided Lambertian flux) and sampleLight.h (double-sided normal, geometric normal as shading normal) — so growing the set has one discoverable set of lockstep edits. No behavior change; all tests pass. --- devices/rtx/device/gpu/lightPickPower.h | 16 ++- devices/rtx/device/gpu/sampleLight.h | 141 ++++++++++++------- devices/rtx/device/material/EmissionPolicy.h | 38 +++++ devices/rtx/device/material/MDL.cpp | 13 +- 4 files changed, 138 insertions(+), 70 deletions(-) create mode 100644 devices/rtx/device/material/EmissionPolicy.h 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..bcfbbc9d3 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 @@ -328,8 +327,9 @@ VISRTX_DEVICE LightSample sampleSpotLight( 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 distribution over a cumulative CDF. + // 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); } @@ -342,10 +342,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 +363,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 +393,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 +413,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 +456,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 +498,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 +525,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 +557,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 +585,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 +606,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 +622,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 +640,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 +650,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 +679,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 +752,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 +798,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 +822,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 +889,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 f8c5e825f..d4fb5e3ee 100644 --- a/devices/rtx/device/material/MDL.cpp +++ b/devices/rtx/device/material/MDL.cpp @@ -33,6 +33,8 @@ #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" @@ -117,17 +119,6 @@ struct MDLValueSource : libmdl::EmissionValueSource } }; -// The renderer's registration policy (ADR 0007). C7 extracts this to a shared -// EmissionPolicy header cross-referenced from the GPU Pick-Power/normal sites. -bool isRegisterable(const libmdl::SlotDescriptor &s) -{ - return s.verdict != libmdl::EmissionVerdict::ProvablyNull - && libmdl::isSubsetOf(s.edfKinds, libmdl::EdfKind::Diffuse) - && s.mode == libmdl::IntensityMode::RadiantExitance - && !s.dependsOnGeometricState - && s.sign == libmdl::EmissionSign::ProvablyNonnegative; -} - } // namespace MDL::MDL(DeviceGlobalState *d) : Material(d) {} From c0ba669bf6b5141b2a0c8d2d3f83fa047e121044 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 16:59:34 +0000 Subject: [PATCH 08/12] test(rtx): integration test for the emission registration policy MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit TestEmissionDescriptorPolicy asserts each faithfulSet gate end-to-end through the ANARI API → compile → fold → policy → light synthesis path, via the world's numLightInstances property: a constant diffuse emitter registers (1), while a non-diffuse spot EDF, provably-negative emission, state::normal-dependent intensity, and power-mode intensity are each described but NOT registered (0) — they render via the forward path, unbiased. Registered here, not radiance: the count is the observable seam for exclusion (a forward-only emitter is pixel-identical to no light). --- devices/rtx/apps/tests/api/CMakeLists.txt | 1 + .../api/TestEmissionDescriptorPolicy.cpp | 177 ++++++++++++++++++ 2 files changed, 178 insertions(+) create mode 100644 devices/rtx/apps/tests/api/TestEmissionDescriptorPolicy.cpp 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; +} From 698a82eebfcf43aad8ecb3a46e90d12e97e429ac Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 17:18:57 +0000 Subject: [PATCH 09/12] fix(rtx): preserve dim-light mass in the geometry-light pick CDF MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The Light Pick CDF was a normalized cumulative float, so a light whose power is below float epsilon of the total quantized to a zero-width interval: the NEE-sample side (inverseSampleCDF, instancePickProbability delta) could never select it, yet the hit-side geometryLightHitPdf still reported a positive analytic pick probability. The deposit's MIS weight then down-weighted the forward path for a light next-event never delivered — the dim emitter rendered too dark in a bright+dim scene. Store and accumulate the CDF in double (masses preserved to ~1e-16 of total) and difference the pick-probability delta in double, so any physically meaningful light stays selectable with a delta that matches the hit-side pNee. The float HDRI/primitive CDFs are untouched: inverseSampleCDF is now templated on the CDF element type. Normalize by the double cumulative total, not the float m_totalLightPower: dividing by the float total can push the last entry above 1.0 when float dropped a dim light's mass, leaving trailing dim lights unselectable while their pNee stays positive. The double total makes the last entry exactly 1.0 and preserves those masses. No test regression (the effect is below the float-epsilon noise floor to reproduce automatically); this is precision hardening. --- devices/rtx/device/gpu/gpu_objects.h | 13 +-- devices/rtx/device/gpu/sampleLight.h | 12 +-- devices/rtx/device/renderer/Quality_ptx.cu | 92 ++++++++++++---------- devices/rtx/device/world/World.cpp | 19 +++-- devices/rtx/device/world/World.h | 2 +- 5 files changed, 82 insertions(+), 56 deletions(-) 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/sampleLight.h b/devices/rtx/device/gpu/sampleLight.h index bcfbbc9d3..2b767199b 100644 --- a/devices/rtx/device/gpu/sampleLight.h +++ b/devices/rtx/device/gpu/sampleLight.h @@ -325,12 +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. diff --git a/devices/rtx/device/renderer/Quality_ptx.cu b/devices/rtx/device/renderer/Quality_ptx.cu index caf91046b..d9d512494 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 @@ -213,14 +219,14 @@ VISRTX_DEVICE float geometryLightHitPdf(const FrameGPUData &frameData, { // 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 +234,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 +254,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; @@ -328,8 +336,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 +571,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,10 +665,11 @@ 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( @@ -668,7 +677,8 @@ VISRTX_GLOBAL void __raygen__() 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 +702,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 +719,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 +783,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/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}; From 989a7359e89b92d7a5a029ab355aa719e6adb099 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 17:22:22 +0000 Subject: [PATCH 10/12] fix(rtx): hit-side geometry-light pNee reads the CDF's magnitude MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit geometryLightHitPdf built its pick probability from the per-hit emission for a constant emitter and from emissionAverage otherwise. But the Light Pick CDF is always built from emissionAverage (the geometry light's configured radiance), so the constant branch only happened to agree — and fed signed, per-point emission that would disagree once emission is negative or spatially varying, biasing the deposit's MIS weight. Always use emissionAverage (the non-negative meanPositive magnitude the CDF was built from) so the hit-side pNee equals the selection pick probability exactly. No change for constant emitters (emission == emissionAverage); closes the signed/textured edge. Drops the now-unused emission argument. Also apply the same (raw > 0 && finite) clamp World::appendLight uses when it builds the CDF: 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 — otherwise wEmission goes NaN/0 and disagrees with selection. --- devices/rtx/device/renderer/Quality_ptx.cu | 31 +++++++++++++--------- 1 file changed, 18 insertions(+), 13 deletions(-) diff --git a/devices/rtx/device/renderer/Quality_ptx.cu b/devices/rtx/device/renderer/Quality_ptx.cu index d9d512494..75509671a 100644 --- a/devices/rtx/device/renderer/Quality_ptx.cu +++ b/devices/rtx/device/renderer/Quality_ptx.cu @@ -212,10 +212,8 @@ 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`/ @@ -299,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; } @@ -672,8 +677,8 @@ VISRTX_GLOBAL void __raygen__() // == 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); } From 1229b9813fb55f2bf3cac92497f1428f602da998 Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 17:26:40 +0000 Subject: [PATCH 11/12] fix(rtx): Interactive deposits unregistered-emitter emission on the bounce MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Interactive reflection bounce shaded the hit surface with ambient only, never its emission, so an emitter reached only by a reflection was invisible there and lit nothing via the forward path — the ADR 0007 render-mode-matrix gap for Interactive. Deposit the bounce surface emission, guarded to UNREGISTERED emitters: a registered (sampleable) one is already covered by the NEE loop, so depositing it here too would double-count, while an unregistered one has no NEE and needs this forward estimate. The two sets are disjoint, so no MIS and no double-count — the parity/Interactive tests stay green. The other two matrix rows are by-design, not bugs: the matte material is diffuse-only and non-emissive per the ANARI spec (its dead continuation ray means no indirect), and Quality finite-depth truncation is inherent (the pick-power accuracy it depended on is fixed by the CDF/pNee commits). So #3 reduces to this Interactive deposit. --- devices/rtx/device/renderer/Interactive_ptx.cu | 9 +++++++++ 1 file changed, 9 insertions(+) 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; From 0ab112d63c541e38176c159f8d91d46c8010d77c Mon Sep 17 00:00:00 2001 From: Thomas Arcila <134677+tarcila@users.noreply.github.com> Date: Wed, 15 Jul 2026 19:43:30 +0000 Subject: [PATCH 12/12] perf(rtx): GPU-drive the Image2D texel reduction MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit computeTextureReduction() now reduces over the resident device texels (Array::data(DEVICE), a real H2D upload) with a single thrust::transform_reduce instead of the host scan, removing the host readback for large emissive textures. The device functor reproduces the per-texel decode exactly — uint8/255 and sRGB->linear on color channels only — so the classifier stats (maxAbs / meanPositive / minValue), and the meanPositive magnitude the emissive Pick Power reads, are bit-for-bit the same policy as the host scan it replaces. If the image has no device residency the reduction stays Unknown (magnitude unit, emission forward-only), never a crash. TextureStats.{h,cu} hold the SDK-free reduction; ungated because averageValue() feeds the non-MDL PBR emissive path too. --- devices/rtx/device/CMakeLists.txt | 1 + devices/rtx/device/sampler/Image2D.cpp | 87 ++++++--------- devices/rtx/device/sampler/TextureStats.cu | 118 +++++++++++++++++++++ devices/rtx/device/sampler/TextureStats.h | 45 ++++++++ 4 files changed, 195 insertions(+), 56 deletions(-) create mode 100644 devices/rtx/device/sampler/TextureStats.cu create mode 100644 devices/rtx/device/sampler/TextureStats.h 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/sampler/Image2D.cpp b/devices/rtx/device/sampler/Image2D.cpp index f3dba8919..0bc6e4cf8 100644 --- a/devices/rtx/device/sampler/Image2D.cpp +++ b/devices/rtx/device/sampler/Image2D.cpp @@ -31,17 +31,20 @@ #include "Image2D.h" +#include "TextureStats.h" #include "utility/AnariTypeHelpers.h" -#include -#include -#include - 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) {} @@ -148,18 +151,14 @@ const Image2D::TextureReduction &Image2D::textureReduction() const return m_reduction; } -// One host scan 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). Reads the retained host pixels; sRGB byte data is -// linearized to match the hardware sampler. Unsupported element types leave the -// reduction Unknown (magnitude stays unit so the emitter is still picked). -// -// TODO(perf): reduce over the resident device texels instead of this host scan -// — the image already uploads a linear device buffer (data(AddressSpace:: -// DEVICE)), so a thrust reduction (cf. light/sampling/CDF.cu) avoids the host -// readback for large emissive textures. The device functor must reproduce this -// per-channel sRGB->linear (color channels only) decode before reducing. +// 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; @@ -169,45 +168,21 @@ Image2D::TextureReduction Image2D::computeTextureReduction() const 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) + 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 bool srgb = isSrgb8(t); - const int colorChannels = (srgb && (nc == 2 || nc == 4)) ? nc - 1 : nc; - - // Per-source-channel accumulators (up to 4). - std::array posSum{}; - std::array maxAbs{}; - std::array minVal{{FLT_MAX, FLT_MAX, FLT_MAX, FLT_MAX}}; - bool finite = true; - - auto accumulate = [&](int c, float v) { - if (!std::isfinite(v)) - finite = false; - posSum[c] += double(std::max(v, 0.0f)); - maxAbs[c] = std::max(maxAbs[c], std::fabs(v)); - minVal[c] = std::min(minVal[c], v); - }; - - if (isFloat32(t)) { - const auto *p = static_cast(host); - for (size_t i = 0; i < count; ++i) - for (int c = 0; c < nc; ++c) - accumulate(c, 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; - accumulate(c, (srgb && c < colorChannels) ? srgbToLinear(v) : v); - } - } else { - return r; // uncommon type for emission ⇒ Unknown, magnitude stays unit - } + 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 @@ -215,13 +190,13 @@ Image2D::TextureReduction Image2D::computeTextureReduction() const 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(posSum[c] / double(count)); - r.maxAbs[rgb] = maxAbs[c]; - r.minValue[rgb] = minVal[c]; + r.meanPositive[rgb] = float(a.posSum[c] / double(count)); + r.maxAbs[rgb] = a.maxAbs[c]; + r.minValue[rgb] = a.minValue[c]; } r.valid = true; - r.finite = finite; + 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 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