diff --git a/devices/rtx/device/frame/Frame.cu b/devices/rtx/device/frame/Frame.cu
index 8ffcbe728..3c04f09ae 100644
--- a/devices/rtx/device/frame/Frame.cu
+++ b/devices/rtx/device/frame/Frame.cu
@@ -316,6 +316,12 @@ void Frame::renderFrame()
       1));
   instrument::rangePop(); // optixLaunch()
 
+  // Increment frameID after rendering completes
+  if (checkerboarding())
+    hd.fb.frameID += int(hd.fb.checkerboardID == 3);
+  else
+    hd.fb.frameID += m_renderer->spp();
+
   if (m_denoise)
     m_denoiser.launch();
 
@@ -567,14 +573,13 @@ void *Frame::mapAlbedoBuffer(bool gpu)
 void *Frame::mapNormalBuffer(bool gpu)
 {
   auto &state = *deviceState();
-  const float invFrameID = m_invFrameID;
   auto begin = thrust::device_pointer_cast<vec3>((vec3 *)m_accumNormal.ptr());
   auto end = begin + numPixels();
   thrust::transform(thrust::cuda::par.on(state.stream),
       begin,
       end,
       thrust::device_pointer_cast<vec3>(m_normalBuffer.dataDevice()),
-      [=] __device__(const vec3 &in) { return in * invFrameID; });
+      [=] __device__(const vec3 &in) { return normalize(in); });
   if (gpu)
     return m_normalBuffer.dataDevice();
   else {
@@ -683,10 +688,6 @@ void Frame::newFrame()
           vec3(0.0f));
     }
   } else {
-    if (checkerboarding())
-      hd.fb.frameID += int(hd.fb.checkerboardID == 3);
-    else
-      hd.fb.frameID += m_renderer->spp();
     hd.fb.checkerboardID =
         checkerboarding() ? ((hd.fb.checkerboardID + 1) & 0x3) : -1;
   }
diff --git a/devices/rtx/device/gpu/computeAO.h b/devices/rtx/device/gpu/computeAO.h
index 059db6e09..77861164b 100644
--- a/devices/rtx/device/gpu/computeAO.h
+++ b/devices/rtx/device/gpu/computeAO.h
@@ -35,13 +35,12 @@
 
 namespace visrtx {
 
-template <typename T>
 VISRTX_DEVICE float computeAO(ScreenSample &ss,
     const Ray &primaryRay,
-    T rayType,
     const Hit &currentHit,
     float dist,
-    int numSamples)
+    int numSamples,
+    float (*surfaceAttenuation)(ScreenSample &, const Ray&))
 {
   float weights = 0.0f;
   float hits = 0.0f;
@@ -55,7 +54,7 @@ VISRTX_DEVICE float computeAO(ScreenSample &ss,
     float weight = max(0.f, dot(aoRay.dir, currentHit.Ns));
     if (weight > 1e-8f) {
       weights += weight;
-      hits += weight * surfaceAttenuation(ss, aoRay, rayType);
+      hits += weight * surfaceAttenuation(ss, aoRay);
     }
   }
 
diff --git a/devices/rtx/device/gpu/gpu_util.h b/devices/rtx/device/gpu/gpu_util.h
index ee1f6ecff..aaef86abb 100644
--- a/devices/rtx/device/gpu/gpu_util.h
+++ b/devices/rtx/device/gpu/gpu_util.h
@@ -133,6 +133,12 @@ VISRTX_DEVICE void accumulateValue(T &a, const T &b, float interp)
   a += b * (1.f - interp);
 }
 
+template <typename T>
+VISRTX_DEVICE void accumulateNormal(T &a, const T &b, float interp)
+{
+  accumulateValue(a, b, interp);
+}
+
 namespace detail {
 
 VISRTX_DEVICE void packPointer(void *ptr, uint32_t &i0, uint32_t &i1)
@@ -372,10 +378,8 @@ VISRTX_DEVICE uint32_t pixelIndex(
   return pixel.x + pixel.y * fb.size.x;
 }
 
-VISRTX_DEVICE void writeOutputColor(const FramebufferGPUData &fb,
-    const vec4 &color,
-    const uint32_t idx,
-    const int frameIDOffset)
+VISRTX_DEVICE void writeOutputColor(
+    const FramebufferGPUData &fb, const vec4 &color, const uint32_t idx)
 {
   if (fb.format == FrameFormat::SRGB) {
     fb.buffers.outColorUint[idx] =
@@ -388,16 +392,45 @@ VISRTX_DEVICE void writeOutputColor(const FramebufferGPUData &fb,
 
 } // namespace detail
 
+VISRTX_DEVICE void setPixelIds(const FramebufferGPUData &fb,
+    const uvec2 &pixel,
+    uint32_t primID,
+    uint32_t objID,
+    uint32_t instID)
+{
+  const uint32_t idx = detail::pixelIndex(fb, pixel);
+
+  if (fb.buffers.primID)
+    fb.buffers.primID[idx] = primID;
+  if (fb.buffers.objID)
+    fb.buffers.objID[idx] = objID;
+  if (fb.buffers.instID)
+    fb.buffers.instID[idx] = instID;
+}
+
+VISRTX_DEVICE void setPixelIds(const FramebufferGPUData &fb,
+    const uvec2 &pixel,
+    const float depth,
+    uint32_t primID,
+    uint32_t objID,
+    uint32_t instID)
+{
+  const uint32_t idx = detail::pixelIndex(fb, pixel);
+  if (detail::accumDepth(fb.buffers.depth, idx, depth)) {
+    if (fb.buffers.primID)
+      fb.buffers.primID[idx] = primID;
+    if (fb.buffers.objID)
+      fb.buffers.objID[idx] = objID;
+    if (fb.buffers.instID)
+      fb.buffers.instID[idx] = instID;
+  }
+}
 
-VISRTX_DEVICE void accumResults(const FrameGPUData &frame,
+VISRTX_DEVICE void accumPixelSample(const FrameGPUData &frame,
     const uvec2 &pixel,
     const vec4 &color,
-    float depth,
     const vec3 &albedo,
     const vec3 &normal,
-    uint32_t primID,
-    uint32_t objID,
-    uint32_t instID,
     const int frameIDOffset = 0)
 {
   const auto &fb = frame.fb;
@@ -406,21 +439,13 @@ VISRTX_DEVICE void accumResults(const FrameGPUData &frame,
 
   // Conditionally apply tonemapping during accumulation
   if (frame.renderer.tonemap)
-    detail::accumValue(fb.buffers.colorAccumulation, idx, detail::tonemap(color));
+    detail::accumValue(
+        fb.buffers.colorAccumulation, idx, detail::tonemap(color));
   else
     detail::accumValue(fb.buffers.colorAccumulation, idx, color);
   detail::accumValue(fb.buffers.albedo, idx, albedo);
   detail::accumValue(fb.buffers.normal, idx, normal);
 
-  if (detail::accumDepth(fb.buffers.depth, idx, depth)) {
-    if (fb.buffers.primID)
-      fb.buffers.primID[idx] = primID;
-    if (fb.buffers.objID)
-      fb.buffers.objID[idx] = objID;
-    if (fb.buffers.instID)
-      fb.buffers.instID[idx] = instID;
-  }
-
   const auto accumColor = fb.buffers.colorAccumulation[idx];
   // Conditionally apply inverse tonemapping on output
   const float frameDivisor = float(fb.frameID + frameIDOffset + 1);
@@ -429,25 +454,22 @@ VISRTX_DEVICE void accumResults(const FrameGPUData &frame,
       ? detail::inverseTonemap(normalizedColor)
       : normalizedColor;
 
-  detail::writeOutputColor(fb, outputColor, idx, frameIDOffset);
+  detail::writeOutputColor(fb, outputColor, idx);
 
   if (fb.checkerboardID == 0 && frameID == 0) {
     auto adjPix = uvec2(pixel.x + 1, pixel.y + 0);
     if (!pixelOutOfFrame(adjPix, fb)) {
-      detail::writeOutputColor(
-          fb, outputColor, detail::pixelIndex(fb, adjPix), frameIDOffset);
+      detail::writeOutputColor(fb, outputColor, detail::pixelIndex(fb, adjPix));
     }
 
     adjPix = uvec2(pixel.x + 0, pixel.y + 1);
     if (!pixelOutOfFrame(adjPix, fb)) {
-      detail::writeOutputColor(
-          fb, outputColor, detail::pixelIndex(fb, adjPix), frameIDOffset);
+      detail::writeOutputColor(fb, outputColor, detail::pixelIndex(fb, adjPix));
     }
 
     adjPix = uvec2(pixel.x + 1, pixel.y + 1);
     if (!pixelOutOfFrame(adjPix, fb)) {
-      detail::writeOutputColor(
-          fb, outputColor, detail::pixelIndex(fb, adjPix), frameIDOffset);
+      detail::writeOutputColor(fb, outputColor, detail::pixelIndex(fb, adjPix));
     }
   }
 }
diff --git a/devices/rtx/device/gpu/intersectRay.h b/devices/rtx/device/gpu/intersectRay.h
index 2c5ca27aa..aede85e96 100644
--- a/devices/rtx/device/gpu/intersectRay.h
+++ b/devices/rtx/device/gpu/intersectRay.h
@@ -102,12 +102,4 @@ VISRTX_DEVICE void intersectVolume(ScreenSample &ss,
   detail::launchRay(ss, r, rayType, false, dataPtr, optixFlags);
 }
 
-template <typename T>
-VISRTX_DEVICE float surfaceAttenuation(ScreenSample &ss, Ray r, T rayType)
-{
-  float a = 0.f;
-  intersectSurface(ss, r, rayType, &a, OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT);
-  return a;
-}
-
 } // namespace visrtx
diff --git a/devices/rtx/device/gpu/renderer/common.h b/devices/rtx/device/gpu/renderer/common.h
new file mode 100644
index 000000000..7431ee0db
--- /dev/null
+++ b/devices/rtx/device/gpu/renderer/common.h
@@ -0,0 +1,52 @@
+/*
+ * Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+ * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+ * POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#pragma once
+
+// Common constants and types used across multiple renderers
+
+namespace visrtx {
+
+// Threshold for considering a surface fully opaque
+constexpr float OPACITY_THRESHOLD = 0.99f;
+
+// Minimum contribution weight to continue tracing
+constexpr float MIN_CONTRIBUTION_EPSILON = 1.0e-8f;
+
+// Standard ray types used across all renderers
+// All renderers should use these standard types for consistency
+enum class RayType
+{
+  PRIMARY = 0, // Primary/shading rays
+  SHADOW = 1, // Shadow/occlusion rays (used for light visibility, AO, etc.)
+};
+
+} // namespace visrtx
diff --git a/devices/rtx/device/gpu/renderer/raygen_helpers.h b/devices/rtx/device/gpu/renderer/raygen_helpers.h
new file mode 100644
index 000000000..4b3274940
--- /dev/null
+++ b/devices/rtx/device/gpu/renderer/raygen_helpers.h
@@ -0,0 +1,207 @@
+/*
+ * Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: BSD-3-Clause
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *
+ * 1. Redistributions of source code must retain the above copyright notice,
+ * this list of conditions and the following disclaimer.
+ *
+ * 2. Redistributions in binary form must reproduce the above copyright notice,
+ * this list of conditions and the following disclaimer in the documentation
+ * and/or other materials provided with the distribution.
+ *
+ * 3. Neither the name of the copyright holder nor the names of its
+ * contributors may be used to endorse or promote products derived from
+ * this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
+ * LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+ * POSSIBILITY OF SUCH DAMAGE.
+ */
+
+#pragma once
+
+#include "gpu/evalShading.h"
+#include "gpu/intersectRay.h"
+#include "gpu/renderer/common.h"
+#include "gpu/shadingState.h"
+
+// Shared helper functions for ray generation across renderers
+
+namespace visrtx {
+
+// Compute surface attenuation for shadow/occlusion rays
+// Uses the standard SHADOW ray type for consistency across all renderers
+VISRTX_DEVICE float surfaceAttenuation(ScreenSample &ss, const Ray &r)
+{
+  float a = 0.0f;
+  intersectSurface(
+      ss, r, RayType::SHADOW, &a, OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT);
+  return a;
+}
+
+// Compute volume attenuation for shadow rays
+// Uses the standard SHADOW ray type for consistency across all renderers
+VISRTX_DEVICE float volumeAttenuation(ScreenSample &ss, const Ray &r)
+{
+  float attenuation = 0.0f;
+  intersectVolume(
+      ss, r, RayType::SHADOW, &attenuation, OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT);
+  return attenuation;
+}
+
+// Evaluate opacity including transmission
+VISRTX_DEVICE float evaluateOpacity(const MaterialShadingState &shadingState)
+{
+  return materialEvaluateOpacity(shadingState)
+      * (1.0f - glm::luminosity(materialEvaluateTransmission(shadingState)));
+}
+
+// Templated rendering loop
+// ShadingPolicy must implement:
+//   static VISRTX_DEVICE vec4 shadeSurface(
+//       const MaterialShadingState &shadingState,
+//       ScreenSample &ss,
+//       const Ray &ray,
+//       const SurfaceHit &hit)
+template <typename ShadingPolicy>
+VISRTX_DEVICE void renderPixel(FrameGPUData &frameData, ScreenSample ss)
+{
+  auto &rendererParams = frameData.renderer;
+
+  for (int i = 0; i < frameData.renderer.numIterations; i++) {
+    // First go with the main ray, pixel centered if first frame.
+    // Jittered samples are produced by next iterations.
+    bool isVeryFirstRay = i == 0 && ss.frameData->fb.frameID == 0;
+    auto ray = makePrimaryRay(ss, isVeryFirstRay);
+    float tmax = ray.t.upper;
+
+    // Output accumulators
+    vec3 outputColor(0.f);
+    vec3 outputAlbedo(0.f);
+    vec3 outputNormal(0.f);
+    float outputOpacity = 0.f;
+
+    // First hit metadata (for picking)
+    float depth = std::numeric_limits<float>::max();
+    uint32_t primID = ~0u;
+    uint32_t objID = ~0u;
+    uint32_t instID = ~0u;
+
+    // Transparency traversal loop
+    while (outputOpacity < OPACITY_THRESHOLD) {
+      ray.t.upper = tmax;
+
+      // Find next surface
+      SurfaceHit surfaceHit;
+      surfaceHit.foundHit = false;
+      intersectSurface(ss,
+          ray,
+          RayType::PRIMARY,
+          &surfaceHit,
+          primaryRayOptiXFlags(rendererParams));
+
+      float hitDist = surfaceHit.foundHit ? surfaceHit.t : ray.t.upper;
+
+      // Ray march volumes up to this surface hit
+      vec3 volumeColor(0.f);
+      float volumeOpacity = 0.f;
+      uint32_t volObjID = ~0u;
+      uint32_t volInstID = ~0u;
+
+      float volumeDepth = rayMarchAllVolumes(ss,
+          ray,
+          RayType::PRIMARY,
+          hitDist,
+          rendererParams.inverseVolumeSamplingRate,
+          volumeColor,
+          volumeOpacity,
+          volObjID,
+          volInstID);
+
+      // Accumulate volume contribution if any
+      if (volumeDepth < depth) {
+        accumulateValue(outputColor, volumeColor, outputOpacity);
+        accumulateValue(outputAlbedo, volumeColor, outputOpacity);
+        accumulateNormal(outputNormal, -ray.dir, outputOpacity);
+        accumulateValue(outputOpacity, volumeOpacity, outputOpacity);
+
+        depth = volumeDepth;
+        objID = volObjID;
+        instID = volInstID;
+        primID = volObjID;
+      }
+
+      // Handle surface hit if any
+      if (surfaceHit.foundHit) {
+        MaterialShadingState shadingState;
+        materialInitShading(
+            &shadingState, frameData, *surfaceHit.material, surfaceHit);
+
+        // Call the renderer-specific shading function
+        const vec4 surfaceColor =
+            ShadingPolicy::shadeSurface(shadingState, ss, ray, surfaceHit);
+
+        // Accumulate surface contribution
+        accumulateValue(
+            outputColor, vec3(surfaceColor) * surfaceColor.a, outputOpacity);
+        accumulateValue(outputAlbedo,
+            materialEvaluateTint(shadingState) * surfaceColor.a,
+            outputOpacity);
+        accumulateNormal(
+            outputNormal, materialEvaluateNormal(shadingState), outputOpacity);
+        accumulateValue(outputOpacity, surfaceColor.a, outputOpacity);
+
+        // Track first hit from surface
+        if (surfaceHit.t < depth) {
+          depth = surfaceHit.t;
+          primID = surfaceHit.primID;
+          objID = surfaceHit.objID;
+          instID = surfaceHit.instID;
+        }
+
+        // Advance ray past this surface for next iteration
+        ray.t.lower = surfaceHit.t + surfaceHit.epsilon;
+      }
+
+      // Record first hit metadata
+      if (isVeryFirstRay) {
+        setPixelIds(frameData.fb, ss.pixel, depth, primID, objID, instID);
+      }
+
+      // Exit if the current ray left the scene
+      if (!surfaceHit.foundHit)
+        break;
+
+      // Otherwise, continue through transparent surface
+    }
+
+    // Accumulate background for remaining transparency
+    const auto bg = getBackground(frameData, ss.screen, ray.dir);
+    const bool premultiplyBg = rendererParams.premultiplyBackground;
+    vec3 bgColor = premultiplyBg ? vec3(bg) * bg.a : vec3(bg);
+
+    accumulateValue(outputColor, bgColor, outputOpacity);
+    accumulateValue(outputOpacity, bg.a, outputOpacity);
+
+    // Write accumulated sample to framebuffer
+    accumPixelSample(frameData,
+        ss.pixel,
+        vec4(outputColor, outputOpacity),
+        outputAlbedo,
+        outputNormal,
+        i);
+  }
+}
+
+} // namespace visrtx
diff --git a/devices/rtx/device/gpu/volumeIntegration.h b/devices/rtx/device/gpu/volumeIntegration.h
index 88beb28cf..3d84ea304 100644
--- a/devices/rtx/device/gpu/volumeIntegration.h
+++ b/devices/rtx/device/gpu/volumeIntegration.h
@@ -123,6 +123,7 @@ VISRTX_DEVICE float _rayMarchVolume(ScreenSample &ss,
   float depth = std::numeric_limits<float>::max();
 
   // Accumulate until full opacity
+  constexpr float MIN_OPACITY_THRESHOLD = 1e-2f;
   constexpr float MAX_OPACITY_THRESHOLD = 0.99f;
   while (opacity < MAX_OPACITY_THRESHOLD && size(interval) >= 0.f) {
     const vec3 p = hit.localRay.org + hit.localRay.dir * interval.lower;
@@ -133,14 +134,15 @@ VISRTX_DEVICE float _rayMarchVolume(ScreenSample &ss,
 
       const float stepAlpha = 1.0f - glm::pow(1.0f - co.w, exponent);
       if (stepAlpha > 0.0f) {
-        // Record depth at first non-zero contribution
-        if (depth == std::numeric_limits<float>::max()) {
-          depth = interval.lower;
-        }
         const float weight = (1.0f - opacity);
         if (color)
           *color += weight * stepAlpha * vec3(co);
         opacity += weight * stepAlpha;
+
+        if (opacity > MIN_OPACITY_THRESHOLD
+            && depth == std::numeric_limits<float>::max()) {
+          depth = interval.lower;
+        }
       }
     }
 
diff --git a/devices/rtx/device/material/shaders/MDLShader_ptx.cu b/devices/rtx/device/material/shaders/MDLShader_ptx.cu
index c8faa9e9c..b812e320b 100644
--- a/devices/rtx/device/material/shaders/MDLShader_ptx.cu
+++ b/devices/rtx/device/material/shaders/MDLShader_ptx.cu
@@ -154,9 +154,7 @@ VISRTX_CALLABLE void __direct_callable__init(MDLShadingState *shadingState,
   shadingState->argBlock = md->argBlock;
 
   // Init
-  mdlInit(&shadingState->state,
-      &shadingState->resData,
-      shadingState->argBlock);
+  mdlInit(&shadingState->state, &shadingState->resData, shadingState->argBlock);
 }
 
 // Signature must match the call inside shaderMDLSurface in MDLShader.cuh.
@@ -202,6 +200,13 @@ VISRTX_CALLABLE
 NextRay __direct_callable__nextRay(
     const MDLShadingState *shadingState, const Ray *ray, RandState *rs)
 {
+  // Before anything, check for opacity. If below, then we just pass through
+  if (curand_uniform(rs) > mdlOpacity(&shadingState->state,
+          &shadingState->resData,
+          shadingState->argBlock)) {
+    return NextRay{ray->dir, vec3(1.0f)};
+  }
+
   // Sample
   BsdfSampleData sample_data = {};
   if (shadingState->isFrontFace) {
@@ -264,15 +269,15 @@ vec3 __direct_callable__evaluateEmission(
 }
 
 VISRTX_CALLABLE
-vec3 __direct_callable__evaluateTransmission(const MDLShadingState *shadingState)
+vec3 __direct_callable__evaluateTransmission(
+    const MDLShadingState *shadingState)
 {
   return mdlTransmission(
       &shadingState->state, &shadingState->resData, shadingState->argBlock);
 }
 
 VISRTX_CALLABLE
-vec3 __direct_callable__evaluateNormal(
-    const MDLShadingState *shadingState)
+vec3 __direct_callable__evaluateNormal(const MDLShadingState *shadingState)
 {
   return make_vec3(shadingState->state.normal);
 }
\ No newline at end of file
diff --git a/devices/rtx/device/material/shaders/PhysicallyBasedShader_ptx.cu b/devices/rtx/device/material/shaders/PhysicallyBasedShader_ptx.cu
index 2041808ca..7abbde0a4 100644
--- a/devices/rtx/device/material/shaders/PhysicallyBasedShader_ptx.cu
+++ b/devices/rtx/device/material/shaders/PhysicallyBasedShader_ptx.cu
@@ -32,7 +32,6 @@
 #include "gpu/gpu_decl.h"
 #include "gpu/gpu_math.h"
 #include "gpu/gpu_objects.h"
-#include "gpu/intersectRay.h"
 #include "gpu/sampleLight.h"
 #include "gpu/shadingState.h"
 #include "gpu/shading_api.h"
@@ -82,35 +81,6 @@ VISRTX_CALLABLE void __direct_callable__init(
       getMaterialParameter(*fd, md->transmission, *hit).x;
 }
 
-VISRTX_CALLABLE NextRay __direct_callable__nextRay(
-    const PhysicallyBasedShadingState *shadingState,
-    const Ray *ray,
-    RandState *rs)
-{
-  // Open cone, along the perfect reflection ray, with a metallic and
-  // roughness-dependent angle
-  const float roughness = shadingState->roughness;
-  const float metalness = shadingState->metallic;
-  const float roughnessSqr = roughness * roughness;
-  const float cosThetaMax = 1.0f - (roughnessSqr * roughnessSqr);
-  const float transmission = shadingState->transmission;
-
-  bool isReflected = curand_uniform(rs) > transmission;
-  auto nextVector = isReflected
-      ? glm::reflect(ray->dir, shadingState->normal)
-      : glm::refract(ray->dir, shadingState->normal, shadingState->ior);
-
-  auto nextRay = computeOrthonormalBasis(normalize(nextVector))
-      * uniformSampleCone(cosThetaMax,
-          vec3(curand_uniform(rs), curand_uniform(rs), curand_uniform(rs)));
-
-  auto nextSampleWeight = isReflected
-      ? shadingState->baseColor * metalness * (1.0f - transmission)
-      : shadingState->baseColor * transmission;
-
-  return NextRay{nextRay, nextSampleWeight};
-}
-
 VISRTX_CALLABLE
 vec3 __direct_callable__evaluateTint(
     const PhysicallyBasedShadingState *shadingState)
@@ -206,3 +176,38 @@ VISRTX_CALLABLE vec3 __direct_callable__shadeSurface(
   return (diffuseBRDF * (1.0f - shadingState->transmission) + specularBRDF)
       * NdotL * lightSample->radiance / lightSample->pdf;
 }
+
+VISRTX_CALLABLE NextRay __direct_callable__nextRay(
+    const PhysicallyBasedShadingState *shadingState,
+    const Ray *ray,
+    RandState *rs)
+{
+  // Before anything, check for opacity. If below, then we just pass through
+  if (curand_uniform(rs) > shadingState->opacity)
+  {
+    return NextRay{ray->dir, vec3(1.0f)};
+  }
+
+  // Open cone, along the perfect reflection ray, with a metallic and
+  // roughness-dependent angle
+  const float roughness = shadingState->roughness;
+  const float metalness = shadingState->metallic;
+  const float roughnessSqr = roughness * roughness;
+  const float cosThetaMax = 1.0f - (roughnessSqr * roughnessSqr);
+  const float transmission = shadingState->transmission;
+
+  bool isReflected = curand_uniform(rs) > transmission;
+  auto nextVector = isReflected
+        ? glm::reflect(ray->dir, shadingState->normal)
+          : glm::refract(ray->dir, shadingState->normal, shadingState->ior);
+
+  auto nextRay = computeOrthonormalBasis(normalize(nextVector))
+        * uniformSampleCone(cosThetaMax,
+            vec3(curand_uniform(rs), curand_uniform(rs), curand_uniform(rs)));
+
+  auto nextSampleWeight = isReflected
+    ? shadingState->baseColor * metalness * (1.0f - transmission)
+    : shadingState->baseColor * transmission;
+
+  return NextRay{nextRay, nextSampleWeight};
+}
\ No newline at end of file
diff --git a/devices/rtx/device/mdl/MaterialRegistry.cpp b/devices/rtx/device/mdl/MaterialRegistry.cpp
index 10c4f5007..e3373d9d7 100644
--- a/devices/rtx/device/mdl/MaterialRegistry.cpp
+++ b/devices/rtx/device/mdl/MaterialRegistry.cpp
@@ -157,6 +157,12 @@ MaterialRegistry::acquireMaterial(
   auto targetCode = make_handle(
       m_core->getPtxTargetCode(compiledMaterial.get(), transaction.get()));
   std::vector<libmdl::TextureDescriptor> textureDescs;
+  if (!targetCode.is_valid_interface()) {
+    m_core->logMessage(mi::base::MESSAGE_SEVERITY_ERROR,
+        "Failed generating PTX target code for material {}",
+        fullMaterialName);
+    return {};
+  }
 
   for (auto i = 1ul; i < targetCode->get_texture_count(); ++i) {
     libmdl::TextureDescriptor textureDesc{i - 1, targetCode->get_texture(i)};
diff --git a/devices/rtx/device/renderer/AmbientOcclusion.cpp b/devices/rtx/device/renderer/AmbientOcclusion.cpp
index 692e9b697..c32dad158 100644
--- a/devices/rtx/device/renderer/AmbientOcclusion.cpp
+++ b/devices/rtx/device/renderer/AmbientOcclusion.cpp
@@ -37,7 +37,7 @@ namespace visrtx {
 
 static const std::vector<HitgroupFunctionNames> g_aoHitNames = {
     {"__closesthit__primary", "__anyhit__primary"},
-    {"__closesthit__ao", "__anyhit__ao"}};
+    {"__closesthit__shadow", "__anyhit__shadow"}};
 
 static const std::vector<std::string> g_aoMissNames = {"__miss__", "__miss__"};
 
diff --git a/devices/rtx/device/renderer/AmbientOcclusion_ptx.cu b/devices/rtx/device/renderer/AmbientOcclusion_ptx.cu
index fd7229567..31d5eb0fd 100644
--- a/devices/rtx/device/renderer/AmbientOcclusion_ptx.cu
+++ b/devices/rtx/device/renderer/AmbientOcclusion_ptx.cu
@@ -29,28 +29,27 @@
  * POSSIBILITY OF SUCH DAMAGE.
  */
 
+#include <limits>
 #include "gpu/evalShading.h"
+#include "gpu/gpu_decl.h"
+#include "gpu/sampleLight.h"
 #include "gpu/shadingState.h"
 #include "gpu/shading_api.h"
+#include "gpu/renderer/common.h"
+#include "gpu/renderer/raygen_helpers.h"
 
 namespace visrtx {
 
-enum class RayType
-{
-  PRIMARY = 0,
-  AO = 1
-};
-
 DECLARE_FRAME_DATA(frameData)
 
 // OptiX programs /////////////////////////////////////////////////////////////
 
-VISRTX_GLOBAL void __closesthit__ao()
+VISRTX_GLOBAL void __closesthit__shadow()
 {
   // no-op
 }
 
-VISRTX_GLOBAL void __anyhit__ao()
+VISRTX_GLOBAL void __anyhit__shadow()
 {
   SurfaceHit hit;
   ray::populateSurfaceHit(hit);
@@ -62,7 +61,7 @@ VISRTX_GLOBAL void __anyhit__ao()
 
   auto &o = ray::rayData<float>();
   accumulateValue(o, materialEvaluateOpacity(shadingState), o);
-  if (o >= 0.99f)
+  if (o >= OPACITY_THRESHOLD)
     optixTerminateRay();
   else
     optixIgnoreIntersection();
@@ -83,145 +82,44 @@ VISRTX_GLOBAL void __miss__()
   // no-op
 }
 
-VISRTX_GLOBAL void __raygen__()
+// AmbientOcclusion shading policy for templated rendering loop /////////////
+
+struct AmbientOcclusionShadingPolicy
 {
-  auto &rendererParams = frameData.renderer;
-  auto &aoParams = rendererParams.params.ao;
+  static VISRTX_DEVICE vec4 shadeSurface(const MaterialShadingState &shadingState,
+      ScreenSample &ss,
+      const Ray &ray,
+      const SurfaceHit &hit)
+  {
+    const auto &rendererParams = frameData.renderer;
+    const auto &aoParams = rendererParams.params.ao;
+
+    const float aoFactor = aoParams.aoSamples > 0
+        ? computeAO(ss,
+              ray,
+              hit,
+              rendererParams.occlusionDistance,
+              aoParams.aoSamples,
+              &surfaceAttenuation)
+        : 1.f;
+
+    auto materialBaseColor = materialEvaluateTint(shadingState);
+    auto materialOpacity = materialEvaluateOpacity(shadingState);
+
+    const auto lighting =
+        aoFactor * rendererParams.ambientIntensity * rendererParams.ambientColor;
+
+    return vec4(materialBaseColor * lighting, materialOpacity);
+  }
+};
 
+VISRTX_GLOBAL void __raygen__()
+{
   auto ss = createScreenSample(frameData);
   if (pixelOutOfFrame(ss.pixel, frameData.fb))
     return;
 
-  for (int i = 0; i < frameData.renderer.numIterations; i++) {
-    auto ray = makePrimaryRay(ss);
-    float tmax = ray.t.upper;
-
-    SurfaceHit surfaceHit;
-    VolumeHit volumeHit;
-    vec3 outputColor(0.f);
-    vec3 outputNormal = ray.dir;
-    float outputOpacity = 0.f;
-    float depth = 1e30f;
-    uint32_t primID = ~0u;
-    uint32_t objID = ~0u;
-    uint32_t instID = ~0u;
-    bool firstHit = true;
-
-    while (outputOpacity < 0.99f) {
-      ray.t.upper = tmax;
-      surfaceHit.foundHit = false;
-      intersectSurface(ss,
-          ray,
-          RayType::PRIMARY,
-          &surfaceHit,
-          primaryRayOptiXFlags(rendererParams));
-
-      vec3 color(0.f);
-      float opacity = 0.f;
-
-      if (surfaceHit.foundHit) {
-        uint32_t vObjID = ~0u;
-        uint32_t vInstID = ~0u;
-        const float vDepth = rayMarchAllVolumes(ss,
-            ray,
-            RayType::PRIMARY,
-            surfaceHit.t,
-            rendererParams.inverseVolumeSamplingRate,
-            color,
-            opacity,
-            vObjID,
-            vInstID);
-
-        if (firstHit) {
-          const bool volumeFirst = vDepth < surfaceHit.t;
-          if (volumeFirst) {
-            outputNormal = -ray.dir;
-            depth = vDepth;
-            primID = 0;
-            objID = vObjID;
-            instID = vInstID;
-          } else {
-            outputNormal = surfaceHit.Ns;
-            depth = surfaceHit.t;
-            primID = computeGeometryPrimId(surfaceHit);
-            objID = surfaceHit.objID;
-            instID = surfaceHit.instID;
-          }
-          firstHit = false;
-        }
-
-        const float aoFactor = aoParams.aoSamples > 0
-            ? computeAO(ss,
-                  ray,
-                  RayType::AO,
-                  surfaceHit,
-                  rendererParams.occlusionDistance,
-                  aoParams.aoSamples)
-            : 1.f;
-
-        MaterialShadingState shadingState;
-        materialInitShading(
-            &shadingState, frameData, *surfaceHit.material, surfaceHit);
-        auto materialBaseColor = materialEvaluateTint(shadingState);
-        auto materialOpacity = materialEvaluateOpacity(shadingState);
-
-        const auto lighting = aoFactor * rendererParams.ambientIntensity
-            * rendererParams.ambientColor;
-        accumulateValue(color, materialBaseColor * lighting, opacity);
-        accumulateValue(opacity, materialOpacity, opacity);
-
-        color *= opacity;
-        accumulateValue(outputColor, color, outputOpacity);
-        accumulateValue(outputOpacity, opacity, outputOpacity);
-
-        ray.t.lower = surfaceHit.t + surfaceHit.epsilon;
-      } else {
-        uint32_t vObjID = ~0u;
-        uint32_t vInstID = ~0u;
-        const float volumeDepth = rayMarchAllVolumes(ss,
-            ray,
-            RayType::PRIMARY,
-            ray.t.upper,
-            rendererParams.inverseVolumeSamplingRate,
-            color,
-            opacity,
-            vObjID,
-            vInstID);
-
-        if (firstHit) {
-          if (opacity > 0.f) {
-            outputNormal = -ray.dir;
-          }
-          depth = min(depth, volumeDepth);
-          primID = 0;
-          objID = vObjID;
-          instID = vInstID;
-        }
-
-        color *= opacity;
-
-        const auto bg = getBackground(frameData, ss.screen, ray.dir);
-        const bool premultiplyBg = rendererParams.premultiplyBackground;
-        accumulateValue(
-            color, premultiplyBg ? vec3(bg) * bg.a : vec3(bg), opacity);
-        accumulateValue(opacity, bg.a, opacity);
-        accumulateValue(outputColor, color, outputOpacity);
-        accumulateValue(outputOpacity, opacity, outputOpacity);
-        break;
-      }
-    }
-
-    accumResults(frameData,
-        ss.pixel,
-        vec4(outputColor, outputOpacity),
-        depth,
-        outputColor,
-        outputNormal,
-        primID,
-        objID,
-        instID,
-        i);
-  }
+  renderPixel<AmbientOcclusionShadingPolicy>(frameData, ss);
 }
 
 } // namespace visrtx
diff --git a/devices/rtx/device/renderer/Debug_ptx.cu b/devices/rtx/device/renderer/Debug_ptx.cu
index e56992bc9..bf729925d 100644
--- a/devices/rtx/device/renderer/Debug_ptx.cu
+++ b/devices/rtx/device/renderer/Debug_ptx.cu
@@ -30,6 +30,7 @@
  */
 
 #include "DebugMethod.h"
+#include "gpu/intersectRay.h"
 #include "gpu/shading_api.h"
 
 namespace visrtx {
@@ -253,15 +254,9 @@ VISRTX_GLOBAL void __raygen__()
     instID = vrd.instance->id;
   }
 
-  accumResults(frameData,
-      ss.pixel,
-      vec4(color, 1.f),
-      depth,
-      color,
-      normal,
-      primID,
-      objID,
-      instID);
+  setPixelIds(frameData.fb, ss.pixel, depth, primID, objID, instID);
+
+  accumPixelSample(frameData, ss.pixel, vec4(color, 1.f), color, normal);
 }
 
 } // namespace visrtx
diff --git a/devices/rtx/device/renderer/DirectLight.cpp b/devices/rtx/device/renderer/DirectLight.cpp
index 9549b656d..3d0970670 100644
--- a/devices/rtx/device/renderer/DirectLight.cpp
+++ b/devices/rtx/device/renderer/DirectLight.cpp
@@ -36,12 +36,13 @@
 namespace visrtx {
 
 static const std::vector<HitgroupFunctionNames> g_directLightHitNames = {
-    {"__closesthit__primary", "__anyhit__primary"},
+    {"__closesthit__shading", "__anyhit__shading"},
     {"__closesthit__shadow", "__anyhit__shadow"},
-    {"__closesthit__bounce", "__anyhit__bounce"}};
+};
 
 static const std::vector<std::string> g_directLightMissNames = {
-    "__miss__", "__miss__", "__miss__"};
+    "__miss__shading", "__miss__shadow"
+};
 
 DirectLight::DirectLight(DeviceGlobalState *s) : Renderer(s, 0.f) {}
 
diff --git a/devices/rtx/device/renderer/DirectLight_ptx.cu b/devices/rtx/device/renderer/DirectLight_ptx.cu
index 8fc5620f2..32c3196c4 100644
--- a/devices/rtx/device/renderer/DirectLight_ptx.cu
+++ b/devices/rtx/device/renderer/DirectLight_ptx.cu
@@ -29,12 +29,12 @@
  * POSSIBILITY OF SUCH DAMAGE.
  */
 
-#include <curand.h>
 #include <cmath>
 #include <glm/common.hpp>
 #include <glm/ext/vector_float4.hpp>
 #include <glm/geometric.hpp>
 #include <glm/vector_relational.hpp>
+#include <limits>
 #include "gpu/evalShading.h"
 #include "gpu/gpu_math.h"
 #include "gpu/gpu_objects.h"
@@ -42,37 +42,21 @@
 #include "gpu/sampleLight.h"
 #include "gpu/shadingState.h"
 #include "gpu/shading_api.h"
+#include "gpu/renderer/common.h"
+#include "gpu/renderer/raygen_helpers.h"
 
 namespace visrtx {
 
-enum class RayType
-{
-  PRIMARY = 0,
-  SHADOW = 1,
-  BOUNCE = 2
-};
-
-struct RayAttenuation
-{
-  const Ray *ray{nullptr};
-  float attenuation{0.f};
-};
-
 DECLARE_FRAME_DATA(frameData)
 
-// Helper functions ///////////////////////////////////////////////////////////
+// DirectLight shading policy for templated rendering loop //////////////////
 
-VISRTX_DEVICE float volumeAttenuation(ScreenSample &ss, Ray r)
+struct DirectLightShadingPolicy
 {
-  RayAttenuation ra;
-  ra.ray = &r;
-  intersectVolume(
-      ss, r, RayType::SHADOW, &ra, OPTIX_RAY_FLAG_DISABLE_CLOSESTHIT);
-  return ra.attenuation;
-}
-
-VISRTX_DEVICE vec4 shadeSurface(
-    ScreenSample &ss, const Ray &ray, const SurfaceHit &hit)
+  static VISRTX_DEVICE vec4 shadeSurface(const MaterialShadingState &shadingState,
+    ScreenSample &ss,
+    const Ray &ray,
+    const SurfaceHit &hit)
 {
   const auto &rendererParams = frameData.renderer;
   const auto &directLightParams = rendererParams.params.directLight;
@@ -83,34 +67,18 @@ VISRTX_DEVICE vec4 shadeSurface(
   const float aoFactor = directLightParams.aoSamples > 0
       ? computeAO(ss,
             ray,
-            RayType::SHADOW,
             hit,
             rendererParams.occlusionDistance,
-            directLightParams.aoSamples)
+            directLightParams.aoSamples,
+            &surfaceAttenuation)
       : 1.f;
 
-  MaterialShadingState shadingState;
-  materialInitShading(&shadingState, frameData, *hit.material, hit);
-
   vec3 contrib = materialEvaluateEmission(shadingState, -ray.dir);
-  float opacity = materialEvaluateOpacity(shadingState);
 
   // Handle ambient light contribution
   if (rendererParams.ambientIntensity > 0.0f) {
-#define USE_SAMPLED_AMBIENT_LIGHT 0
-#if USE_SAMPLED_AMBIENT_LIGHT
-    const LightSample ls = {
-        .radiance =
-            rendererParams.ambientColor * rendererParams.ambientIntensity,
-        .dir = sampleHemisphere(ss.rs, hit.Ns),
-        .dist = 1e20f,
-        .pdf = 1.0f,
-    };
-    contrib = materialShadeSurface(shadingState, hit, ls, -ray.dir);
-#else
     contrib += rendererParams.ambientColor * rendererParams.ambientIntensity
         * materialEvaluateTint(shadingState);
-#endif
   }
 
   // Handle all lights contributions
@@ -125,26 +93,25 @@ VISRTX_DEVICE vec4 shadeSurface(
       if (lightSample.pdf == 0.0f)
         continue;
 
+      // Shadowing
       const Ray shadowRay = {
           hit.hitpoint + hit.Ng * hit.epsilon,
           lightSample.dir,
           {hit.epsilon, lightSample.dist},
       };
 
-      const float surface_o =
-          1.f - surfaceAttenuation(ss, shadowRay, RayType::SHADOW);
-      const float volume_o = 1.f - volumeAttenuation(ss, shadowRay);
-      const float attenuation = surface_o * volume_o;
+      const float surface_attenuation =
+          1.0f - surfaceAttenuation(ss, shadowRay);
+      const float volume_attenuation = 1.0f - volumeAttenuation(ss, shadowRay);
+      const float attenuation = surface_attenuation * volume_attenuation;
 
-      if (attenuation <= 1.0e-12f)
+      // Complete occlusion?
+      if (attenuation <= MIN_CONTRIBUTION_EPSILON)
         continue;
 
       const vec3 thisLightContrib =
           materialShadeSurface(shadingState, hit, lightSample, -ray.dir);
 
-      if (glm::any(glm::isnan(thisLightContrib)))
-        continue;
-
       contrib += thisLightContrib * attenuation;
     }
   }
@@ -152,92 +119,47 @@ VISRTX_DEVICE vec4 shadeSurface(
   // Take AO in account
   contrib *= aoFactor;
 
-  // Then proceed with the bounce rays
-  vec3 nextRayContrib(0.0f);
-  vec3 nextRayContribWeight = vec3(1.f);
-
-  Ray bounceRay = ray;
+  // Then proceed with single bounce ray for indirect lighting
   SurfaceHit bounceHit = hit;
-
-  for (int depth = 0; depth < frameData.renderer.maxRayDepth; ++depth) {
-    NextRay nextRay = materialNextRay(shadingState, bounceRay, ss.rs);
-
-    if (glm::all(glm::lessThan(
-            glm::abs(vec3(nextRay.direction)), glm::vec3(1.0e-8f))))
-      break;
-
-    nextRayContribWeight *= vec3(nextRay.contributionWeight);
-    if (glm::all(glm::lessThan(nextRayContribWeight, glm::vec3(1.0e-8f))))
-      break;
-
-    bounceRay = {
+  NextRay nextRay = materialNextRay(shadingState, ray, ss.rs);
+  if (glm::any(glm::greaterThan(
+          nextRay.contributionWeight, glm::vec3(MIN_CONTRIBUTION_EPSILON)))) {
+    Ray bounceRay = {
         bounceHit.hitpoint
             + bounceHit.Ng
                 * std::copysignf(
                     bounceHit.epsilon, dot(bounceHit.Ns, nextRay.direction)),
-        normalize(vec3(nextRay.direction)),
+        nextRay.direction,
     };
 
+    // Only check for intersecting surfaces and background as secondary light
+    // interactions
     bounceHit.foundHit = false;
-    bounceHit.isFrontFace = true;
-    intersectSurface(ss, bounceRay, RayType::BOUNCE, &bounceHit);
+    intersectSurface(ss, bounceRay, RayType::PRIMARY, &bounceHit);
 
-    // We hit something. Gather its contribution.
     if (bounceHit.foundHit) {
-      // This HDRI search is not ideal. It does not account for light instance
-      // transformations and should be reworked later on.
-      auto hdri = (frameData.world.hdri != -1)
-          ? &frameData.registry.lights[frameData.world.hdri]
-          : nullptr;
-
-      LightSample lightSample;
-      // If we have an active HDRI, sample it.
-      if (hdri && hdri->hdri.visible) {
-        lightSample = detail::sampleHDRILight(
-            *hdri, glm::identity<mat4>(), bounceHit, ss.rs);
-      } else {
-        // Otherwise fallback to some simple background probing.
-        lightSample = {
-            rendererParams.ambientColor * rendererParams.ambientIntensity,
-            bounceHit.Ns,
-            std::numeric_limits<float>::max(),
-            1.0f / (4.0f * float(M_PI)),
-        };
-      }
+      // We hit something. Gather its contribution, cosine weighted diffuse
+      // only, we want this to be lightweight.
+      MaterialShadingState bounceShadingState;
       materialInitShading(
-          &shadingState, frameData, *bounceHit.material, bounceHit);
-      nextRayContrib = materialShadeSurface(
-          shadingState, bounceHit, lightSample, bounceRay.dir);
-
-      if (glm::any(glm::isnan(nextRayContrib))) {
-        break;
-      }
-      contrib += nextRayContrib * nextRayContribWeight;
+          &bounceShadingState, frameData, *bounceHit.material, bounceHit);
+
+      auto sampleDir = randomDir(ss.rs, bounceHit.Ns);
+      auto cosineT = dot(bounceHit.Ns, sampleDir);
+      auto color = materialEvaluateTint(bounceShadingState) * cosineT;
+      contrib += color * nextRay.contributionWeight;
     } else {
-      // This HDRI search is not ideal. It does not account for light instance
-      // transformations and should be reworked later on.
-      auto hdri = (frameData.world.hdri != -1)
-          ? &frameData.registry.lights[frameData.world.hdri]
-          : nullptr;
-      // No hit, get background contribution.
-      vec3 radiance;
-      // If we have an active HDRI, sample it.
-      if (hdri && hdri->hdri.visible) {
-        radiance =
-            detail::sampleHDRILight(*hdri, glm::identity<mat4>(), bounceRay.dir)
-                .radiance;
-      } else {
-        radiance =
-            rendererParams.ambientColor * rendererParams.ambientIntensity;
-      }
-      nextRayContrib = radiance;
-      contrib += nextRayContrib * nextRayContribWeight;
-      break;
+      // No hit, get background contribution directly (no surface to weight
+      // against)
+      const auto color = getBackground(frameData, ss.screen, bounceRay.dir);
+      contrib += vec3(color) * nextRay.contributionWeight;
     }
   }
 
+  float opacity = evaluateOpacity(shadingState);
   return vec4(contrib, opacity);
-}
+  }
+};
 
 // OptiX programs /////////////////////////////////////////////////////////////
 
@@ -246,6 +168,11 @@ VISRTX_GLOBAL void __closesthit__shadow()
   // no-op
 }
 
+VISRTX_GLOBAL void __miss__shadow()
+{
+  // no-op
+}
+
 VISRTX_GLOBAL void __anyhit__shadow()
 {
   auto &rendererParams = frameData.renderer.params;
@@ -256,180 +183,57 @@ VISRTX_GLOBAL void __anyhit__shadow()
 
     MaterialShadingState shadingState;
     materialInitShading(&shadingState, frameData, *hit.material, hit);
-    auto opacity = materialEvaluateOpacity(shadingState);
+    auto opacity = evaluateOpacity(shadingState);
 
     auto &o = ray::rayData<float>();
 
     accumulateValue(o, opacity, o);
 
-    if (o >= 0.99f)
+    if (o >= OPACITY_THRESHOLD)
       optixTerminateRay();
     else
       optixIgnoreIntersection();
   } else {
-    auto &ra = ray::rayData<RayAttenuation>();
+    auto &attenuation = ray::rayData<float>();
     VolumeHit hit;
     ray::populateVolumeHit(hit);
     rayMarchVolume(ray::screenSample(),
         hit,
-        ra.attenuation,
+        attenuation,
         rendererParams.directLight.inverseVolumeSamplingRateShadows);
-    if (ra.attenuation < 0.99f)
+    if (attenuation < OPACITY_THRESHOLD)
       optixIgnoreIntersection();
   }
 }
 
-VISRTX_GLOBAL void __anyhit__primary()
+VISRTX_GLOBAL void __anyhit__shading()
 {
   ray::cullbackFaces();
 }
 
-VISRTX_GLOBAL void __closesthit__primary()
+VISRTX_GLOBAL void __closesthit__shading()
 {
   ray::populateHit();
 }
 
-VISRTX_GLOBAL void __anyhit__bounce()
+VISRTX_GLOBAL void __miss__shading()
 {
-  ray::cullbackFaces();
-}
-
-VISRTX_GLOBAL void __closesthit__bounce()
-{
-  ray::populateHit();
-}
-
-VISRTX_GLOBAL void __miss__()
-{
-  // TODO
+  if (ray::isIntersectingSurfaces()) {
+    auto &hit = ray::rayData<SurfaceHit>();
+    hit.foundHit = false;
+  } else {
+    auto &hit = ray::rayData<VolumeHit>();
+    hit.foundHit = false;
+  }
 }
 
 VISRTX_GLOBAL void __raygen__()
 {
-  auto &rendererParams = frameData.renderer;
-
   auto ss = createScreenSample(frameData);
   if (pixelOutOfFrame(ss.pixel, frameData.fb))
     return;
 
-  for (int i = 0; i < frameData.renderer.numIterations; i++) {
-    auto ray = makePrimaryRay(ss);
-    float tmax = ray.t.upper;
-
-    SurfaceHit surfaceHit;
-    VolumeHit volumeHit;
-    vec3 outputColor(0.f);
-    vec3 outputNormal = ray.dir;
-    float outputOpacity = 0.f;
-    float depth = 1e30f;
-    uint32_t primID = ~0u;
-    uint32_t objID = ~0u;
-    uint32_t instID = ~0u;
-    bool firstHit = true;
-
-    while (outputOpacity < 0.99f) {
-      ray.t.upper = tmax;
-      surfaceHit.foundHit = false;
-      intersectSurface(ss,
-          ray,
-          RayType::PRIMARY,
-          &surfaceHit,
-          primaryRayOptiXFlags(rendererParams));
-
-      vec3 color(0.f);
-      float opacity = 0.f;
-
-      if (surfaceHit.foundHit) {
-        uint32_t vObjID = ~0u;
-        uint32_t vInstID = ~0u;
-        const float vDepth = rayMarchAllVolumes(ss,
-            ray,
-            RayType::PRIMARY,
-            surfaceHit.t,
-            rendererParams.inverseVolumeSamplingRate,
-            color,
-            opacity,
-            vObjID,
-            vInstID);
-
-        if (firstHit) {
-          const bool volumeFirst = vDepth < surfaceHit.t;
-          if (volumeFirst) {
-            outputNormal = -ray.dir;
-            depth = vDepth;
-            primID = 0;
-            objID = vObjID;
-            instID = vInstID;
-          } else {
-            outputNormal = surfaceHit.Ns;
-            depth = surfaceHit.t;
-            primID = computeGeometryPrimId(surfaceHit);
-            objID = surfaceHit.objID;
-            instID = surfaceHit.instID;
-          }
-          firstHit = false;
-        }
-
-        const vec4 shadingResult = shadeSurface(ss, ray, surfaceHit);
-        if (glm::any(glm::isnan(vec3(shadingResult)))) {
-          color = vec3(0.f);
-          opacity = 0.f;
-        }
-        accumulateValue(color, vec3(shadingResult), opacity);
-        accumulateValue(opacity, shadingResult.w, opacity);
-
-        color *= opacity;
-        accumulateValue(outputColor, color, outputOpacity);
-        accumulateValue(outputOpacity, opacity, outputOpacity);
-
-        ray.t.lower = surfaceHit.t + surfaceHit.epsilon;
-      } else {
-        uint32_t vObjID = ~0u;
-        uint32_t vInstID = ~0u;
-        const float volumeDepth = rayMarchAllVolumes(ss,
-            ray,
-            RayType::PRIMARY,
-            ray.t.upper,
-            rendererParams.inverseVolumeSamplingRate,
-            color,
-            opacity,
-            vObjID,
-            vInstID);
-
-        if (firstHit) {
-          if (opacity > 0.f) {
-            outputNormal = -ray.dir;
-          }
-          depth = min(depth, volumeDepth);
-          primID = 0;
-          objID = vObjID;
-          instID = vInstID;
-        }
-
-        color *= opacity;
-
-        const auto bg = getBackground(frameData, ss.screen, ray.dir);
-        const bool premultiplyBg = rendererParams.premultiplyBackground;
-        accumulateValue(
-            color, premultiplyBg ? vec3(bg) * bg.a : vec3(bg), opacity);
-        accumulateValue(opacity, bg.a, opacity);
-        accumulateValue(outputColor, color, outputOpacity);
-        accumulateValue(outputOpacity, opacity, outputOpacity);
-        break;
-      }
-    }
-
-    accumResults(frameData,
-        ss.pixel,
-        vec4(outputColor, outputOpacity),
-        depth,
-        outputColor,
-        outputNormal,
-        primID,
-        objID,
-        instID,
-        i);
-  }
+  renderPixel<DirectLightShadingPolicy>(frameData, ss);
 }
 
 } // namespace visrtx
diff --git a/devices/rtx/device/renderer/PathTracer_ptx.cu b/devices/rtx/device/renderer/PathTracer_ptx.cu
index 646cda0f3..167c4d26d 100644
--- a/devices/rtx/device/renderer/PathTracer_ptx.cu
+++ b/devices/rtx/device/renderer/PathTracer_ptx.cu
@@ -33,6 +33,7 @@
 
 #include "gpu/evalShading.h"
 #include "gpu/gpu_debug.h"
+#include "gpu/sampleLight.h"
 #include "gpu/shadingState.h"
 #include "gpu/shading_api.h"
 namespace visrtx {
@@ -201,16 +202,10 @@ VISRTX_GLOBAL void __raygen__()
       color = vec3(1) - color;
     if (debug())
       printf("========== END: FrameID %i ==========\n", frameData.fb.frameID);
-    accumResults(frameData,
-        ss.pixel,
-        vec4(color, 1.f),
-        outDepth,
-        outColor,
-        outNormal,
-        primID,
-        objID,
-        instID,
-        i);
+    setPixelIds(frameData.fb, ss.pixel, outDepth, primID, objID, instID);
+
+    accumPixelSample(
+        frameData, ss.pixel, vec4(color, 1.f), outColor, outNormal, i);
   }
 }
 
diff --git a/devices/rtx/device/renderer/Raycast_ptx.cu b/devices/rtx/device/renderer/Raycast_ptx.cu
index 9ab4f9a99..85e87eb0a 100644
--- a/devices/rtx/device/renderer/Raycast_ptx.cu
+++ b/devices/rtx/device/renderer/Raycast_ptx.cu
@@ -30,16 +30,14 @@
  */
 
 #include "gpu/evalShading.h"
+#include "gpu/sampleLight.h"
 #include "gpu/shadingState.h"
 #include "gpu/shading_api.h"
+#include "gpu/renderer/common.h"
+#include "gpu/renderer/raygen_helpers.h"
 
 namespace visrtx {
 
-enum class RayType
-{
-  PRIMARY
-};
-
 DECLARE_FRAME_DATA(frameData)
 
 VISRTX_GLOBAL void __anyhit__primary()
@@ -57,130 +55,36 @@ VISRTX_GLOBAL void __miss__()
   // no-op
 }
 
-VISRTX_GLOBAL void __raygen__()
+// Raycast shading policy for templated rendering loop //////////////////////
+
+struct RaycastShadingPolicy
 {
-  auto &rendererParams = frameData.renderer;
+  static VISRTX_DEVICE vec4 shadeSurface(
+      const MaterialShadingState &shadingState,
+      ScreenSample &ss,
+      const Ray &ray,
+      const SurfaceHit &hit)
+  {
+    const auto &rendererParams = frameData.renderer;
+
+    // Simple cosine-weighted diffuse lighting
+    auto materialBaseColor = materialEvaluateTint(shadingState);
+    auto materialOpacity = materialEvaluateOpacity(shadingState);
+
+    const auto lighting = glm::abs(glm::dot(ray.dir, hit.Ns))
+        * rendererParams.ambientColor;
+
+    return vec4(materialBaseColor * lighting, materialOpacity);
+  }
+};
 
+VISRTX_GLOBAL void __raygen__()
+{
   auto ss = createScreenSample(frameData);
   if (pixelOutOfFrame(ss.pixel, frameData.fb))
     return;
-  auto ray = makePrimaryRay(ss, true /*pixel centered*/);
-  float tmax = ray.t.upper;
-
-  vec3 outputColor(0.f);
-  vec3 outputNormal = ray.dir;
-  float outputOpacity = 0.f;
-  float depth = 1e30f;
-  uint32_t primID = ~0u;
-  uint32_t objID = ~0u;
-  uint32_t instID = ~0u;
-  bool firstHit = true;
-
-  while (outputOpacity < 0.99f) {
-    SurfaceHit surfaceHit;
-    ray.t.upper = tmax;
-    surfaceHit.foundHit = false;
-    intersectSurface(ss,
-        ray,
-        RayType::PRIMARY,
-        &surfaceHit,
-        primaryRayOptiXFlags(rendererParams));
-
-    vec3 color(0.f);
-    float opacity = 0.f;
-
-    if (surfaceHit.foundHit) {
-      uint32_t vObjID = ~0u;
-      uint32_t vInstID = ~0u;
-      const float vDepth = rayMarchAllVolumes(ss,
-          ray,
-          RayType::PRIMARY,
-          surfaceHit.t,
-          rendererParams.inverseVolumeSamplingRate,
-          color,
-          opacity,
-          vObjID,
-          vInstID);
-
-      if (firstHit) {
-        const bool volumeFirst = vDepth < surfaceHit.t;
-        if (volumeFirst) {
-          outputNormal = -ray.dir;
-          depth = vDepth;
-          primID = 0;
-          objID = vObjID;
-          instID = vInstID;
-        } else {
-          outputNormal = surfaceHit.Ng;
-          depth = surfaceHit.t;
-          primID = computeGeometryPrimId(surfaceHit);
-          objID = surfaceHit.objID;
-          instID = surfaceHit.instID;
-        }
-        firstHit = false;
-      }
-
-      MaterialShadingState shadingState;
-      materialInitShading(
-          &shadingState, frameData, *surfaceHit.material, surfaceHit);
-      auto materialBaseColor = materialEvaluateTint(shadingState);
-      auto materialOpacity = materialEvaluateOpacity(shadingState);
-
-      const auto lighting = glm::abs(glm::dot(ray.dir, surfaceHit.Ns))
-          * rendererParams.ambientColor;
-      accumulateValue(color, materialBaseColor * lighting, opacity);
-      accumulateValue(opacity, materialOpacity, opacity);
-
-      color *= opacity;
-      accumulateValue(outputColor, color, outputOpacity);
-      accumulateValue(outputOpacity, opacity, outputOpacity);
-
-      ray.t.lower = surfaceHit.t + surfaceHit.epsilon;
-    } else {
-      uint32_t vObjID = ~0u;
-      uint32_t vInstID = ~0u;
-      const float volumeDepth = rayMarchAllVolumes(ss,
-          ray,
-          RayType::PRIMARY,
-          ray.t.upper,
-          rendererParams.inverseVolumeSamplingRate,
-          color,
-          opacity,
-          vObjID,
-          vInstID);
-
-      if (firstHit) {
-        if (opacity > 0.f) {
-          outputNormal = -ray.dir;
-        }
-        depth = min(depth, volumeDepth);
-        primID = 0;
-        objID = vObjID;
-        instID = vInstID;
-      }
-
-      color *= opacity;
-
-      const auto bg = getBackground(frameData, ss.screen, ray.dir);
-      const bool premultiplyBg = rendererParams.premultiplyBackground;
-      accumulateValue(
-          color, premultiplyBg ? vec3(bg) * bg.a : vec3(bg), opacity);
-      accumulateValue(opacity, bg.a, opacity);
-      accumulateValue(outputColor, color, outputOpacity);
-      accumulateValue(outputOpacity, opacity, outputOpacity);
-      break;
-    }
-  }
 
-  accumResults(frameData,
-      ss.pixel,
-      vec4(outputColor, outputOpacity),
-      depth,
-      outputColor,
-      outputNormal,
-      primID,
-      objID,
-      instID);
+  renderPixel<RaycastShadingPolicy>(frameData, ss);
 }
 
 } // namespace visrtx
diff --git a/devices/rtx/device/renderer/Test_ptx.cu b/devices/rtx/device/renderer/Test_ptx.cu
index a854f569e..b2d3de740 100644
--- a/devices/rtx/device/renderer/Test_ptx.cu
+++ b/devices/rtx/device/renderer/Test_ptx.cu
@@ -57,15 +57,9 @@ VISRTX_GLOBAL void __raygen__()
     return;
   auto ray = makePrimaryRay(ss);
 
-  accumResults(frameData,
-      ss.pixel,
-      vec4(ray.dir, 1.f),
-      1.f,
-      ray.dir,
-      -ray.dir,
-      ~0u,
-      ~0u,
-      ~0u);
+  setPixelIds(frameData.fb, ss.pixel, 1.0f, ~0u, ~0u, ~0u);
+
+  accumPixelSample(frameData, ss.pixel, vec4(ray.dir, 1.f), ray.dir, -ray.dir);
 }
 
 } // namespace visrtx