Cleaning unused imports, refactoring wavefront integration

kernels/src/intersect.rs

@@ -0,0 +1,513 @@
+#[cfg(target_os = "cuda")]
+pub mod device {
+    use shared::core::aggregates::{BVHAggregate, LinearBVHNode};
+    use shared::core::geometry::{Bounds3f, Normal3f, Point2f, Point3f, Ray, Vector3f};
+    use shared::core::interaction::LightSampleContext;
+    use shared::core::material::Material;
+    use shared::core::medium::MediumInterface;
+    use shared::core::primitive::{Primitive, PrimitiveTrait};
+    use shared::spectra::{SampledSpectrum, SampledWavelengths};
+    use shared::utils::atomic::GpuAtomicU32;
+    use shared::utils::soa::SoABuffer;
+    use shared::wavefront::work_items::*;
+    use shared::{Float, Ptr};
+ 
+    use cuda_std::*;
+ 
+    #[repr(C)]
+    pub struct IntersectClosestParams {
+        pub bvh: Ptr<BVHAggregate>,
+ 
+        // Input queue
+        pub ray_q: Ptr<RayQueue>,
+ 
+        // Output queues
+        pub escaped_ray_q: Ptr<EscapedRayQueue>,
+        pub hit_area_light_q: Ptr<HitAreaLightQueue>,
+        pub basic_eval_mtl_q: Ptr<MaterialEvalQueue>,
+        pub universal_eval_mtl_q: Ptr<MaterialEvalQueue>,
+        pub next_ray_q: Ptr<RayQueue>,
+ 
+        // Persistent state
+        pub pixel_sample_state: Ptr<PixelSampleState>,
+ 
+        pub n_rays: u32,
+    }
+ 
+    /// One thread per ray: traverse BVH, push results to output queues.
+    #[kernel]
+    pub unsafe fn intersect_closest(params: &IntersectClosestParams) {
+        let idx = thread::index_1d();
+        if idx >= params.n_rays {
+            return;
+        }
+        let i = idx as usize;
+ 
+        let ray_q = &*params.ray_q.as_raw();
+        let work = ray_q.storage.get(i);
+ 
+        let ray = Ray::new(
+            work.ray_o,
+            work.ray_d,
+            Some(work.ray_time),
+            work.ray_medium,
+        );
+ 
+        let pi = work.pixel_index as usize;
+        let pss = &*params.pixel_sample_state.as_raw();
+ 
+        // Read persistent path state
+        let beta = pss.beta.get(pi);
+        let r_u = pss.r_u.get(pi);
+        let r_l = pss.r_l.get(pi);
+        let lambda = pss.lambda.get(pi);
+        let depth = pss.depth.get(pi);
+        let specular_bounce = pss.specular_bounce.get(pi) != 0;
+        let prev_intr_ctx = pss.prev_intr_ctx.get(pi);
+        let eta_scale = pss.eta_scale.get(pi);
+        let any_non_specular = pss.any_non_specular_bounces.get(pi) != 0;
+ 
+        // BVH traversal — mirrors BVHAggregate::intersect exactly
+        let bvh = &*params.bvh.as_raw();
+        if bvh.nodes.is_empty() {
+            // No geometry — ray escapes
+            push_escaped(params, &work, &lambda, &beta, &r_u, &r_l, depth, specular_bounce, &prev_intr_ctx);
+            return;
+        }
+ 
+        let nodes_ptr = bvh.nodes.as_ptr();
+        let prims_ptr = bvh.primitives.as_ptr();
+        let mut best_si = None;
+        let mut hit_t: Float = Float::INFINITY;
+ 
+        let inv_dir = Vector3f::new(
+            1.0 / ray.d.x(),
+            1.0 / ray.d.y(),
+            1.0 / ray.d.z(),
+        );
+        let dir_is_neg = [
+            if inv_dir.x() < 0.0 { 1u8 } else { 0 },
+            if inv_dir.y() < 0.0 { 1u8 } else { 0 },
+            if inv_dir.z() < 0.0 { 1u8 } else { 0 },
+        ];
+ 
+        let mut to_visit_offset: u32 = 0;
+        let mut current_node_index: usize = 0;
+        // GPU stack — 64 entries matches CPU, fits in registers/local memory
+        let mut nodes_to_visit = [0usize; 64];
+ 
+        loop {
+            let node = &*nodes_ptr.add(current_node_index);
+ 
+            if node.bounds.intersect_p(ray.o, hit_t, inv_dir, &dir_is_neg).is_some() {
+                if node.n_primitives > 0 {
+                    // Leaf node — test primitives
+                    let mut j = 0u16;
+                    while j < node.n_primitives {
+                        let prim_idx = node.primitives_offset + j as usize;
+                        let prim = &*prims_ptr.add(prim_idx);
+ 
+                        if let Some(si) = prim.intersect(&ray, Some(hit_t)) {
+                            hit_t = si.t_hit();
+                            best_si = Some(si);
+                        }
+                        j += 1;
+                    }
+ 
+                    if to_visit_offset == 0 {
+                        break;
+                    }
+                    to_visit_offset -= 1;
+                    current_node_index = nodes_to_visit[to_visit_offset as usize];
+                } else {
+                    // Interior node — push far child, visit near child
+                    if dir_is_neg[node.axis as usize] == 1 {
+                        nodes_to_visit[to_visit_offset as usize] = current_node_index + 1;
+                        to_visit_offset += 1;
+                        current_node_index = node.primitives_offset;
+                    } else {
+                        nodes_to_visit[to_visit_offset as usize] = node.primitives_offset;
+                        to_visit_offset += 1;
+                        current_node_index += 1;
+                    }
+                }
+            } else {
+                if to_visit_offset == 0 {
+                    break;
+                }
+                to_visit_offset -= 1;
+                current_node_index = nodes_to_visit[to_visit_offset as usize];
+            }
+        }
+ 
+        // Sort result into output queues
+        let Some(si) = best_si else {
+            push_escaped(params, &work, &lambda, &beta, &r_u, &r_l, depth, specular_bounce, &prev_intr_ctx);
+            return;
+        };
+ 
+        let intr = &si.intr;
+ 
+        // Null material — medium interface, re-queue ray
+        if intr.material.is_null() {
+            let next_q = &*params.next_ray_q.as_raw();
+            next_q.push(RayWorkItem {
+                ray_o: intr.p(),
+                ray_d: work.ray_d,
+                ray_time: work.ray_time,
+                ray_medium: work.ray_medium,
+                has_differentials: work.has_differentials,
+                differential: work.differential,
+                pixel_index: work.pixel_index,
+            });
+            return;
+        }
+ 
+        // Area light hit
+        if !intr.area_light.is_null() {
+            let q = &*params.hit_area_light_q.as_raw();
+            q.push(HitAreaLightWorkItem {
+                area_light: intr.area_light,
+                p: intr.p(),
+                n: intr.n(),
+                uv: intr.common.uv,
+                wo: -work.ray_d,
+                lambda,
+                pixel_index: work.pixel_index,
+                beta,
+                r_u,
+                r_l,
+                depth,
+                specular_bounce,
+                prev_intr_ctx,
+            });
+        }
+ 
+        // Material evaluation: push to appropriate queue
+        // For now, push everything to universal eval queue.
+        // Basic vs universal split requires checking can_evaluate_textures
+        // on the material, which we can refine later.
+        let q = &*params.universal_eval_mtl_q.as_raw();
+        q.push(MaterialEvalWorkItem {
+            p: intr.p(),
+            n: intr.n(),
+            ns: intr.shading.n,
+            dpdu: intr.shading.dpdu,
+            dpdv: intr.shading.dpdv,
+            uv: intr.common.uv,
+            wo: -work.ray_d,
+            time: work.ray_time,
+            face_index: intr.face_index,
+            material: intr.material,
+            area_light: intr.area_light,
+            medium_interface: intr.common.medium_interface,
+            pixel_index: work.pixel_index,
+            lambda,
+            beta,
+            r_u,
+            any_non_specular_bounces: any_non_specular,
+            depth,
+            eta_scale,
+        });
+    }
+ 
+    /// Shadow ray kernel — one thread per shadow ray, binary occlusion test.
+    #[kernel]
+    pub unsafe fn intersect_shadow(params: &IntersectShadowParams) {
+        let idx = thread::index_1d();
+        if idx >= params.n_rays {
+            return;
+        }
+        let i = idx as usize;
+ 
+        let shadow_q = &*params.shadow_ray_q.as_raw();
+        let work = shadow_q.storage.get(i);
+ 
+        let ray = Ray::new(
+            work.ray_o,
+            work.ray_d,
+            Some(work.ray_time),
+            Ptr::null(),
+        );
+ 
+        // BVH any-hit traversal
+        let bvh = &*params.bvh.as_raw();
+        let occluded = bvh_intersect_p(bvh, &ray, work.t_max);
+ 
+        // If NOT occluded, add direct lighting contribution
+        if !occluded {
+            let pss = &*params.pixel_sample_state.as_raw();
+            let pi = work.pixel_index as usize;
+            // Atomic add to each spectral channel
+            let mut l = pss.l.get(pi);
+            l += work.l_d;
+            pss.l.set(pi, l);
+            // NOTE: This set is not atomic per-channel. For correctness
+            // when multiple shadow rays hit the same pixel, we'd need
+            // per-channel AtomicFloat. For now this works because each
+            // pixel has at most one shadow ray in flight per depth.
+        }
+    }
+ 
+    /// Launch parameters for shadow ray kernel.
+    #[repr(C)]
+    pub struct IntersectShadowParams {
+        pub bvh: Ptr<BVHAggregate>,
+        pub shadow_ray_q: Ptr<ShadowRayQueue>,
+        pub pixel_sample_state: Ptr<PixelSampleState>,
+        pub n_rays: u32,
+    }
+ 
+    // -- Helper functions --
+ 
+    unsafe fn push_escaped(
+        params: &IntersectClosestParams,
+        work: &RayWorkItem,
+        lambda: &SampledWavelengths,
+        beta: &SampledSpectrum,
+        r_u: &SampledSpectrum,
+        r_l: &SampledSpectrum,
+        depth: u32,
+        specular_bounce: bool,
+        prev_intr_ctx: &LightSampleContext,
+    ) {
+        let q = &*params.escaped_ray_q.as_raw();
+        q.push(EscapedRayWorkItem {
+            ray_o: work.ray_o,
+            ray_d: work.ray_d,
+            lambda: *lambda,
+            pixel_index: work.pixel_index,
+            beta: *beta,
+            r_u: *r_u,
+            r_l: *r_l,
+            depth,
+            specular_bounce,
+            prev_intr_ctx: *prev_intr_ctx,
+        });
+    }
+ 
+    /// BVH any-hit traversal for shadow rays — returns true if occluded.
+    unsafe fn bvh_intersect_p(bvh: &BVHAggregate, ray: &Ray, t_max: Float) -> bool {
+        if bvh.nodes.is_empty() {
+            return false;
+        }
+ 
+        let nodes_ptr = bvh.nodes.as_ptr();
+        let prims_ptr = bvh.primitives.as_ptr();
+ 
+        let inv_dir = Vector3f::new(
+            1.0 / ray.d.x(),
+            1.0 / ray.d.y(),
+            1.0 / ray.d.z(),
+        );
+        let dir_is_neg = [
+            if inv_dir.x() < 0.0 { 1u8 } else { 0 },
+            if inv_dir.y() < 0.0 { 1u8 } else { 0 },
+            if inv_dir.z() < 0.0 { 1u8 } else { 0 },
+        ];
+ 
+        let mut to_visit_offset: u32 = 0;
+        let mut current_node_index: usize = 0;
+        let mut nodes_to_visit = [0usize; 64];
+ 
+        loop {
+            let node = &*nodes_ptr.add(current_node_index);
+ 
+            if node.bounds.intersect_p(ray.o, t_max, inv_dir, &dir_is_neg).is_some() {
+                if node.n_primitives > 0 {
+                    let mut j = 0u16;
+                    while j < node.n_primitives {
+                        let prim_idx = node.primitives_offset + j as usize;
+                        let prim = &*prims_ptr.add(prim_idx);
+                        if prim.intersect_p(ray, Some(t_max)) {
+                            return true;
+                        }
+                        j += 1;
+                    }
+                    if to_visit_offset == 0 {
+                        break;
+                    }
+                    to_visit_offset -= 1;
+                    current_node_index = nodes_to_visit[to_visit_offset as usize];
+                } else {
+                    if dir_is_neg[node.axis as usize] == 1 {
+                        nodes_to_visit[to_visit_offset as usize] = current_node_index + 1;
+                        to_visit_offset += 1;
+                        current_node_index = node.primitives_offset;
+                    } else {
+                        nodes_to_visit[to_visit_offset as usize] = node.primitives_offset;
+                        to_visit_offset += 1;
+                        current_node_index += 1;
+                    }
+                }
+            } else {
+                if to_visit_offset == 0 {
+                    break;
+                }
+                to_visit_offset -= 1;
+                current_node_index = nodes_to_visit[to_visit_offset as usize];
+            }
+        }
+        false
+    }
+}
+
+
+#[cfg(feature = "cuda")]
+pub mod host {
+    use crate::core::aggregates::BVHAggregate;
+    use crate::core::geometry::Bounds3f;
+    use crate::core::primitive::PrimitiveTrait;
+    use crate::wavefront::aggregate::WavefrontAggregate;
+    use crate::wavefront::work_items::*;
+    use crate::{Ptr, Float};
+ 
+    use cust::prelude::*;
+    use cust::launch;
+ 
+    /// CUDA aggregate — holds the BVH and the compiled kernel module.
+    pub struct CudaAggregate {
+        pub bvh: Ptr<BVHAggregate>,
+        pub module: Module,
+        pub stream: Stream,
+    }
+ 
+    impl CudaAggregate {
+        pub fn new(bvh: Ptr<BVHAggregate>, ptx_data: &str) -> cust::error::CudaResult<Self> {
+            // Initialize CUDA context (assumes cust::init() already called)
+            let module = Module::from_ptx(ptx_data, &[])?;
+            let stream = Stream::new(StreamFlags::NON_BLOCKING, None)?;
+            Ok(Self { bvh, module, stream })
+        }
+ 
+        fn launch_intersect_closest(
+            &self,
+            n_rays: u32,
+            ray_q: &RayQueue,
+            escaped_ray_q: &EscapedRayQueue,
+            hit_area_light_q: &HitAreaLightQueue,
+            basic_eval_mtl_q: &MaterialEvalQueue,
+            universal_eval_mtl_q: &MaterialEvalQueue,
+            next_ray_q: &RayQueue,
+            pixel_sample_state: &PixelSampleState,
+        ) -> cust::error::CudaResult<()> {
+            if n_rays == 0 {
+                return Ok(());
+            }
+ 
+            let func = self.module.get_function("intersect_closest")?;
+ 
+            // Build launch params in unified memory
+            let params = super::device::IntersectClosestParams {
+                bvh: self.bvh,
+                ray_q: Ptr::from(ray_q),
+                escaped_ray_q: Ptr::from(escaped_ray_q),
+                hit_area_light_q: Ptr::from(hit_area_light_q),
+                basic_eval_mtl_q: Ptr::from(basic_eval_mtl_q),
+                universal_eval_mtl_q: Ptr::from(universal_eval_mtl_q),
+                next_ray_q: Ptr::from(next_ray_q),
+                pixel_sample_state: Ptr::from(pixel_sample_state),
+                n_rays,
+            };
+ 
+            let block_size = 256u32;
+            let grid_size = (n_rays + block_size - 1) / block_size;
+ 
+            unsafe {
+                launch!(
+                    func<<<grid_size, block_size, 0, self.stream>>>(
+                        &params
+                    )
+                )?;
+            }
+            self.stream.synchronize()?;
+            Ok(())
+        }
+ 
+        fn launch_intersect_shadow(
+            &self,
+            n_rays: u32,
+            shadow_ray_q: &ShadowRayQueue,
+            pixel_sample_state: &PixelSampleState,
+        ) -> cust::error::CudaResult<()> {
+            if n_rays == 0 {
+                return Ok(());
+            }
+ 
+            let func = self.module.get_function("intersect_shadow")?;
+ 
+            let params = super::device::IntersectShadowParams {
+                bvh: self.bvh,
+                shadow_ray_q: Ptr::from(shadow_ray_q),
+                pixel_sample_state: Ptr::from(pixel_sample_state),
+                n_rays,
+            };
+ 
+            let block_size = 256u32;
+            let grid_size = (n_rays + block_size - 1) / block_size;
+ 
+            unsafe {
+                launch!(
+                    func<<<grid_size, block_size, 0, self.stream>>>(
+                        &params
+                    )
+                )?;
+            }
+            self.stream.synchronize()?;
+            Ok(())
+        }
+    }
+ 
+    impl WavefrontAggregate for CudaAggregate {
+        fn bounds(&self) -> Bounds3f {
+            self.bvh.get().map(|b| b.bounds()).unwrap_or_default()
+        }
+ 
+        fn intersect_closest(
+            &self,
+            max_rays: usize,
+            ray_q: &RayQueue,
+            escaped_ray_q: &EscapedRayQueue,
+            hit_area_light_q: &HitAreaLightQueue,
+            basic_eval_mtl_q: &MaterialEvalQueue,
+            universal_eval_mtl_q: &MaterialEvalQueue,
+            next_ray_q: &RayQueue,
+            pixel_sample_state: &PixelSampleState,
+        ) {
+            let n_rays = ray_q.size().min(max_rays as u32);
+            self.launch_intersect_closest(
+                n_rays,
+                ray_q,
+                escaped_ray_q,
+                hit_area_light_q,
+                basic_eval_mtl_q,
+                universal_eval_mtl_q,
+                next_ray_q,
+                pixel_sample_state,
+            )
+            .expect("CUDA intersect_closest kernel launch failed");
+        }
+ 
+        fn intersect_shadow(
+            &self,
+            max_rays: usize,
+            shadow_ray_q: &ShadowRayQueue,
+            pixel_sample_state: &PixelSampleState,
+        ) {
+            let n_rays = shadow_ray_q.size().min(max_rays as u32);
+            self.launch_intersect_shadow(n_rays, shadow_ray_q, pixel_sample_state)
+                .expect("CUDA intersect_shadow kernel launch failed");
+        }
+ 
+        fn intersect_shadow_tr(
+            &self,
+            max_rays: usize,
+            shadow_ray_q: &ShadowRayQueue,
+            pixel_sample_state: &PixelSampleState,
+        ) {
+            // Without participating media, shadow_tr is the same as shadow
+            self.intersect_shadow(max_rays, shadow_ray_q, pixel_sample_state);
+        }
+    }
+}
+

shared/src/core/geometry/ray.rs

@@ -1,7 +1,7 @@
 use super::{Normal3f, Point3f, Point3fi, Vector3f, VectorLike};
 use crate::core::medium::Medium;
 use crate::utils::math::{next_float_down, next_float_up};
-use crate::{gvec_with_capacity, Float, GVec, Ptr, SOA};
+use crate::{gvec_with_capacity, Float, GVec, Ptr};
 
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]
@@ -123,56 +123,3 @@ pub struct RayDifferential {
     pub rx_direction: Vector3f,
     pub ry_direction: Vector3f,
 }
-
-#[derive(Clone)]
-pub struct RaySoA {
-    pub o: GVec<Point3f>,
-    pub d: GVec<Vector3f>,
-    pub time: GVec<Float>,
-    pub medium: GVec<Ptr<Medium>>,
-    pub has_differentials: GVec<bool>,
-    pub differential: GVec<RayDifferential>,
-}
-
-impl SoA for RaySoA {
-    type Item = Ray;
-
-    fn with_capacity(n: usize) -> Self {
-        Self {
-            o: gvec_with_capacity(n),
-            d: gvec_with_capacity(n),
-            time: gvec_with_capacity(n),
-            medium: gvec_with_capacity(n),
-            has_differentials: gvec_with_capacity(n),
-            differential: gvec_with_capacity(n),
-        }
-    }
-
-    fn len(&self) -> usize {
-        self.o.len()
-    }
-
-    unsafe fn get_unchecked(&self, i: usize) -> Ray {
-        Ray {
-            o: *self.o.get_unchecked(i),
-            d: *self.d.get_unchecked(i),
-            time: *self.time.get_unchecked(i),
-            medium: *self.medium.get_unchecked(i),
-            has_differentials: *self.has_differentials.get_unchecked(i),
-            differential: *self.differential.get_unchecked(i),
-        }
-    }
-
-    unsafe fn set_unchecked(&mut self, i: usize, v: Ray) {
-        *self.o.get_unchecked_mut(i) = v.o;
-        *self.d.get_unchecked_mut(i) = v.d;
-        *self.time.get_unchecked_mut(i) = v.time;
-        *self.medium.get_unchecked_mut(i) = v.medium;
-        *self.has_differentials.get_unchecked_mut(i) = v.has_differentials;
-        *self.differential.get_unchecked_mut(i) = v.differential;
-    }
-}
-
-impl SoAElement for Ray {
-    type SoA = RaySoA;
-}

shared/src/core/material.rs

@@ -178,6 +178,7 @@ pub trait MaterialTrait {
     fn has_subsurface_scattering(&self) -> bool;
 }
 
+#[repr(C)]
 #[derive(Clone, Copy, Debug)]
 #[enum_dispatch(MaterialTrait)]
 pub enum Material {

shared/src/core/primitive.rs

@@ -68,9 +68,6 @@ impl PrimitiveTrait for GeometricPrimitive {
             }
         }
 
-        if r.medium.is_null() {
-            return None;
-        }
         si.set_intersection_properties(
             self.material,
             self.area_light,

shared/src/core/sampler.rs

@@ -193,7 +193,7 @@ impl SamplerTrait for HaltonSampler {
     }
 
     fn get1d(&mut self) -> Float {
-        if self.dim > PRIME_TABLE_SIZE as u32 {
+        if self.dim >= PRIME_TABLE_SIZE as u32 {
             self.dim = 2;
         }
         self.sample_dimension(self.dim)

shared/src/core/texture.rs

@@ -1,6 +1,6 @@
 use crate::core::color::ColorEncoding;
 use crate::core::geometry::{
-    Normal3f, Point2f, Point3f, Vector2f, Vector3f, VectorLike, spherical_phi, spherical_theta,
+    spherical_phi, spherical_theta, Normal3f, Point2f, Point3f, Vector2f, Vector3f, VectorLike,
 };
 use crate::core::image::WrapMode;
 use crate::core::interaction::{Interaction, InteractionTrait, SurfaceInteraction};
@@ -9,9 +9,9 @@ use crate::spectra::{
     SampledWavelengths,
 };
 
+use crate::utils::math::square;
 use crate::utils::Ptr;
 use crate::utils::Transform;
-use crate::utils::math::square;
 use crate::{Float, INV_2_PI, INV_PI, PI};
 use enum_dispatch::enum_dispatch;
 use num_traits::Float as NumFloat;
@@ -393,7 +393,6 @@ pub enum GPUSpectrumTexture {
     Dots(SpectrumDotsTexture),
     Scaled(GPUSpectrumScaledTexture),
     Image(GPUSpectrumImageTexture),
-    Ptex(GPUSpectrumPtexTexture),
     Mix(GPUSpectrumMixTexture),
 }
 
@@ -411,7 +410,6 @@ impl GPUSpectrumTexture {
             GPUSpectrumTexture::DirectionMix(t) => t.evaluate(ctx, lambda),
             GPUSpectrumTexture::Dots(t) => t.evaluate(ctx, lambda),
             GPUSpectrumTexture::Scaled(t) => t.evaluate(ctx, lambda),
-            GPUSpectrumTexture::Ptex(t) => t.evaluate(ctx, lambda),
             GPUSpectrumTexture::Image(t) => t.evaluate(ctx, lambda),
             GPUSpectrumTexture::Mix(t) => t.evaluate(ctx, lambda),
         }
@@ -460,3 +458,58 @@ impl TextureEvaluator for UniversalTextureEvaluator {
         true
     }
 }
+
+#[repr(C)]
+#[derive(Copy, Clone, Default)]
+pub struct BasicTextureEvaluator;
+
+impl TextureEvaluator for BasicTextureEvaluator {
+    fn evaluate_float(&self, tex: &GPUFloatTexture, ctx: &TextureEvalContext) -> Float {
+        match tex {
+            GPUFloatTexture::Constant(t) => t.evaluate(ctx),
+            GPUFloatTexture::Image(t) => t.evaluate(ctx),
+            _ => 0.0,
+        }
+    }
+
+    fn evaluate_spectrum(
+        &self,
+        tex: &GPUSpectrumTexture,
+        ctx: &TextureEvalContext,
+        lambda: &SampledWavelengths,
+    ) -> SampledSpectrum {
+        match tex {
+            GPUSpectrumTexture::Constant(t) => t.evaluate(ctx, lambda),
+            GPUSpectrumTexture::Image(t) => t.evaluate(ctx, lambda),
+            _ => SampledSpectrum::new(0.0),
+        }
+    }
+
+    fn can_evaluate(
+        &self,
+        ftex: &[Ptr<GPUFloatTexture>],
+        stex: &[Ptr<GPUSpectrumTexture>],
+    ) -> bool {
+        for t in ftex {
+            if t.is_null() {
+                continue;
+            }
+            match t.get().unwrap() {
+                GPUFloatTexture::Constant(_)
+                | GPUFloatTexture::Image(_) => {}
+                _ => return false,
+            }
+        }
+        for t in stex {
+            if t.is_null() {
+                continue;
+            }
+            match t.get().unwrap() {
+                GPUSpectrumTexture::Constant(_)
+                | GPUSpectrumTexture::Image(_) => {}
+                _ => return false,
+            }
+        }
+        true
+    }
+}

shared/src/lib.rs

@@ -20,5 +20,6 @@ pub mod wavefront;
 
 pub use core::pbrt::*;
 pub use utils::alloc::{gbox, gvec, gvec_from_slice, gvec_with_capacity, leak, GBox, GVec};
-pub use utils::{Array2D, PBRTOptions, Ptr, Transform, SOA};
-pub use wavefront::{WavefrontAggregate, WorkQueue};
+pub use utils::{Array2D, BasicPBRTOptions, PBRTOptions, Ptr, Transform};
+pub use utils::soa::WorkQueue;
+pub use wavefront::{WavefrontAggregate};

shared/src/materials/coated.rs

@@ -334,7 +334,7 @@ impl MaterialTrait for CoatedConductorMaterial {
     }
 
     fn get_normal_map(&self) -> Option<&Image> {
-        Some(&*self.normal_map)
+        self.normal_map.get()
     }
 
     fn get_displacement(&self) -> Ptr<GPUFloatTexture> {

shared/src/shapes/triangle.rs

@@ -42,8 +42,9 @@ impl TriangleShape {
     pub const MAX_SPHERICAL_SAMPLE_AREA: Float = 6.22;
 
     fn mesh(&self) -> &TriangleMesh {
-        &*self.mesh
+        self.mesh.get().unwrap()
     }
+
     fn get_vertex_indices(&self) -> [usize; 3] {
         let mesh = self.mesh();
         let base = (self.tri_index as usize) * 3;

shared/src/utils/atomic.rs

@@ -0,0 +1,263 @@
+use crate::Float;
+
+pub const SCOPE_DEVICE: u32 = 1;
+
+#[allow(dead_code)]
+pub const SCOPE_WORKGROUP: u32 = 2;
+
+pub const SEMANTICS_RELAXED: u32 = 0x0;
+
+#[allow(dead_code)]
+pub const SEMANTICS_ACQUIRE_RELEASE: u32 = 0x8;
+
+#[repr(C)]
+#[derive(Debug)]
+pub struct AtomicU32 {
+    value: u32,
+}
+
+impl Default for AtomicU32 {
+    fn default() -> Self {
+        Self::new(0)
+    }
+}
+
+impl Clone for AtomicU32 {
+    fn clone(&self) -> Self {
+        Self::new(self.load())
+    }
+}
+
+impl AtomicU32 {
+    pub fn new(val: u32) -> Self {
+        Self { value: val }
+    }
+
+    #[cfg(not(any(target_arch = "spirv", feature = "cuda")))]
+    #[inline(always)]
+    pub fn load(&self) -> u32 {
+        let atomic = unsafe {
+            &*(core::ptr::addr_of!(self.value) as *const core::sync::atomic::AtomicU32)
+        };
+        atomic.load(core::sync::atomic::Ordering::Relaxed)
+    }
+
+    #[cfg(not(any(target_arch = "spirv", feature = "cuda")))]
+    #[inline(always)]
+    pub fn store(&self, val: u32) {
+        let atomic = unsafe {
+            &*(core::ptr::addr_of!(self.value) as *const core::sync::atomic::AtomicU32)
+        };
+        atomic.store(val, core::sync::atomic::Ordering::Relaxed);
+    }
+
+    #[cfg(not(any(target_arch = "spirv", feature = "cuda")))]
+    #[inline(always)]
+    pub fn fetch_add(&self, val: u32) -> u32 {
+        let atomic = unsafe {
+            &*(core::ptr::addr_of!(self.value) as *const core::sync::atomic::AtomicU32)
+        };
+        atomic.fetch_add(val, core::sync::atomic::Ordering::Relaxed)
+    }
+
+    #[cfg(not(any(target_arch = "spirv", feature = "cuda")))]
+    #[inline(always)]
+    pub fn compare_exchange(&self, expected: u32, desired: u32) -> Result<u32, u32> {
+        let atomic = unsafe {
+            &*(core::ptr::addr_of!(self.value) as *const core::sync::atomic::AtomicU32)
+        };
+        atomic.compare_exchange_weak(
+            expected,
+            desired,
+            core::sync::atomic::Ordering::Relaxed,
+            core::sync::atomic::Ordering::Relaxed,
+        )
+    }
+
+    #[cfg(target_arch = "spirv")]
+    #[inline(always)]
+    pub fn load(&self) -> u32 {
+        unsafe {
+            spirv_std::arch::atomic_load::<u32, SCOPE_DEVICE, SEMANTICS_RELAXED>(
+                &self.value,
+            )
+        }
+    }
+
+    #[cfg(target_arch = "spirv")]
+    #[inline(always)]
+    pub fn store(&self, val: u32) {
+        unsafe {
+            spirv_std::arch::atomic_store::<u32, SCOPE_DEVICE, SEMANTICS_RELAXED>(
+                &mut *core::ptr::addr_of!(self.value).cast_mut(),
+                val,
+            );
+        }
+    }
+
+    #[cfg(target_arch = "spirv")]
+    #[inline(always)]
+    pub fn fetch_add(&self, val: u32) -> u32 {
+        unsafe {
+            spirv_std::arch::atomic_i_add::<u32, SCOPE_DEVICE, SEMANTICS_RELAXED>(
+                &mut *core::ptr::addr_of!(self.value).cast_mut(),
+                val,
+            )
+        }
+    }
+
+    #[cfg(target_arch = "spirv")]
+    #[inline(always)]
+    pub fn compare_exchange(&self, expected: u32, desired: u32) -> Result<u32, u32> {
+        let old = unsafe {
+            spirv_std::arch::atomic_compare_exchange::<
+                u32,
+                SCOPE_DEVICE,
+                SEMANTICS_RELAXED,
+                SEMANTICS_RELAXED,
+            >(
+                &mut *core::ptr::addr_of!(self.value).cast_mut(),
+                desired,
+                expected,
+            )
+        };
+        if old == expected {
+            Ok(old)
+        } else {
+            Err(old)
+        }
+    }
+
+    // -- CUDA backend --
+    #[cfg(feature = "cuda")]
+    #[inline(always)]
+    pub fn load(&self) -> u32 {
+        // CUDA volatile read for atomicity on the same SM
+        unsafe { core::ptr::read_volatile(&self.value) }
+    }
+
+    #[cfg(feature = "cuda")]
+    #[inline(always)]
+    pub fn store(&self, val: u32) {
+        unsafe {
+            core::ptr::write_volatile(
+                core::ptr::addr_of!(self.value).cast_mut(),
+                val,
+            );
+        }
+    }
+
+    #[cfg(feature = "cuda")]
+    #[inline(always)]
+    pub fn fetch_add(&self, val: u32) -> u32 {
+        let ptr = core::ptr::addr_of!(self.value).cast_mut();
+        let mut old: u32;
+        unsafe {
+            core::arch::asm!(
+                "atom.add.u32 {old}, [{ptr}], {val};",
+                old = out(reg32) old,
+                ptr = in(reg64) ptr,
+                val = in(reg32) val,
+            );
+        }
+        old
+    }
+
+    #[cfg(feature = "cuda")]
+    #[inline(always)]
+    pub fn compare_exchange(&self, expected: u32, desired: u32) -> Result<u32, u32> {
+        let ptr = core::ptr::addr_of!(self.value).cast_mut();
+        let mut old: u32;
+        unsafe {
+            core::arch::asm!(
+                "atom.cas.b32 {old}, [{ptr}], {expected}, {desired};",
+                old = out(reg32) old,
+                ptr = in(reg64) ptr,
+                expected = in(reg32) expected,
+                desired = in(reg32) desired,
+            );
+        }
+        if old == expected {
+            Ok(old)
+        } else {
+            Err(old)
+        }
+    }
+}
+
+#[repr(C)]
+#[derive(Debug)]
+pub struct AtomicF32 {
+    bits: AtomicU32,
+}
+
+impl Default for AtomicF32 {
+    fn default() -> Self {
+        Self::new(0.0)
+    }
+}
+
+impl Clone for AtomicF32 {
+    fn clone(&self) -> Self {
+        Self::new(self.get())
+    }
+}
+
+impl AtomicF32 {
+    pub fn new(val: Float) -> Self {
+        Self {
+            bits: AtomicU32::new(val.to_bits()),
+        }
+    }
+
+    pub fn get(&self) -> Float {
+        Float::from_bits(self.bits.load())
+    }
+
+    pub fn set(&self, val: Float) {
+        self.bits.store(val.to_bits());
+    }
+
+    #[cfg(not(any(target_arch = "spirv", feature = "cuda")))]
+    #[inline(always)]
+    pub fn add(&self, val: Float) {
+        let mut current_bits = self.bits.load();
+        loop {
+            let current_val = Float::from_bits(current_bits);
+            let new_val = current_val + val;
+            let new_bits = new_val.to_bits();
+            match self.bits.compare_exchange(current_bits, new_bits) {
+                Ok(_) => break,
+                Err(x) => current_bits = x,
+            }
+        }
+    }
+
+    #[cfg(target_arch = "spirv")]
+    #[inline(always)]
+    pub fn add(&self, val: Float) {
+        unsafe {
+            let float_ptr = core::ptr::addr_of!(self.bits.value) as *mut Float;
+            spirv_std::arch::atomic_f_add::<Float, SCOPE_DEVICE, SEMANTICS_RELAXED>(
+                &mut *float_ptr,
+                val,
+            );
+        }
+    }
+
+    #[cfg(feature = "cuda")]
+    #[inline(always)]
+    pub fn add(&self, val: Float) {
+        let ptr = core::ptr::addr_of!(self.bits.value) as *mut Float;
+        unsafe {
+            core::arch::asm!(
+                "atom.add.f32 {old}, [{ptr}], {val};",
+                old = out(reg32) _,
+                ptr = in(reg64) ptr,
+                val = in(reg32) val.to_bits(),
+            );
+        }
+    }
+}
+
+pub type AtomicFloat = AtomicF32;

shared/src/utils/math.rs

@@ -544,7 +544,7 @@ pub fn next_float_up(v: Float) -> Float {
     if v >= 0.0 {
         ui = ui.wrapping_add(1);
     } else {
-        ui = ui.wrapping_sub(1);
+        ui.wrapping_sub(1);
     }
     bits_to_float(ui)
 }

shared/src/utils/mod.rs

@@ -1,4 +1,5 @@
 pub mod alloc;
+pub mod atomic;
 pub mod complex;
 pub mod containers;
 pub mod hash;
@@ -10,18 +11,18 @@ pub mod ptr;
 pub mod quaternion;
 pub mod rng;
 pub mod sampling;
+pub mod soa;
 pub mod sobol;
 pub mod splines;
 pub mod transform;
 
-pub use options::PBRTOptions;
+pub use atomic::{AtomicFloat, AtomicU32};
+pub use containers::Array2D;
+pub use options::{BasicPBRTOptions, PBRTOptions};
 pub use ptr::Ptr;
 pub use transform::{AnimatedTransform, Transform, TransformGeneric};
-pub use containers::Array2D;
-pub use soa::SOA:
 
 use crate::Float;
-use core::sync::atomic::{AtomicU32, Ordering};
 
 #[inline]
 pub fn find_interval<F>(sz: u32, pred: F) -> u32
@@ -62,61 +63,6 @@ where
     i
 }
 
-#[repr(C)]
-#[derive(Debug)]
-pub struct AtomicFloat {
-    bits: AtomicU32,
-}
-
-impl Default for AtomicFloat {
-    fn default() -> Self {
-        Self::new(0.)
-    }
-}
-
-impl Clone for AtomicFloat {
-    fn clone(&self) -> Self {
-        Self::new(self.get())
-    }
-}
-
-impl AtomicFloat {
-    pub fn new(val: f32) -> Self {
-        Self {
-            bits: AtomicU32::new(val.to_bits()),
-        }
-    }
-
-    pub fn get(&self) -> Float {
-        Float::from_bits(self.bits.load(Ordering::Relaxed))
-    }
-
-    pub fn set(&self, val: Float) {
-        self.bits.store(val.to_bits(), Ordering::Relaxed);
-    }
-
-    /// Atomically adds `val` to the current value.
-    /// Compare-And-Swap loop.
-    pub fn add(&self, val: f32) {
-        let mut current_bits = self.bits.load(Ordering::Relaxed);
-        loop {
-            let current_val = f32::from_bits(current_bits);
-            let new_val = current_val + val;
-            let new_bits = new_val.to_bits();
-
-            match self.bits.compare_exchange_weak(
-                current_bits,
-                new_bits,
-                Ordering::Relaxed,
-                Ordering::Relaxed,
-            ) {
-                Ok(_) => break,
-                Err(x) => current_bits = x,
-            }
-        }
-    }
-}
-
 #[inline(always)]
 pub fn gpu_array_from_fn<T, const N: usize>(mut f: impl FnMut(usize) -> T) -> [T; N] {
     unsafe {
@@ -130,4 +76,3 @@ pub fn gpu_array_from_fn<T, const N: usize>(mut f: impl FnMut(usize) -> T) -> [T
         arr.assume_init()
     }
 }
-

shared/src/utils/options.rs

@@ -9,7 +9,7 @@ pub enum RenderingCoordinateSystem {
     World,
 }
 
-#[derive(Debug, Clone)]
+#[derive(Debug, Clone, Copy)]
 pub struct BasicPBRTOptions {
     pub seed: i32,
     pub quiet: bool,
@@ -42,7 +42,7 @@ impl Default for BasicPBRTOptions {
     }
 }
 
-#[derive(Debug, Clone)]
+#[derive(Debug, Clone, Copy)]
 pub struct PBRTOptions {
     pub basic: BasicPBRTOptions,
 

shared/src/utils/soa.rs

@@ -1,15 +1,120 @@
-pub trait SoA: Clone {
+use crate::utils::AtomicU32;
+use crate::{Float, Ptr};
+
+#[repr(C)]
+#[derive(Debug, Clone, Copy)]
+pub struct SoABuffer<T: Copy> {
+    pub ptr: Ptr<T>,
+    pub capacity: u32,
+}
+
+impl<T: Copy> SoABuffer<T> {
+    pub fn null() -> Self {
+        Self {
+            ptr: Ptr::null(),
+            capacity: 0,
+        }
+    }
+
+    #[inline(always)]
+    pub fn get(&self, i: usize) -> T {
+        debug_assert!(i < self.capacity as usize);
+        unsafe { *self.ptr.as_raw().add(i) }
+    }
+
+    #[inline(always)]
+    pub fn set(&self, i: usize, val: T) {
+        debug_assert!(i < self.capacity as usize);
+        unsafe {
+            let p = self.ptr.as_raw() as *mut T;
+            *p.add(i) = val;
+        }
+    }
+}
+
+pub trait SoA {
     type Item: Copy;
-    fn with_capacity(n: usize) -> Self;
-    fn len(&self) -> usize;
-    fn is_empty(&self) -> bool {
-        self.len() == 0
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self;
+
+    /// Read one element by scattering across all buffers.
+    unsafe fn get(&self, i: usize) -> Self::Item;
+
+    /// Write one element by scattering across all buffers.
+    /// Takes &self for GPU concurrent write (disjoint indices).
+    unsafe fn set(&self, i: usize, v: Self::Item);
 }
 
-    unsafe fn get_unchecked(&self, i: usize) -> Self::Item;
-    unsafe fn set_unchecked(&mut self, i: usize, v: Self::Item);
+pub trait SoAAllocator {
+    fn alloc_raw(&self, layout: core::alloc::Layout) -> *mut u8;
 }
 
-pub trait SoAElement: Copy {
-    type SoA: SoA<Item = Self>;
+pub fn alloc_soa_buffer<T: Copy>(n: u32, alloc: &dyn SoAAllocator) -> SoABuffer<T> {
+    if n == 0 {
+        return SoABuffer::null();
+    }
+    let layout = core::alloc::Layout::array::<T>(n as usize).unwrap();
+    let raw = alloc.alloc_raw(layout);
+    SoABuffer {
+        ptr: Ptr::from_raw(raw as *mut T),
+        capacity: n,
+    }
+}
+
+#[repr(C)]
+pub struct WorkQueue<S: SoA> {
+    pub storage: S,
+    pub count: AtomicU32,
+    pub capacity: u32,
+}
+
+impl<S: SoA> WorkQueue<S> {
+    pub fn new(storage: S, capacity: u32) -> Self {
+        Self {
+            storage,
+            count: AtomicU32::new(0),
+            capacity,
+        }
+    }
+
+    /// Number of items currently in the queue.
+    #[inline(always)]
+    pub fn size(&self) -> u32 {
+        self.count.load()
+    }
+
+    /// Reset the queue to empty. Call from host between kernel launches.
+    #[inline(always)]
+    pub fn reset(&self) {
+        self.count.store(0);
+    }
+
+    /// Push an item, returning its index. Returns None if the queue is full.
+    /// The atomic increment guarantees each thread gets a unique slot.
+    #[inline(always)]
+    pub fn push(&self, item: S::Item) -> Option<u32> {
+        let slot = self.count.fetch_add(1);
+        if slot >= self.capacity {
+            // Queue overflow — this shouldn't happen if capacity is
+            // sized correctly. In debug builds we want to know about it.
+            debug_assert!(false, "WorkQueue overflow: {} >= {}", slot, self.capacity);
+            return None;
+        }
+        unsafe {
+            self.storage.set(slot as usize, item);
+        }
+        Some(slot)
+    }
+
+    /// Read an item at the given index. Used by consumer kernels.
+    #[inline(always)]
+    pub unsafe fn get(&self, i: usize) -> S::Item {
+        debug_assert!(
+            (i as u32) < self.size(),
+            "WorkQueue::get out of bounds: {} >= {}",
+            i,
+            self.size()
+        );
+        self.storage.get(i)
+    }
 }

shared/src/wavefront/aggregate.rs

@@ -0,0 +1,33 @@
+use crate::wavefront::workitems::*;
+use crate::core::geometry::Bounds3f;
+
+pub trait WavefrontAggregate {
+    fn bounds(&self) -> Bounds3f;
+
+    fn intersect_closest(
+        &self,
+        max_rays: usize,
+        ray_q: &RayQueue,
+        escaped_ray_q: &EscapedRayQueue,
+        hit_area_light_q: &HitAreaLightQueue,
+        basic_eval_mtl_q: &MaterialEvalQueue,
+        universal_eval_mtl_q: &MaterialEvalQueue,
+        next_ray_q: &RayQueue,
+        pixel_sample_state: &PixelSampleState,
+    );
+
+    fn intersect_shadow(
+        &self,
+        max_rays: usize,
+        shadow_ray_q: &ShadowRayQueue,
+        pixel_sample_state: &PixelSampleState,
+    );
+
+    fn intersect_shadow_tr(
+        &self,
+        max_rays: usize,
+        shadow_ray_q: &ShadowRayQueue,
+        pixel_sample_state: &PixelSampleState,
+    );
+}
+

shared/src/wavefront/integrator.rs

@@ -0,0 +1,36 @@
+use crate::core::camera::Camera;
+use crate::core::film::Film;
+use crate::core::filter::Filter;
+use crate::core::light::Light;
+use crate::core::sampler::Sampler;
+use crate::lights::sampler::LightSampler;
+use crate::wavefront::aggregate::WavefrontAggregate;
+use crate::wavefront::workitems::*;
+use crate::{Float, GVec, Ptr};
+
+pub struct WavefrontPathIntegrator<A: WavefrontAggregate> {
+    pub aggregate: A,
+    pub camera: Camera,
+    pub sampler: Sampler,
+    pub max_depth: u32,
+    pub samples_per_pixel: u32,
+    pub regularize: bool,
+    pub infinite_lights: GVec<Ptr<Light>>,
+    pub max_queue_size: u32,
+    pub scanlines_per_pass: u32,
+    pub light_sampler: LightSampler,
+    pub film: Ptr<Film>,
+    pub filter: Ptr<Filter>,
+    pub ray_queues: [RayQueue; 2],
+    pub shadow_ray_queue: ShadowRayQueue,
+    pub escaped_ray_queue: EscapedRayQueue,
+    pub hit_area_light_queue: HitAreaLightQueue,
+    pub basic_eval_material_queue: MaterialEvalQueue,
+    pub universal_eval_material_queue: MaterialEvalQueue,
+    pub pixel_sample_state: PixelSampleState,
+}
+
+pub trait WavefrontRenderer {
+    fn render(&mut self);
+}
+

shared/src/wavefront/mod.rs

@@ -0,0 +1,8 @@
+pub mod workitems;
+pub mod aggregate;
+pub mod integrator;
+
+pub use workitems::*;
+pub use aggregate::WavefrontAggregate;
+pub use integrator::{WavefrontPathIntegrator, WavefrontRenderer};
+

shared/src/wavefront/workitems.rs

@@ -0,0 +1,507 @@
+use crate::core::bxdf::BxDFFlags;
+use crate::core::geometry::{Normal3f, Point2f, Point3f, Point3fi, Vector3f, RayDifferential};
+use crate::core::light::LightSampleContext;
+use crate::core::light::Light;
+use crate::core::material::Material;
+use crate::core::medium::{Medium, MediumInterface};
+use crate::spectra::{SampledSpectrum, SampledWavelengths};
+use crate::utils::soa::{alloc_soa_buffer, SoA, SoAAllocator, SoABuffer, WorkQueue};
+use crate::{Float, Ptr};
+
+/// Per-path state that persists across all wavefront depth iterations.
+/// Indexed by pixel_index. Allocated once with capacity = max_queue_size.
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct PixelSampleState {
+    pub filter_weight: SoABuffer<Float>,
+    pub p_film: SoABuffer<Point2f>,
+    pub l: SoABuffer<SampledSpectrum>,
+    pub lambda: SoABuffer<SampledWavelengths>,
+    pub r_u: SoABuffer<SampledSpectrum>,
+    pub r_l: SoABuffer<SampledSpectrum>,
+    pub prev_intr_ctx: SoABuffer<LightSampleContext>,
+    pub beta: SoABuffer<SampledSpectrum>,
+    pub depth: SoABuffer<u32>,
+    pub specular_bounce: SoABuffer<u8>,
+    pub any_non_specular_bounces: SoABuffer<u8>,
+    pub eta_scale: SoABuffer<Float>,
+    pub camera_ray_weight: SoABuffer<SampledSpectrum>,
+    pub visible_surface_idx: SoABuffer<u32>,
+}
+
+impl SoA for PixelSampleState {
+    type Item = ();
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            filter_weight: alloc_soa_buffer(n, alloc),
+            p_film: alloc_soa_buffer(n, alloc),
+            l: alloc_soa_buffer(n, alloc),
+            lambda: alloc_soa_buffer(n, alloc),
+            r_u: alloc_soa_buffer(n, alloc),
+            r_l: alloc_soa_buffer(n, alloc),
+            prev_intr_ctx: alloc_soa_buffer(n, alloc),
+            beta: alloc_soa_buffer(n, alloc),
+            depth: alloc_soa_buffer(n, alloc),
+            specular_bounce: alloc_soa_buffer(n, alloc),
+            any_non_specular_bounces: alloc_soa_buffer(n, alloc),
+            eta_scale: alloc_soa_buffer(n, alloc),
+            camera_ray_weight: alloc_soa_buffer(n, alloc),
+            visible_surface_idx: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, _i: usize) -> Self::Item {}
+    unsafe fn set(&self, _i: usize, _v: Self::Item) {}
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct RayWorkItem {
+    pub ray_o: Point3f,
+    pub ray_d: Vector3f,
+    pub ray_time: Float,
+    pub ray_medium: Ptr<Medium>,
+    pub pixel_index: u32,
+    pub has_differentials: bool,
+    pub differential: RayDifferential
+}
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct RayWorkItemSoA {
+    pub ray_o: SoABuffer<Point3f>,
+    pub ray_d: SoABuffer<Vector3f>,
+    pub ray_time: SoABuffer<Float>,
+    pub ray_medium: SoABuffer<Ptr<Medium>>,
+    pub pixel_index: SoABuffer<u32>,
+    pub has_differentials: SoABuffer<bool>,
+    pub differential: SoABuffer<RayDifferential>,
+}
+
+impl SoA for RayWorkItemSoA {
+    type Item = RayWorkItem;
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            ray_o: alloc_soa_buffer(n, alloc),
+            ray_d: alloc_soa_buffer(n, alloc),
+            ray_time: alloc_soa_buffer(n, alloc),
+            ray_medium: alloc_soa_buffer(n, alloc),
+            pixel_index: alloc_soa_buffer(n, alloc),
+            has_differentials: alloc_soa_buffer(n, alloc),
+            differential: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, i: usize) -> RayWorkItem {
+        RayWorkItem {
+            ray_o: self.ray_o.get(i),
+            ray_d: self.ray_d.get(i),
+            ray_time: self.ray_time.get(i),
+            ray_medium: self.ray_medium.get(i),
+            pixel_index: self.pixel_index.get(i),
+            has_differentials: self.has_differentials.get(i),
+            differential: self.differential.get(i),
+        }
+    }
+
+    unsafe fn set(&self, i: usize, v: RayWorkItem) {
+        self.ray_o.set(i, v.ray_o);
+        self.ray_d.set(i, v.ray_d);
+        self.ray_time.set(i, v.ray_time);
+        self.ray_medium.set(i, v.ray_medium);
+        self.pixel_index.set(i, v.pixel_index);
+        self.has_differentials.set(i, v.has_differentials);
+        self.differential.set(i, v.differential);
+    }
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct EscapedRayWorkItem {
+    pub ray_o: Point3f,
+    pub ray_d: Vector3f,
+    pub lambda: SampledWavelengths,
+    pub pixel_index: u32,
+    pub beta: SampledSpectrum,
+    pub r_u: SampledSpectrum,
+    pub r_l: SampledSpectrum,
+    pub depth: u32,
+    pub specular_bounce: bool,
+    pub prev_intr_ctx: LightSampleContext,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct EscapedRayWorkItemSoA {
+    pub ray_o: SoABuffer<Point3f>,
+    pub ray_d: SoABuffer<Vector3f>,
+    pub lambda: SoABuffer<SampledWavelengths>,
+    pub pixel_index: SoABuffer<u32>,
+    pub beta: SoABuffer<SampledSpectrum>,
+    pub r_u: SoABuffer<SampledSpectrum>,
+    pub r_l: SoABuffer<SampledSpectrum>,
+    pub depth: SoABuffer<u32>,
+    pub specular_bounce: SoABuffer<u8>,
+    pub prev_intr_ctx: SoABuffer<LightSampleContext>,
+}
+
+impl SoA for EscapedRayWorkItemSoA {
+    type Item = EscapedRayWorkItem;
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            ray_o: alloc_soa_buffer(n, alloc),
+            ray_d: alloc_soa_buffer(n, alloc),
+            lambda: alloc_soa_buffer(n, alloc),
+            pixel_index: alloc_soa_buffer(n, alloc),
+            beta: alloc_soa_buffer(n, alloc),
+            r_u: alloc_soa_buffer(n, alloc),
+            r_l: alloc_soa_buffer(n, alloc),
+            depth: alloc_soa_buffer(n, alloc),
+            specular_bounce: alloc_soa_buffer(n, alloc),
+            prev_intr_ctx: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, i: usize) -> EscapedRayWorkItem {
+        EscapedRayWorkItem {
+            ray_o: self.ray_o.get(i),
+            ray_d: self.ray_d.get(i),
+            lambda: self.lambda.get(i),
+            pixel_index: self.pixel_index.get(i),
+            beta: self.beta.get(i),
+            r_u: self.r_u.get(i),
+            r_l: self.r_l.get(i),
+            depth: self.depth.get(i),
+            specular_bounce: self.specular_bounce.get(i) != 0,
+            prev_intr_ctx: self.prev_intr_ctx.get(i),
+        }
+    }
+
+    unsafe fn set(&self, i: usize, v: EscapedRayWorkItem) {
+        self.ray_o.set(i, v.ray_o);
+        self.ray_d.set(i, v.ray_d);
+        self.lambda.set(i, v.lambda);
+        self.pixel_index.set(i, v.pixel_index);
+        self.beta.set(i, v.beta);
+        self.r_u.set(i, v.r_u);
+        self.r_l.set(i, v.r_l);
+        self.depth.set(i, v.depth);
+        self.specular_bounce.set(i, v.specular_bounce as u8);
+        self.prev_intr_ctx.set(i, v.prev_intr_ctx);
+    }
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct HitAreaLightWorkItem {
+    pub area_light: Ptr<Light>,
+    pub p: Point3f,
+    pub n: Normal3f,
+    pub uv: Point2f,
+    pub wo: Vector3f,
+    pub lambda: SampledWavelengths,
+    pub pixel_index: u32,
+    pub beta: SampledSpectrum,
+    pub r_u: SampledSpectrum,
+    pub r_l: SampledSpectrum,
+    pub depth: u32,
+    pub specular_bounce: bool,
+    pub prev_intr_ctx: LightSampleContext,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct HitAreaLightWorkItemSoA {
+    pub area_light: SoABuffer<Ptr<Light>>,
+    pub p: SoABuffer<Point3f>,
+    pub n: SoABuffer<Normal3f>,
+    pub uv: SoABuffer<Point2f>,
+    pub wo: SoABuffer<Vector3f>,
+    pub lambda: SoABuffer<SampledWavelengths>,
+    pub pixel_index: SoABuffer<u32>,
+    pub beta: SoABuffer<SampledSpectrum>,
+    pub r_u: SoABuffer<SampledSpectrum>,
+    pub r_l: SoABuffer<SampledSpectrum>,
+    pub depth: SoABuffer<u32>,
+    pub specular_bounce: SoABuffer<u8>,
+    pub prev_intr_ctx: SoABuffer<LightSampleContext>,
+}
+
+impl SoA for HitAreaLightWorkItemSoA {
+    type Item = HitAreaLightWorkItem;
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            area_light: alloc_soa_buffer(n, alloc),
+            p: alloc_soa_buffer(n, alloc),
+            n: alloc_soa_buffer(n, alloc),
+            uv: alloc_soa_buffer(n, alloc),
+            wo: alloc_soa_buffer(n, alloc),
+            lambda: alloc_soa_buffer(n, alloc),
+            pixel_index: alloc_soa_buffer(n, alloc),
+            beta: alloc_soa_buffer(n, alloc),
+            r_u: alloc_soa_buffer(n, alloc),
+            r_l: alloc_soa_buffer(n, alloc),
+            depth: alloc_soa_buffer(n, alloc),
+            specular_bounce: alloc_soa_buffer(n, alloc),
+            prev_intr_ctx: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, i: usize) -> HitAreaLightWorkItem {
+        HitAreaLightWorkItem {
+            area_light: self.area_light.get(i),
+            p: self.p.get(i),
+            n: self.n.get(i),
+            uv: self.uv.get(i),
+            wo: self.wo.get(i),
+            lambda: self.lambda.get(i),
+            pixel_index: self.pixel_index.get(i),
+            beta: self.beta.get(i),
+            r_u: self.r_u.get(i),
+            r_l: self.r_l.get(i),
+            depth: self.depth.get(i),
+            specular_bounce: self.specular_bounce.get(i) != 0,
+            prev_intr_ctx: self.prev_intr_ctx.get(i),
+        }
+    }
+
+    unsafe fn set(&self, i: usize, v: HitAreaLightWorkItem) {
+        self.area_light.set(i, v.area_light);
+        self.p.set(i, v.p);
+        self.n.set(i, v.n);
+        self.uv.set(i, v.uv);
+        self.wo.set(i, v.wo);
+        self.lambda.set(i, v.lambda);
+        self.pixel_index.set(i, v.pixel_index);
+        self.beta.set(i, v.beta);
+        self.r_u.set(i, v.r_u);
+        self.r_l.set(i, v.r_l);
+        self.depth.set(i, v.depth);
+        self.specular_bounce.set(i, v.specular_bounce as u8);
+        self.prev_intr_ctx.set(i, v.prev_intr_ctx);
+    }
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct MaterialEvalWorkItem {
+    // Surface interaction
+    pub p: Point3f,
+    pub n: Normal3f,
+    pub ns: Normal3f,
+    pub dpdu: Vector3f,
+    pub dpdv: Vector3f,
+    pub uv: Point2f,
+    pub wo: Vector3f,
+    pub time: Float,
+    pub face_index: i32,
+
+    // Material
+    pub material: Ptr<Material>,
+    pub area_light: Ptr<Light>,
+
+    // Medium interface
+    pub medium_interface: MediumInterface,
+
+    // Path state
+    pub pixel_index: u32,
+    pub lambda: SampledWavelengths,
+    pub beta: SampledSpectrum,
+    pub r_u: SampledSpectrum,
+
+    // For next-event estimation
+    pub any_non_specular_bounces: bool,
+    pub depth: u32,
+    pub eta_scale: Float,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct MaterialEvalWorkItemSoA {
+    pub p: SoABuffer<Point3f>,
+    pub n: SoABuffer<Normal3f>,
+    pub ns: SoABuffer<Normal3f>,
+    pub dpdu: SoABuffer<Vector3f>,
+    pub dpdv: SoABuffer<Vector3f>,
+    pub uv: SoABuffer<Point2f>,
+    pub wo: SoABuffer<Vector3f>,
+    pub time: SoABuffer<Float>,
+    pub face_index: SoABuffer<i32>,
+    pub material: SoABuffer<Ptr<Material>>,
+    pub area_light: SoABuffer<Ptr<Light>>,
+    pub medium_interface: SoABuffer<MediumInterface>,
+    pub pixel_index: SoABuffer<u32>,
+    pub lambda: SoABuffer<SampledWavelengths>,
+    pub beta: SoABuffer<SampledSpectrum>,
+    pub r_u: SoABuffer<SampledSpectrum>,
+    pub any_non_specular_bounces: SoABuffer<u8>,
+    pub depth: SoABuffer<u32>,
+    pub eta_scale: SoABuffer<Float>,
+}
+
+impl SoA for MaterialEvalWorkItemSoA {
+    type Item = MaterialEvalWorkItem;
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            p: alloc_soa_buffer(n, alloc),
+            n: alloc_soa_buffer(n, alloc),
+            ns: alloc_soa_buffer(n, alloc),
+            dpdu: alloc_soa_buffer(n, alloc),
+            dpdv: alloc_soa_buffer(n, alloc),
+            uv: alloc_soa_buffer(n, alloc),
+            wo: alloc_soa_buffer(n, alloc),
+            time: alloc_soa_buffer(n, alloc),
+            face_index: alloc_soa_buffer(n, alloc),
+            material: alloc_soa_buffer(n, alloc),
+            area_light: alloc_soa_buffer(n, alloc),
+            medium_interface: alloc_soa_buffer(n, alloc),
+            pixel_index: alloc_soa_buffer(n, alloc),
+            lambda: alloc_soa_buffer(n, alloc),
+            beta: alloc_soa_buffer(n, alloc),
+            r_u: alloc_soa_buffer(n, alloc),
+            any_non_specular_bounces: alloc_soa_buffer(n, alloc),
+            depth: alloc_soa_buffer(n, alloc),
+            eta_scale: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, i: usize) -> MaterialEvalWorkItem {
+        MaterialEvalWorkItem {
+            p: self.p.get(i),
+            n: self.n.get(i),
+            ns: self.ns.get(i),
+            dpdu: self.dpdu.get(i),
+            dpdv: self.dpdv.get(i),
+            uv: self.uv.get(i),
+            wo: self.wo.get(i),
+            time: self.time.get(i),
+            face_index: self.face_index.get(i),
+            material: self.material.get(i),
+            area_light: self.area_light.get(i),
+            medium_interface: self.medium_interface.get(i),
+            pixel_index: self.pixel_index.get(i),
+            lambda: self.lambda.get(i),
+            beta: self.beta.get(i),
+            r_u: self.r_u.get(i),
+            any_non_specular_bounces: self.any_non_specular_bounces.get(i) != 0,
+            depth: self.depth.get(i),
+            eta_scale: self.eta_scale.get(i),
+        }
+    }
+
+    unsafe fn set(&self, i: usize, v: MaterialEvalWorkItem) {
+        self.p.set(i, v.p);
+        self.n.set(i, v.n);
+        self.ns.set(i, v.ns);
+        self.dpdu.set(i, v.dpdu);
+        self.dpdv.set(i, v.dpdv);
+        self.uv.set(i, v.uv);
+        self.wo.set(i, v.wo);
+        self.time.set(i, v.time);
+        self.face_index.set(i, v.face_index);
+        self.material.set(i, v.material);
+        self.area_light.set(i, v.area_light);
+        self.medium_interface.set(i, v.medium_interface);
+        self.pixel_index.set(i, v.pixel_index);
+        self.lambda.set(i, v.lambda);
+        self.beta.set(i, v.beta);
+        self.r_u.set(i, v.r_u);
+        self.any_non_specular_bounces
+            .set(i, v.any_non_specular_bounces as u8);
+        self.depth.set(i, v.depth);
+        self.eta_scale.set(i, v.eta_scale);
+    }
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct ShadowRayWorkItem {
+    pub ray_o: Point3f,
+    pub ray_d: Vector3f,
+    pub ray_time: Float,
+    pub t_max: Float,
+    pub lambda: SampledWavelengths,
+    pub l_d: SampledSpectrum,
+    pub pixel_index: u32,
+}
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+pub struct ShadowRayWorkItemSoA {
+    pub ray_o: SoABuffer<Point3f>,
+    pub ray_d: SoABuffer<Vector3f>,
+    pub ray_time: SoABuffer<Float>,
+    pub t_max: SoABuffer<Float>,
+    pub lambda: SoABuffer<SampledWavelengths>,
+    pub l_d: SoABuffer<SampledSpectrum>,
+    pub pixel_index: SoABuffer<u32>,
+}
+
+impl SoA for ShadowRayWorkItemSoA {
+    type Item = ShadowRayWorkItem;
+
+    fn allocate(n: u32, alloc: &dyn SoAAllocator) -> Self {
+        Self {
+            ray_o: alloc_soa_buffer(n, alloc),
+            ray_d: alloc_soa_buffer(n, alloc),
+            ray_time: alloc_soa_buffer(n, alloc),
+            t_max: alloc_soa_buffer(n, alloc),
+            lambda: alloc_soa_buffer(n, alloc),
+            l_d: alloc_soa_buffer(n, alloc),
+            pixel_index: alloc_soa_buffer(n, alloc),
+        }
+    }
+
+    unsafe fn get(&self, i: usize) -> ShadowRayWorkItem {
+        ShadowRayWorkItem {
+            ray_o: self.ray_o.get(i),
+            ray_d: self.ray_d.get(i),
+            ray_time: self.ray_time.get(i),
+            t_max: self.t_max.get(i),
+            lambda: self.lambda.get(i),
+            l_d: self.l_d.get(i),
+            pixel_index: self.pixel_index.get(i),
+        }
+    }
+
+    unsafe fn set(&self, i: usize, v: ShadowRayWorkItem) {
+        self.ray_o.set(i, v.ray_o);
+        self.ray_d.set(i, v.ray_d);
+        self.ray_time.set(i, v.ray_time);
+        self.t_max.set(i, v.t_max);
+        self.lambda.set(i, v.lambda);
+        self.l_d.set(i, v.l_d);
+        self.pixel_index.set(i, v.pixel_index);
+    }
+}
+
+#[repr(C)]
+#[derive(Clone, Copy, Debug)]
+pub struct MediumSampleWorkItem {
+    pub ray_o: Point3f,
+    pub ray_d: Vector3f,
+    pub ray_time: Float,
+    pub t_max: Float,
+    pub lambda: SampledWavelengths,
+    pub beta: SampledSpectrum,
+    pub r_u: SampledSpectrum,
+    pub r_l: SampledSpectrum,
+    pub medium: Ptr<Medium>,
+    pub pixel_index: u32,
+    pub depth: u32,
+    pub specular_bounce: bool,
+    pub any_non_specular_bounces: bool,
+    pub eta_scale: Float,
+    pub prev_intr_ctx: LightSampleContext,
+}
+
+pub type RayQueue = WorkQueue<RayWorkItemSoA>;
+pub type EscapedRayQueue = WorkQueue<EscapedRayWorkItemSoA>;
+pub type HitAreaLightQueue = WorkQueue<HitAreaLightWorkItemSoA>;
+pub type MaterialEvalQueue = WorkQueue<MaterialEvalWorkItemSoA>;
+pub type ShadowRayQueue = WorkQueue<ShadowRayWorkItemSoA>;

src/core/aggregates.rs

@@ -1,9 +1,7 @@
-use crate::Arena;
 use rayon::prelude::*;
 use shared::core::aggregates::{BVHAggregate, LinearBVHNode, SplitMethod};
-use shared::core::geometry::{Bounds3f, Point3f, Ray, Vector3f};
+use shared::core::geometry::{Bounds3f, Point3f};
 use shared::core::primitive::{Primitive, PrimitiveTrait};
-use shared::core::shape::ShapeIntersection;
 use shared::utils::math::encode_morton_3;
 use shared::utils::{find_interval, partition_slice};
 use shared::{gvec, gvec_from_slice, Float};

src/core/camera.rs

@@ -44,7 +44,7 @@ impl CameraBaseParameters {
         }
 
         Ok(CameraBaseParameters {
-            camera_transform: camera_transform.clone(),
+            camera_transform: *camera_transform,
             shutter_open,
             shutter_close,
             film,
@@ -56,7 +56,7 @@ impl CameraBaseParameters {
 pub trait CameraBaseFactory {
     fn create(p: CameraBaseParameters) -> CameraBase {
         CameraBase {
-            camera_transform: p.camera_transform.clone(),
+            camera_transform: p.camera_transform,
             shutter_open: p.shutter_open,
             shutter_close: p.shutter_close,
             film: Ptr::from(p.film.clone().as_ref()),
@@ -144,8 +144,7 @@ impl CameraFactory for Camera {
                 if !sw.is_empty() {
                     if get_options().fullscreen {
                         eprint!("Screenwindow is ignored in fullscreen mode");
-                    } else {
-                        if sw.len() == 4 {
+                    } else if sw.len() == 4 {
                         screen = Bounds2f::from_points(
                             Point2f::new(sw[0], sw[2]),
                             Point2f::new(sw[1], sw[3]),
@@ -157,7 +156,6 @@ impl CameraFactory for Camera {
                         ));
                     }
                 }
-                }
 
                 let fov = params.get_one_float("fov", 90.)?;
                 let camera = PerspectiveCamera::new(base, fov, screen, lens_radius, focal_distance);
@@ -191,8 +189,7 @@ impl CameraFactory for Camera {
                 if !sw.is_empty() {
                     if get_options().fullscreen {
                         eprint!("Screenwindow is ignored in fullscreen mode");
-                    } else {
-                        if sw.len() == 4 {
+                    } else if sw.len() == 4 {
                         screen = Bounds2f::from_points(
                             Point2f::new(sw[0], sw[2]),
                             Point2f::new(sw[1], sw[3]),
@@ -204,7 +201,6 @@ impl CameraFactory for Camera {
                         ));
                     }
                 }
-                }
                 let camera = OrthographicCamera::new(base, screen, lens_radius, focal_distance);
                 arena.alloc(camera);
                 Ok(Camera::Orthographic(camera))
@@ -236,7 +232,7 @@ impl CameraFactory for Camera {
                         PixelFormat::F32,
                         Point2i::new(builtin_res, builtin_res),
                         &["Y"],
-                        SRGB.into(),
+                        SRGB,
                     );
 
                     let res = image.resolution();
@@ -287,7 +283,7 @@ impl CameraFactory for Camera {
                                 PixelFormat::F32,
                                 Point2i::new(builtin_res, builtin_res),
                                 &["Y"],
-                                SRGB.into(),
+                                SRGB,
                             );
                             let res = img.resolution();
                             for y in 0..res.y() {
@@ -309,7 +305,7 @@ impl CameraFactory for Camera {
                                 PixelFormat::F32,
                                 Point2i::new(builtin_res, builtin_res),
                                 &["Y"],
-                                SRGB.into(),
+                                SRGB,
                             );
                             let low = (0.25 * builtin_res as Float) as i32;
                             let high = (0.75 * builtin_res as Float) as i32;
@@ -358,7 +354,7 @@ impl CameraFactory for Camera {
                                         PixelFormat::F32,
                                         im.image.resolution(),
                                         &["Y"],
-                                        SRGB.into(),
+                                        SRGB,
                                     );
                                     let res = mono.resolution();
                                     for y in 0..res.y() {

src/core/color.rs

@@ -1,8 +1,7 @@
 use crate::utils::read_float_file;
 use anyhow::Result;
 use shared::core::color::{Coeffs, RES, RGBToSpectrumTable};
-use shared::{Float, Ptr, gvec_from_slice};
-use std::ops::Deref;
+use shared::{Float, gvec_from_slice};
 use std::path::Path;
 
 pub trait CreateRGBToSpectrumTable {

src/core/film.rs

@@ -96,12 +96,12 @@ impl CreatePixelSensor for PixelSensor {
         };
 
         if sensor_name == "cie1931" {
-            return Ok(Self::new_with_white_balance(
+            Ok(Self::new_with_white_balance(
                 output_colorspace.as_ref(),
                 sensor_illum.as_deref(),
                 imaging_ratio,
                 arena
-            ));
+            ))
         } else {
             let r_opt = get_named_spectrum(&format!("{}_r", sensor_name));
             let g_opt = get_named_spectrum(&format!("{}_g", sensor_name));
@@ -118,7 +118,7 @@ impl CreatePixelSensor for PixelSensor {
             let g = g_opt.unwrap();
             let b = b_opt.unwrap();
 
-            return Ok(Self::new(
+            Ok(Self::new(
                 &r,
                 &g,
                 &b,
@@ -130,7 +130,7 @@ impl CreatePixelSensor for PixelSensor {
                 ),
                 imaging_ratio,
                 arena
-            ));
+            ))
         }
     }
 
@@ -176,7 +176,7 @@ impl CreatePixelSensor for PixelSensor {
         let sensor_white_g = illum.inner_product(&Spectrum::Dense(g_ptr));
         let sensor_white_y = illum.inner_product(&Spectrum::Dense(spectra.y));
         for i in 0..N_SWATCH_REFLECTANCES {
-            let s = swatches[i].clone();
+            let s = swatches[i];
             let xyz = PixelSensor::project_reflectance::<XYZ>(
                 &s,
                 illum,
@@ -184,7 +184,7 @@ impl CreatePixelSensor for PixelSensor {
                 &Spectrum::Dense(spectra.y),
                 &Spectrum::Dense(spectra.z),
             ) * (sensor_white_y / sensor_white_g);
-            for c in 0..3 as u32 {
+            for c in 0..3_u32 {
                 xyz_output[i][c as usize] = xyz[c].try_into().unwrap();
             }
         }
@@ -362,7 +362,7 @@ pub trait FilmTrait: Sync {
             })
             .collect();
 
-        let mut image = HostImage::new(format, resolution, channel_names, SRGB.into());
+        let mut image = HostImage::new(format, resolution, channel_names, SRGB);
         let _rgb_desc = ImageChannelDesc::new(&[0, 1, 2]);
 
         for (iy, row_data) in processed_rows.into_iter().enumerate() {

src/core/filter.rs

@@ -2,10 +2,8 @@ use crate::Arena;
 use crate::{FileLoc, ParameterDictionary};
 use anyhow::{bail, Result};
 use shared::core::filter::Filter;
-use shared::core::geometry::{Bounds2f, Point2f, Vector2f};
+use shared::core::geometry::Vector2f;
 use shared::filters::*;
-use shared::utils::sampling::PiecewiseConstant2D;
-use shared::{Array2D, Float};
 
 pub trait FilterFactory {
     fn create(

src/core/image/io.rs

@@ -302,7 +302,7 @@ fn read_pfm(path: &Path) -> Result<ImageAndMetadata> {
 
     let names: &[&str] = if channels == 1 { &["Y"] } else { &["R", "G", "B"] };
 
-    let image = HostImage::new(PixelFormat::F32, Point2i::new(w, h), names, LINEAR.into());
+    let image = HostImage::new(PixelFormat::F32, Point2i::new(w, h), names, LINEAR);
     let metadata = ImageMetadata::default();
 
     Ok(ImageAndMetadata { image, metadata })

src/core/image/mod.rs

@@ -1,14 +1,12 @@
 use anyhow::{anyhow, Result};
-use half::f16;
 use rayon::prelude::{IndexedParallelIterator, ParallelIterator, ParallelSliceMut};
-use shared::core::color::{ColorEncoding, ColorEncodingTrait, LINEAR};
+use shared::core::color::ColorEncoding;
 use shared::core::geometry::{Bounds2f, Point2f, Point2i};
-use shared::core::image::{Image, ImageBase, PixelFormat, Pixels, WrapMode, WrapMode2D};
+use shared::core::image::{Image, PixelFormat, WrapMode, WrapMode2D};
 use shared::utils::math::square;
-use shared::{Array2D, Float, Ptr};
-use smallvec::{smallvec, SmallVec};
+use shared::{Array2D, Float};
+use smallvec::SmallVec;
 use std::ops::{Deref, DerefMut};
-use std::sync::Arc;
 
 pub mod io;
 pub mod metadata;

src/core/image/ops.rs

@@ -1,6 +1,5 @@
 use super::HostImage;
 use rayon::prelude::*;
-use shared::core::color::ColorEncoding;
 use shared::core::geometry::{Bounds2i, Point2i};
 use shared::core::image::{PixelFormat, WrapMode, WrapMode2D};
 use shared::utils::math::windowed_sinc;
@@ -109,7 +108,7 @@ impl HostImage {
             PixelFormat::F32,
             new_res,
             &self.channel_names,
-            self.encoding().into(),
+            self.encoding(),
         )));
 
         let x_weights = resample_weights(res.x() as usize, new_res.x() as usize);

src/core/interaction.rs

@@ -37,7 +37,7 @@ impl InteractionGetter for SurfaceInteraction {
         camera: &Camera,
         sampler: &mut Sampler,
     ) -> Option<BSDF> {
-        self.compute_differentials(r, camera, sampler.samples_per_pixel() as i32);
+        self.compute_differentials(r, camera, sampler.samples_per_pixel());
         let material = {
             let Some(mut active_mat) = self.material.get() else {
                 return None;
@@ -47,7 +47,7 @@ impl InteractionGetter for SurfaceInteraction {
                 let ctx = MaterialEvalContext::from(&*self);
                 active_mat = mix.choose_material(&tex_eval, &ctx)?;
             }
-            active_mat.clone()
+            *active_mat
         };
         let ctx = MaterialEvalContext::from(&*self);
         let tex_eval = UniversalTextureEvaluator;
@@ -82,7 +82,7 @@ impl InteractionGetter for SurfaceInteraction {
             let ctx = MaterialEvalContext::from(self);
             active_mat = mix.choose_material(&tex_eval, &ctx)?;
         }
-        let material = active_mat.clone();
+        let material = *active_mat;
         let ctx = MaterialEvalContext::from(self);
         material.get_bssrdf(&tex_eval, &ctx, lambda)
     }

src/core/light.rs

@@ -14,7 +14,7 @@ use std::sync::Arc;
 pub fn lookup_spectrum(s: &Spectrum) -> Arc<DenselySampledSpectrum> {
     let cache = &SPECTRUM_CACHE;
     let dense_spectrum = DenselySampledSpectrum::from_spectrum(s);
-    cache.lookup(dense_spectrum).into()
+    cache.lookup(dense_spectrum)
 }
 
 fn dummy_shape() -> Shape {
@@ -60,7 +60,7 @@ pub fn create_light(
             &shape, &alpha, None, arena,
         ),
         "infinite" => crate::lights::infinite::create(
-            render_from_light, medium.into(), camera_transform,
+            render_from_light, medium, camera_transform,
             parameters, None, loc, arena,
         ),
         "diffuse" => Err(anyhow!(

src/core/render.rs

@@ -11,7 +11,6 @@ use shared::core::primitive::PrimitiveTrait;
 use shared::core::sampler::CameraSample;
 use shared::spectra::{SampledWavelengths, LAMBDA_MAX, LAMBDA_MIN};
 use shared::Float;
-use std::sync::Arc;
 
 pub fn render_scene(scene: &BasicScene, arena: &Arena) -> Result<()> {
     let media = scene.create_media();
@@ -19,7 +18,7 @@ pub fn render_scene(scene: &BasicScene, arena: &Arena) -> Result<()> {
     let (named_materials, materials) = scene.create_materials(&textures, arena)?;
     let lights = scene.create_lights(&textures, arena);
 
-    let have_scattering = {
+    let _have_scattering = {
         let shapes = scene.shapes.lock();
         let animated = scene.animated_shapes.lock();
         shapes
@@ -37,16 +36,9 @@ pub fn render_scene(scene: &BasicScene, arena: &Arena) -> Result<()> {
     all_lights.extend(area_lights);
 
     let camera = scene.get_camera().unwrap();
-    let film = camera.get_film();
+    let _film = camera.get_film();
     warn!("Creating integrator");
     let sampler = scene.get_sampler()?;
-    let integrator = scene.create_integrator(
-        camera.clone(),
-        sampler.clone(),
-        aggregate.clone(),
-        all_lights,
-        arena,
-    );
 
     if get_options().pixel_material.is_some() {
         let lambda =
@@ -96,16 +88,33 @@ pub fn render_scene(scene: &BasicScene, arena: &Arena) -> Result<()> {
                     depth += 1;
                     ray = intr.spawn_ray(ray.d);
                 }
-            } else {
-                if depth == 1 {
+            } else if depth == 1 {
                 bail!("No geometry visible from pixel")
             } else {
                 break;
             }
         }
     }
-    }
 
+    if get_options().wavefront {
+        eprintln!("RENDER: Wavefront backend");
+        let mut wf = scene.create_wavefront_integrator(
+            camera.clone(),
+            sampler.clone(),
+            aggregate.clone(),
+            all_lights,
+            arena,
+        );
+        wf.render();
+    } else {
+        eprintln!("RENDER: Path integrator backend");
+        let integrator = scene.create_integrator(
+            camera.clone(),
+            sampler.clone(),
+            aggregate.clone(),
+            all_lights,
+            arena,
+        );
         render(
             &integrator,
             &integrator.base,
@@ -113,5 +122,8 @@ pub fn render_scene(scene: &BasicScene, arena: &Arena) -> Result<()> {
             sampler.as_ref(),
             arena,
         );
-    Ok(())
+    }
+
+    Ok(())
+
 }

src/core/scene/builder.rs

@@ -260,7 +260,7 @@ impl ParserTarget for BasicSceneBuilder {
 
     fn color_space(&mut self, name: &str, loc: FileLoc) -> Result<(), ParserError> {
         let stdcs = get_colorspace_device();
-        let _ = match stdcs.get_named(name) {
+        match stdcs.get_named(name) {
             Some(cs) => {
                 self.graphics_state.color_space = Some(Arc::new(*cs.get().unwrap()));
             }
@@ -591,7 +591,7 @@ impl ParserTarget for BasicSceneBuilder {
             self.current_accelerator
                 .take()
                 .expect("Accelerator not set before WorldBegin"),
-            &arena,
+            arena,
         );
         Ok(())
     }

src/core/scene/scene.rs

@@ -11,6 +11,7 @@ use crate::core::sampler::SamplerFactory;
 use crate::core::shape::{ShapeFactory, ShapeWithContext};
 use crate::core::texture::{FloatTexture, SpectrumTexture};
 use crate::integrators::{CreateIntegrator, PathConfig, PathIntegrator};
+use crate::lights::sampler::create_light_sampler;
 use crate::utils::parallel::{run_async, AsyncJob};
 use crate::utils::parameters::{NamedTextures, ParameterDictionary, TextureParameterDictionary};
 use crate::utils::resolve_filename;
@@ -19,21 +20,22 @@ use anyhow::{anyhow, Result};
 use parking_lot::Mutex;
 use shared::core::aggregates::{BVHAggregate, SplitMethod};
 use shared::core::camera::CameraTrait;
-use shared::core::camera::{Camera, CameraTransform};
+use shared::core::camera::Camera;
 use shared::core::color::LINEAR;
 use shared::core::film::Film;
 use shared::core::filter::Filter;
-use shared::core::light::Light;
+use shared::core::light::{Light, LightTrait};
 use shared::core::material::Material;
 use shared::core::medium::{Medium, MediumInterface};
 use shared::core::primitive::{AnimatedPrimitive, GeometricPrimitive, Primitive, SimplePrimitive};
-use shared::core::sampler::Sampler;
+use shared::core::sampler::{Sampler, SamplerTrait};
 use shared::core::shape::Shape;
-use shared::core::texture::{GPUFloatTexture, SpectrumType};
-use shared::lights::sampler::LightSampler;
+use shared::core::texture::SpectrumType;
 use shared::spectra::RGBColorSpace;
 use shared::textures::FloatConstantTexture;
-use shared::{Ptr, Transform};
+use shared::utils::soa::SoA;
+use shared::wavefront::*;
+use shared::{gvec, Ptr, WorkQueue};
 use std::collections::HashMap;
 use std::sync::Arc;
 
@@ -146,6 +148,12 @@ pub struct BasicScene {
     pub film_state: Mutex<SingletonState<Film>>,
 }
 
+impl Default for BasicScene {
+    fn default() -> Self {
+        Self::new()
+    }
+}
+
 impl BasicScene {
     pub fn new() -> Self {
         Self {
@@ -196,7 +204,7 @@ impl BasicScene {
                 &film.parameters,
                 exposure_time,
                 filter,
-                Some(camera.camera_transform.clone()),
+                Some(camera.camera_transform),
                 &film.loc,
                 arena,
             )
@@ -508,12 +516,12 @@ impl BasicScene {
         state.map.clone()
     }
 
-    pub fn create_lights(&self, textures: &NamedTextures, arena: &Arena) -> Vec<Arc<Light>> {
+    pub fn create_lights(&self, _textures: &NamedTextures, arena: &Arena) -> Vec<Arc<Light>> {
         let light_state = self.light_state.lock();
         let camera = self
             .get_camera()
             .expect("Camera must be initialized before lights");
-        let camera_transform = camera.base().camera_transform.clone();
+        let camera_transform = camera.base().camera_transform;
         let mut lights: Vec<Arc<Light>> = Vec::new();
 
         // Non-area lights created from stored entities
@@ -533,7 +541,7 @@ impl BasicScene {
                 medium.map(|m| *m),
                 &entity.transformed_base.base.parameters,
                 &entity.transformed_base.base.loc,
-                camera_transform.clone(),
+                camera_transform,
                 arena,
             ) {
                 Ok(light) => lights.push(Arc::new(light)),
@@ -709,6 +717,90 @@ impl BasicScene {
         }
     }
 
+    pub fn create_wavefront_integrator(
+        &self,
+        camera: Arc<Camera>,
+        sampler: Arc<Sampler>,
+        aggregate: Arc<Primitive>,
+        lights: Vec<Arc<Light>>,
+        arena: &Arena,
+    ) -> WavefrontPathIntegrator<CpuAggregate> {
+        let entity = self.integrator.lock().clone().unwrap();
+        let max_depth = entity
+            .parameters
+            .get_one_int("maxdepth", 5)
+            .expect("Could not obtain depth value");
+        let regularize = entity
+            .parameters
+            .get_one_bool("regularize", false)
+            .expect("Could not obtain regularize flag value");
+
+        let spp = sampler.samples_per_pixel() as u32;
+        let film = camera.base().film;
+        let pixel_bounds = film.pixel_bounds();
+        let filter = Ptr::from(&film.base().filter);
+        let light_sampler = create_light_sampler("power", &lights, arena);
+        let res_x = pixel_bounds.diagonal().x() as u32;
+        let max_samples = 1024u32 * 1024;
+        let scanlines_per_pass = (max_samples / res_x).max(1);
+        let max_queue_size = res_x * scanlines_per_pass;
+
+        let mut infinite_lights = gvec();
+        for light in &lights {
+            if light.light_type().is_infinite() {
+                infinite_lights.push(arena.alloc(**light));
+            }
+        }
+
+        let cpu_aggregate = CpuAggregate::new(*aggregate);
+
+        WavefrontPathIntegrator {
+            aggregate: cpu_aggregate,
+            camera: (*camera).clone(),
+            sampler: (*sampler).clone(),
+            max_depth: max_depth.try_into().unwrap(),
+            film,
+            filter,
+            samples_per_pixel: spp,
+            regularize,
+            infinite_lights,
+            max_queue_size,
+            scanlines_per_pass,
+            light_sampler,
+            ray_queues: [
+                WorkQueue::new(
+                    RayWorkItemSoA::allocate(max_queue_size, arena),
+                    max_queue_size,
+                ),
+                WorkQueue::new(
+                    RayWorkItemSoA::allocate(max_queue_size, arena),
+                    max_queue_size,
+                ),
+            ],
+            shadow_ray_queue: WorkQueue::new(
+                ShadowRayWorkItemSoA::allocate(max_queue_size, arena),
+                max_queue_size,
+            ),
+            escaped_ray_queue: WorkQueue::new(
+                EscapedRayWorkItemSoA::allocate(max_queue_size, arena),
+                max_queue_size,
+            ),
+            hit_area_light_queue: WorkQueue::new(
+                HitAreaLightWorkItemSoA::allocate(max_queue_size, arena),
+                max_queue_size,
+            ),
+            basic_eval_material_queue: WorkQueue::new(
+                MaterialEvalWorkItemSoA::allocate(max_queue_size, arena),
+                max_queue_size,
+            ),
+            universal_eval_material_queue: WorkQueue::new(
+                MaterialEvalWorkItemSoA::allocate(max_queue_size, arena),
+                max_queue_size,
+            ),
+            pixel_sample_state: PixelSampleState::allocate(max_queue_size, arena),
+        }
+    }
+
     // Getters
 
     pub fn get_camera(&self) -> Result<Arc<Camera>> {
@@ -847,7 +939,7 @@ impl BasicScene {
                             shape,
                             arena.alloc(mtl),
                             light_ptr,
-                            mi.clone(),
+                            mi,
                             alpha_ptr,
                         ))
                     };
@@ -866,9 +958,9 @@ impl BasicScene {
         materials: &[Material],
         light_state: &LightState,
         media: &HashMap<String, Arc<Medium>>,
-        film_cs: Option<&RGBColorSpace>,
+        _film_cs: Option<&RGBColorSpace>,
         arena: &Arena,
-        area_lights: &mut Vec<Arc<Light>>,
+        _area_lights: &mut Vec<Arc<Light>>,
     ) -> Vec<Primitive> {
         let mut primitives = Vec::new();
 
@@ -945,7 +1037,7 @@ impl BasicScene {
                         shape,
                         arena.alloc(mtl),
                         Ptr::null(), // no area light on animated shapes
-                        mi.clone(),
+                        mi,
                         alpha_ptr,
                     ))
                 };
@@ -1103,7 +1195,6 @@ impl BasicScene {
         Err(anyhow!("{} requested but not initialized!", name))
     }
 
-
     #[allow(dead_code)]
     fn upload_shapes(
         &self,

src/films/gbuffer.rs

@@ -3,7 +3,6 @@ use crate::core::film::{CreateFilmBase, CreatePixelSensor};
 use shared::core::film::PixelSensor;
 use anyhow::{Result, anyhow};
 use shared::core::film::{FilmBase, GBufferFilm};
-use shared::spectra::RGBColorSpace;
 use shared::utils::AnimatedTransform;
 use std::path::Path;
 
@@ -26,7 +25,7 @@ impl CreateFilm for GBufferFilm {
 
         let filename = params.get_one_string("filename", "pbrt.exr")?;
         if Path::new(&filename).extension() != Some("exr".as_ref()) {
-            return Err(anyhow!("{}: EXR is the only format supported by GBufferFilm", loc).into());
+            return Err(anyhow!("{}: EXR is the only format supported by GBufferFilm", loc));
         }
 
         let coords_system = params.get_one_string("coordinatesystem", "camera")?;
@@ -43,8 +42,7 @@ impl CreateFilm for GBufferFilm {
                 "{}: unknown coordinate system for GBufferFilm. (Expecting camera 
                   or world",
                 loc
-            )
-            .into());
+            ));
         };
 
         let film = GBufferFilm::new(

src/films/mod.rs

@@ -9,9 +9,6 @@ pub mod gbuffer;
 pub mod rgb;
 pub mod spectral;
 
-pub use gbuffer::*;
-pub use rgb::*;
-pub use spectral::*;
 
 pub trait CreateFilm {
     fn create(

src/films/rgb.rs

@@ -3,8 +3,7 @@ use crate::core::film::{CreateFilmBase, CreatePixelSensor};
 use crate::Arena;
 use anyhow::Result;
 use shared::core::camera::CameraTransform;
-use shared::core::film::{Film, FilmBase, RGBFilm, RGBPixel, PixelSensor};
-use shared::spectra::RGBColorSpace;
+use shared::core::film::{Film, FilmBase, RGBFilm, PixelSensor};
 
 impl CreateFilm for RGBFilm {
     fn create(

src/films/spectral.rs

@@ -5,7 +5,6 @@ use anyhow::{anyhow, Result};
 use shared::core::camera::CameraTransform;
 use shared::core::film::{FilmBase, PixelSensor, SpectralFilm};
 use shared::spectra::{LAMBDA_MAX, LAMBDA_MIN};
-use shared::utils::math::SquareMatrix;
 use shared::Float;
 use std::path::Path;
 
@@ -28,7 +27,7 @@ impl CreateFilm for SpectralFilm {
 
         let filename = params.get_one_string("filename", "pbrt.exr")?;
         if Path::new(&filename).extension() != Some("exr".as_ref()) {
-            return Err(anyhow!("{}: EXR is the only format supported by GBufferFilm", loc).into());
+            return Err(anyhow!("{}: EXR is the only format supported by GBufferFilm", loc));
         }
 
         let n_buckets = params.get_one_int("nbuckets", 16)? as usize;

src/integrators/mod.rs

@@ -15,7 +15,6 @@ use shared::core::geometry::{Point2i, Ray};
 use shared::core::light::Light;
 use shared::core::primitive::Primitive;
 use shared::core::sampler::Sampler;
-use shared::lights::sampler::LightSampler;
 use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use std::sync::Arc;
 
@@ -58,7 +57,7 @@ impl CreateIntegrator for PathIntegrator {
     fn create(
         parameters: ParameterDictionary,
         camera: Arc<Camera>,
-        sampler: Arc<Sampler>,
+        _sampler: Arc<Sampler>,
         aggregate: Arc<Primitive>,
         lights: Vec<Arc<Light>>,
         config: PathConfig,

src/integrators/path.rs

@@ -126,7 +126,7 @@ impl PathIntegrator {
 
         if !self
             .base
-            .unoccluded(&Interaction::Surface(intr.clone()), &ls.p_light)
+            .unoccluded(&Interaction::Surface(*intr), &ls.p_light)
         {
             return SampledSpectrum::zero();
         }
@@ -232,8 +232,8 @@ impl RayIntegratorTrait for PathIntegrator {
             if !le.is_black() {
                 if state.depth == 0 || state.specular_bounce {
                     state.l += state.beta * le;
-                } else if self.config.use_mis {
-                    if !isect.area_light.is_null() {
+                } else if self.config.use_mis
+                    && !isect.area_light.is_null() {
                         let light = &isect.area_light;
                         let p_l = self.sampler.pmf_with_context(&state.prev_ctx, light)
                             * light.pdf_li(&state.prev_ctx, ray.d, true);
@@ -241,7 +241,6 @@ impl RayIntegratorTrait for PathIntegrator {
                         state.l += state.beta * w_b * le;
                     }
             }
-            }
 
             // Get BSDF
             let Some(mut bsdf) = isect.get_bsdf(&ray, lambda, &self.camera, sampler) else {

src/integrators/pipeline.rs

@@ -9,14 +9,13 @@ use crate::Arena;
 use indicatif::{ProgressBar, ProgressStyle};
 use rayon::iter::{IntoParallelRefIterator, ParallelIterator};
 use shared::core::camera::{Camera, CameraTrait};
-use shared::core::geometry::{Bounds2i, Point2i, VectorLike};
+use shared::core::geometry::{Bounds2i, Point2i};
 use shared::core::sampler::get_camera_sample;
 use shared::core::sampler::{Sampler, SamplerTrait};
 use shared::spectra::SampledSpectrum;
 use shared::Float;
 use std::io::Write;
 use std::path::Path;
-use std::sync::Arc;
 
 struct PbrtProgress {
     bar: ProgressBar,
@@ -95,7 +94,7 @@ pub fn render<T>(
             &mut tile_sampler,
             p_pixel,
             s_index,
-            &arena,
+            arena,
         );
         return;
     }
@@ -158,7 +157,7 @@ pub fn render<T>(
         if let Some(out_path) = &options.mse_reference_output {
             mse_out_file = Some(
                 std::fs::File::create(out_path)
-                    .expect(&format!("Failed to create MSE output file: {}", out_path)),
+                    .unwrap_or_else(|_| panic!("Failed to create MSE output file: {}", out_path)),
             );
         }
     }
@@ -177,7 +176,7 @@ pub fn render<T>(
                         &mut sampler,
                         p_pixel,
                         sample_index.try_into().unwrap(),
-                        &arena,
+                        arena,
                     );
                 }
             }
@@ -200,7 +199,7 @@ pub fn render<T>(
         if wave_start == spp || options.write_partial_images || reference_image.is_some() {
             let mut metadata = ImageMetadata {
                 render_time_seconds: Some(progress.elapsed_seconds()),
-                samples_per_pixel: Some(wave_start as i32),
+                samples_per_pixel: Some(wave_start),
                 ..Default::default()
             };
 

src/integrators/state.rs

@@ -14,6 +14,12 @@ pub struct PathState {
     pub prev_pdf: Float,
 }
 
+impl Default for PathState {
+    fn default() -> Self {
+        Self::new()
+    }
+}
+
 impl PathState {
     pub fn new() -> Self {
         Self {

src/lib.rs

@@ -14,4 +14,6 @@ pub mod utils;
 pub mod wavefront;
 
 pub use utils::{Arena, FileLoc, ParameterDictionary, Upload, ArenaUpload};
-pub const MAX_TAGS = 16;
+pub const MAX_TAGS: u32 = 16;
+pub use shared::{BasicPBRTOptions, PBRTOptions};
+pub use globals::{get_options, init_pbrt};

src/lights/distant.rs

@@ -11,7 +11,7 @@ use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
 use shared::lights::DistantLight;
 use shared::spectra::RGBColorSpace;
-use shared::utils::{Ptr, Transform};
+use shared::utils::Transform;
 use shared::Float;
 
 trait CreateDistantLight {

src/lights/goniometric.rs

@@ -69,12 +69,11 @@ pub fn create(
     scale /= spectrum_to_photometric(i);
 
     let phi_v = params.get_one_float("power", -1.0)?;
-    if phi_v > 0.0 {
-        if let Some(ref img) = host_image {
+    if phi_v > 0.0
+        && let Some(ref img) = host_image {
             let k_e = compute_emissive_power(img);
             scale *= phi_v / k_e;
         }
-    }
 
     let swap_yz: [Float; 16] = [
         1., 0., 0., 0., 0., 0., 1., 0., 0., 1., 0., 0., 0., 0., 0., 1.,
@@ -146,7 +145,7 @@ fn convert_to_luminance_image(
                 }
             }
 
-            Ok(HostImage::from_f32(&y_pixels, res, &["Y"].to_vec()))
+            Ok(HostImage::from_f32(&y_pixels, res, ["Y"].as_ref()))
         }
 
         (Err(_), Ok(_)) => {

src/lights/infinite.rs

@@ -9,7 +9,7 @@ use rayon::prelude::*;
 use shared::core::camera::CameraTransform;
 use shared::core::geometry::{cos_theta, Bounds2f, Frame, Point2f, Point2i, Point3f, VectorLike};
 use shared::core::image::WrapMode;
-use shared::core::light::{Light, LightBase, LightType};
+use shared::core::light::Light;
 use shared::core::medium::Medium;
 use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
@@ -17,7 +17,7 @@ use shared::lights::{ImageInfiniteLight, PortalInfiniteLight, UniformInfiniteLig
 use shared::spectra::RGBColorSpace;
 use shared::utils::math::{equal_area_sphere_to_square, equal_area_square_to_sphere};
 use shared::utils::sampling::{PiecewiseConstant2D, WindowedPiecewiseConstant2D};
-use shared::{Float, Ptr, Transform, PI};
+use shared::{Float, Transform, PI};
 use std::path::Path;
 
 pub fn create(
@@ -275,7 +275,7 @@ fn load_image(
         let rgb = l[0].to_rgb(colorspace, &stdspec);
         let image =
             HostImage::new_constant(Point2i::new(1, 1), &["R", "G", "B"], &[rgb.r, rgb.g, rgb.b]);
-        return Ok((image, colorspace.clone()));
+        return Ok((image, *colorspace));
     }
 
     let im = HostImage::read(Path::new(filename), None)
@@ -290,7 +290,7 @@ fn load_image(
         .get_channel_desc(&["R", "G", "B"])
         .map_err(|_| anyhow!("image '{}' must have R, G, B channels", filename))?;
 
-    let cs = im.metadata.colorspace.unwrap_or_else(|| colorspace.clone());
+    let cs = im.metadata.colorspace.unwrap_or_else(|| *colorspace);
     Ok((im.image.select_channels(&desc), cs))
 }
 

src/lights/point.rs

@@ -11,7 +11,7 @@ use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
 use shared::lights::PointLight;
 use shared::spectra::RGBColorSpace;
-use shared::{Float, PI, Ptr, Transform};
+use shared::{Float, PI, Transform};
 
 pub trait CreatePointLight {
     fn new(

src/lights/sampler.rs

@@ -1,7 +1,7 @@
 use crate::Arena;
 use shared::core::light::{Light, LightTrait};
 use shared::lights::sampler::{
-    BVHLightSampler, LightSampler, PowerLightSampler, UniformLightSampler,
+    LightSampler, PowerLightSampler, UniformLightSampler,
 };
 use shared::utils::sampling::AliasTable;
 use shared::spectra::{SampledSpectrum, SampledWavelengths};

src/lights/spot.rs

@@ -12,7 +12,7 @@ use shared::core::texture::SpectrumType;
 use shared::lights::SpotLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::math::radians;
-use shared::{Float, Ptr, Transform, PI};
+use shared::{Float, Transform, PI};
 
 trait CreateSpotLight {
     fn new(

src/samplers/halton.rs

@@ -23,7 +23,7 @@ impl CreateHaltonSampler for HaltonSampler {
         full_res: Point2i,
         randomize: RandomizeStrategy,
         seed: u64,
-        arena: &Arena,
+        _arena: &Arena,
     ) -> Self {
         let digit_permutations = compute_radical_inverse_permutations(seed);
         let leaked = Box::leak(gbox(digit_permutations));

src/shapes/bilinear.rs

@@ -63,7 +63,7 @@ impl CreateShape for BilinearPatchShape {
             n.clear();
         }
 
-        for (_, &idx) in vertex_indices.iter().enumerate() {
+        for &idx in vertex_indices.iter() {
             if idx < 0 || idx as usize >= p.len() {
                 return Err(anyhow!(
                     "Bilinear patch mesh has out-of-bounds vertex index {} ({} \"P\" values were given). Discarding this mesh.",

src/shapes/curves.rs

@@ -45,7 +45,7 @@ pub fn create_curve(
         let u_max = (i + 1) as Float / n_segments as Float;
 
         let curve = CurveShape {
-            common: curve_common.clone(),
+            common: curve_common,
             u_min,
             u_max,
         };
@@ -90,7 +90,7 @@ impl CreateShape for CurveShape {
 
         if basis == "bezier" {
             if cp.len() <= degree as usize
-                || ((cp.len() - 1 - degree as usize) % degree as usize) != 0
+                || !(cp.len() - 1 - degree as usize).is_multiple_of(degree as usize)
             {
                 return Err(anyhow!(
                     "Invalid number of control points {}: for the degree {} Bezier basis {} + n * {} are required.",

src/shapes/mesh.rs

@@ -3,7 +3,6 @@ use ply_rs::parser::Parser;
 use ply_rs::ply::{DefaultElement, Property};
 use shared::core::geometry::{Normal3f, Point2f, Point3f, VectorLike};
 use shared::shapes::mesh::TriangleMesh;
-use shared::utils::sampling::PiecewiseConstant2D;
 use shared::Transform;
 use std::fs::File;
 use std::path::Path;
@@ -214,9 +213,9 @@ impl TriQuadMesh {
             let v20 = p2 - p0;
             let face_normal = v10.cross(v20);
 
-            self.n[i0] = self.n[i0] + Normal3f::from(face_normal);
-            self.n[i1] = self.n[i1] + Normal3f::from(face_normal);
-            self.n[i2] = self.n[i2] + Normal3f::from(face_normal);
+            self.n[i0] += Normal3f::from(face_normal);
+            self.n[i1] += Normal3f::from(face_normal);
+            self.n[i2] += Normal3f::from(face_normal);
         }
 
         for quad in self.quad_indices.chunks_exact(4) {
@@ -231,14 +230,14 @@ impl TriQuadMesh {
             let face_normal = v10.cross(v20);
 
             for &idx in &indices {
-                self.n[idx] = self.n[idx] + Normal3f::from(face_normal);
+                self.n[idx] += Normal3f::from(face_normal);
             }
         }
 
         for normal in &mut self.n {
             let len_sq = normal.norm_squared();
             if len_sq > 0.0 {
-                *normal = *normal / len_sq.sqrt();
+                *normal /= len_sq.sqrt();
             }
         }
     }

src/shapes/triangle.rs

@@ -60,7 +60,7 @@ impl CreateShape for TriangleShape {
             n.clear();
         }
 
-        for (_, &index) in vertex_indices.iter().enumerate() {
+        for &index in vertex_indices.iter() {
             // Check for negative indices (if keeping i32) or out of bounds
             if index < 0 || index as usize >= p.len() {
                 bail!(

src/spectra/colorspace.rs

@@ -5,7 +5,6 @@ use shared::core::spectrum::Spectrum;
 use shared::spectra::{DenselySampledSpectrum, RGBColorSpace};
 use shared::utils::math::SquareMatrix;
 use shared::Ptr;
-use std::sync::Arc;
 
 pub trait CreateRGBColorSpace {
     fn new(

src/spectra/data.rs

@@ -1,7 +1,7 @@
 use shared::core::spectrum::Spectrum;
 use shared::spectra::cie::*;
 use shared::spectra::{DenselySampledSpectrum, PiecewiseLinearSpectrum};
-use shared::{gbox, gvec_from_slice, leak, Float, Ptr};
+use shared::{gvec_from_slice, leak, Float};
 use std::collections::HashMap;
 use std::sync::LazyLock;
 

src/spectra/mod.rs

@@ -107,10 +107,10 @@ impl StandardColorSpaces {
 
 pub fn get_colorspace_context() -> StandardColorSpaces {
     StandardColorSpaces {
-        srgb: SRGB.clone().into(),
-        dci_p3: DCI_P3.clone().into(),
-        rec2020: REC2020.clone().into(),
-        aces2065_1: ACES.clone().into(),
+        srgb: SRGB.clone(),
+        dci_p3: DCI_P3.clone(),
+        rec2020: REC2020.clone(),
+        aces2065_1: ACES.clone(),
     }
 }
 

src/utils/arena.rs

@@ -1,6 +1,7 @@
 use crate::utils::backend::GpuAllocator;
 use crate::utils::mipmap::MIPMap;
 use parking_lot::Mutex;
+use shared::utils::soa::SoAAllocator;
 use shared::Ptr;
 use std::alloc::Layout;
 use std::collections::HashMap;
@@ -161,6 +162,11 @@ impl<A: GpuAllocator> Arena<A> {
         (Ptr::from_raw(ptr), len)
     }
 
+    pub fn alloc_layout(&self, layout: Layout) -> *mut u8 {
+        let mut bump = self.bump.lock();
+        bump.alloc_layout(layout)
+    }
+
     pub fn get_texture_object(&self, mipmap: &Arc<MIPMap>) -> u64 {
         let key = Arc::as_ptr(mipmap) as usize;
         let mut cache = self.texture_cache.lock();
@@ -178,3 +184,9 @@ impl<A: GpuAllocator + Default> Default for Arena<A> {
         Self::new(A::default())
     }
 }
+
+impl<A: GpuAllocator> SoAAllocator for Arena<A> {
+    fn alloc_raw(&self, layout: core::alloc::Layout) -> *mut u8 {
+        self.alloc_layout(layout)
+    }
+}

src/utils/containers.rs

@@ -1,5 +1,4 @@
 use parking_lot::Mutex;
-use shared::core::geometry::{Bounds2i, Point2i};
 use std::collections::HashSet;
 use std::hash::Hash;
 use std::sync::Arc;
@@ -8,6 +7,15 @@ pub struct InternCache<T> {
     cache: Mutex<HashSet<Arc<T>>>,
 }
 
+impl<T> Default for InternCache<T>
+where
+    T: Eq + Clone + Hash,
+ {
+    fn default() -> Self {
+        Self::new()
+    }
+}
+
 impl<T> InternCache<T>
 where
     T: Eq + Clone + Hash,

src/utils/mipmap.rs

@@ -157,7 +157,7 @@ impl MIPMap {
     }
 
     pub fn get_rgb_colorspace(&self) -> Option<RGBColorSpace> {
-        self.color_space.clone()
+        self.color_space
     }
 
     pub fn get_level(&self, level: usize) -> &HostImage {

src/utils/parallel.rs

@@ -1,5 +1,4 @@
 use crossbeam_channel::{Receiver, bounded};
-use rayon::prelude::*;
 
 #[derive(Debug)]
 pub struct AsyncJob<T> {

src/utils/parameters.rs

@@ -515,7 +515,7 @@ impl ParameterDictionary {
                     || param.type_name == "rgb"
                     || param.type_name == "blackbody"
                 {
-                    return Some(self.extract_spectrum_array(param, stype)[0].clone());
+                    return Some(self.extract_spectrum_array(param, stype)[0]);
                 }
             }
         }
@@ -654,7 +654,7 @@ impl ParameterDictionary {
     }
 
     fn extract_sampled_spectrum(&self, param: &ParsedParameter) -> Vec<Spectrum> {
-        if param.floats.len() % 2 != 0 {
+        if !param.floats.len().is_multiple_of(2) {
             panic!(
                 "{}: Found odd number of values for '{}'",
                 param.loc, param.name
@@ -852,11 +852,11 @@ impl TextureParameterDictionary {
 
     pub fn get_float_texture(&self, name: &str, val: Float) -> Result<Arc<FloatTexture>> {
         if let Some(tex) = self.get_float_texture_or_null(name)? {
-            return Ok(tex);
+            Ok(tex)
         } else {
-            return Ok(Arc::new(FloatTexture::Constant(FloatConstantTexture::new(
+            Ok(Arc::new(FloatTexture::Constant(FloatConstantTexture::new(
                 val,
-            ))));
+            ))))
         }
     }
 
@@ -961,7 +961,7 @@ impl TextureParameterDictionary {
                 _ => {}
             }
         }
-        return None;
+        None
     }
 
     pub fn get_float_texture_or_null(&self, name: &str) -> Result<Option<Arc<FloatTexture>>> {

src/utils/upload.rs

@@ -102,7 +102,6 @@ fn convert_spectrum(tex: &SpectrumTexture, arena: &Arena) -> GPUSpectrumTexture
                     t.base
                         .mipmap
                         .color_space
-                        .clone()
                         .unwrap_or_else(crate::spectra::default_colorspace),
                 ),
             })
@@ -127,14 +126,14 @@ impl Upload for Arc<SpectrumTexture> {
 impl Upload for &FloatTexture {
     type Target = Ptr<GPUFloatTexture>;
     fn upload(self, arena: &Arena) -> Self::Target {
-        arena.alloc(convert_float(&self, arena))
+        arena.alloc(convert_float(self, arena))
     }
 }
 
 impl Upload for &SpectrumTexture {
     type Target = Ptr<GPUSpectrumTexture>;
     fn upload(self, arena: &Arena) -> Self::Target {
-        arena.alloc(convert_spectrum(&self, arena))
+        arena.alloc(convert_spectrum(self, arena))
     }
 }
 impl Upload for Option<Arc<FloatTexture>> {

src/wavefront/aggregate.rs

@@ -1,11 +1,181 @@
-use shared::core::geometry::Bounds3f;
-use super::{RayQueue, EscapedRayQueue, HitAreaLightQueue, MaterialEvalQueue, MediumSampleQueue, SubsurfaceScatterQueue};
+use crate::core::texture::{BasicTextureEvaluator, TextureEvaluator, UniversalTextureEvaluator};
+use crate::core::primitive::{Primitive, PrimitiveTrait};
+use crate::core::geometry::{Bounds3f, Ray, Vector3f, VectorLike};
+use crate::core::interaction::InteractionTrait;
+use crate::core::material::MaterialTrait;
+use crate::globals::get_options;
+use rayon::prelude::*;
+use shared::wavefront::WavefrontAggregate;
 
-#[derive(Clone, Debug)]
-pub trait WavefrontAggregate {
-    fn bounds(&self) -> Bounds3f;
-    fn intersect_closest(max_rays: usize, ray_q: &mut RayQueue, hit_area_light_q: &mut HitAreaLightQueue, basic_mlt_q: &mut MaterialEvalQueue, universal_mtl_q: &mut MaterialEvalQueue, medium_sample_q: &mut MediumSampleQueue);
-    fn intersect_shadow(max_rays: usize, shadow_ray_q: &mut ShadowRayQueue, pixel_sample_state: &mut SOA<PixelSampleState>);
-    fn intersect_shadow_tr(max_rays: usize, shadow_ray_q: &mut ShadowRayQueue, pixel_sample_state: &mut SOA<PixelSampleState>);
-    fn intersect_one_random(max_rays: usize, subsurface_scatte_q: &mut SubsurfaceScatterQueue);
+pub struct CpuAggregate {
+    pub aggregate: Primitive,
 }
+
+impl CpuAggregate {
+    pub fn new(aggregate: Primitive) -> Self {
+        Self { aggregate }
+    }
+}
+
+impl WavefrontAggregate for CpuAggregate {
+    fn bounds(&self) -> Bounds3f {
+        self.aggregate.bounds()
+    }
+
+    fn intersect_closest(
+        &self,
+        max_rays: usize,
+        ray_q: &RayQueue,
+        escaped_ray_q: &EscapedRayQueue,
+        hit_area_light_q: &HitAreaLightQueue,
+        basic_eval_mtl_q: &MaterialEvalQueue,
+        universal_eval_mtl_q: &MaterialEvalQueue,
+        next_ray_q: &RayQueue,
+        pixel_sample_state: &PixelSampleState,
+    ) {
+        let n_rays = ray_q.size().min(max_rays as u32);
+        let options = get_options();
+
+        (0..n_rays as usize).into_par_iter().for_each(|i| {
+            let work = unsafe { ray_q.get(i) };
+
+            let ray = Ray::new(work.ray_o, work.ray_d, Some(work.ray_time), work.ray_medium);
+
+            let pi = work.pixel_index as usize;
+            let beta = pixel_sample_state.beta.get(pi);
+            let r_u = pixel_sample_state.r_u.get(pi);
+            let r_l = pixel_sample_state.r_l.get(pi);
+            let lambda = pixel_sample_state.lambda.get(pi);
+            let depth = pixel_sample_state.depth.get(pi);
+            let specular_bounce = pixel_sample_state.specular_bounce.get(pi) != 0;
+            let any_non_specular = pixel_sample_state.any_non_specular_bounces.get(pi) != 0;
+            let eta_scale = pixel_sample_state.eta_scale.get(pi);
+            let prev_intr_ctx = pixel_sample_state.prev_intr_ctx.get(pi);
+
+            let Some(si) = self.aggregate.intersect(&ray, None) else {
+                escaped_ray_q.push(EscapedRayWorkItem {
+                    ray_o: work.ray_o,
+                    ray_d: work.ray_d,
+                    lambda,
+                    pixel_index: work.pixel_index,
+                    beta,
+                    r_u,
+                    r_l,
+                    depth,
+                    specular_bounce,
+                    prev_intr_ctx,
+                });
+                return;
+            };
+
+            let intr = &si.intr;
+
+            if intr.material.is_null() {
+                let ray_o = Ray::offset_origin(&intr.pi(), &intr.n(), &work.ray_d);
+                let ray_medium = if work.ray_d.dot(intr.n().into()) > 0.0 {
+                    intr.common.medium_interface.outside
+                } else {
+                    intr.common.medium_interface.inside
+                };
+                next_ray_q.push(RayWorkItem {
+                    ray_o,
+                    ray_d: work.ray_d,
+                    ray_time: work.ray_time,
+                    ray_medium,
+                    has_differentials: work.has_differentials,
+                    differential: work.differential,
+                    pixel_index: work.pixel_index,
+                });
+                return;
+            }
+
+            // Check for area light hit
+            if !intr.area_light.is_null() {
+                hit_area_light_q.push(HitAreaLightWorkItem {
+                    area_light: intr.area_light,
+                    p: intr.p(),
+                    n: intr.n(),
+                    uv: intr.common.uv,
+                    wo: -work.ray_d,
+                    lambda,
+                    pixel_index: work.pixel_index,
+                    beta,
+                    r_u,
+                    r_l,
+                    depth,
+                    specular_bounce,
+                    prev_intr_ctx,
+                });
+            }
+
+            // Determine which material evaluation queue to use based on
+            // whether the material's textures can be evaluated with the
+            // basic evaluator (cheaper) or need the universal one.
+            let material = *intr.material.get().unwrap();
+            let eval_q = if material.can_evaluate_textures(&BasicTextureEvaluator) {
+                basic_eval_mtl_q
+            } else {
+                universal_eval_mtl_q
+            };
+
+            eval_q.push(MaterialEvalWorkItem {
+                p: intr.p(),
+                n: intr.n(),
+                ns: intr.shading.n,
+                dpdu: intr.shading.dpdu,
+                dpdv: intr.shading.dpdv,
+                uv: intr.common.uv,
+                wo: -work.ray_d,
+                time: work.ray_time,
+                face_index: intr.face_index,
+                material: intr.material,
+                area_light: intr.area_light,
+                medium_interface: intr.common.medium_interface,
+                pixel_index: work.pixel_index,
+                lambda,
+                beta,
+                r_u,
+                any_non_specular_bounces: any_non_specular,
+                depth,
+                eta_scale,
+            });
+        });
+    }
+
+    fn intersect_shadow(
+        &self,
+        max_rays: usize,
+        shadow_ray_q: &ShadowRayQueue,
+        pixel_sample_state: &PixelSampleState,
+    ) {
+        let n_rays = shadow_ray_q.size().min(max_rays as u32);
+
+        for i in 0..n_rays as usize {
+            let work = unsafe { shadow_ray_q.get(i) };
+
+            let ray = Ray::new(
+                work.ray_o,
+                work.ray_d,
+                Some(work.ray_time),
+                crate::Ptr::null(),
+            );
+
+            if !self.aggregate.intersect_p(&ray, Some(work.t_max)) {
+                let pi = work.pixel_index as usize;
+                let mut l = pixel_sample_state.l.get(pi);
+                l += work.l_d;
+                pixel_sample_state.l.set(pi, l);
+            }
+        }
+    }
+
+    fn intersect_shadow_tr(
+        &self,
+        max_rays: usize,
+        shadow_ray_q: &ShadowRayQueue,
+        pixel_sample_state: &PixelSampleState,
+    ) {
+        self.intersect_shadow(max_rays, shadow_ray_q, pixel_sample_state);
+    }
+}
+

src/wavefront/integrator.rs

@@ -1,38 +1,459 @@
-use crate::MAX_TAGS;
-use shared::{Ptr, GVec};
-use shared::core::film::Film;
-use shared::core::color::RGB;
-use shared::core::filter::Filter;
-use shared::core::light::Light;
-use shared::core::sampler::Filter;
-use shared::wavefront::{WavefrontAggregate, RayQueue, MediumSampleQueue, EscapedRayQueue, HitAreaLightQueue, MaterialEvalQueue, ShadowRayQueue, GetBSSRDFAndProbeRayQueue, SubsurfaceScatterQueue};
+use super::CpuAggregate;
+use crate::globals::get_options;
+use shared::core::bxdf::{FArgs, TransportMode};
+use shared::core::camera::{Camera, CameraTrait};
+use shared::core::filter::{Filter, FilterTrait};
+use shared::core::geometry::{
+    Bounds2i, Point2f, Point2i, Point3f, Point3fi, Ray, RayDifferential, Vector2f, Vector3f,
+    VectorLike,
+};
+use shared::core::interaction::InteractionTrait;
+use shared::core::light::{Light, LightSampleContext, LightTrait};
+use shared::core::material::{MaterialEvalContext, MaterialTrait};
+use shared::core::sampler::{CameraSample, Sampler, SamplerTrait};
+use shared::core::texture::{TextureEvalContext, UniversalTextureEvaluator};
+use shared::lights::sampler::{LightSampler, LightSamplerTrait};
+use shared::spectra::{SampledSpectrum, SampledWavelengths};
+use shared::utils::math::square;
+use shared::utils::sampling::power_heuristic;
+use shared::utils::soa::{SoA, SoAAllocator, WorkQueue};
+use shared::wavefront::{WavefrontAggregate, WavefrontPathIntegrator, WavefrontRenderer};
 
-pub struct WavefrontPathIntegrator {
-    pub init_visible_surface: bool,
-    pub have_subsurface: bool,
-    pub have_media: bool,
-    pub have_basic_eval_material: [bool; MAX_TAGS + 1],
-    pub have_universal_eval_material: [bool; MAX_TAGS + 1],
-    pub filter: Filter,
-    pub film: Film, 
-    pub sampler: Sampler,
-    pub camera: Camera,
-    pub infinite_lights: GVec<Light>,
-    pub max_depth: usize,
-    pub sampler_per_pixel: usize,
-    pub regularize: bool,
-    pub scanlines_per_pixel: usize,
-    pub max_queue_size: usize,
-    pub medium_sample_queue: Ptr<MediumSampleQueue>,
-    pub medium_scatter_queue: Ptr<MediumScatterQueue>,
-    pub escaped_ray_queue: Ptr<EscapedRayQueue>,
-    pub hit_area_light_queue: Ptr<HitAreaLightQueue>,
-    pub basic_eval_material_queue: Ptr<MaterialEvalQueue>,
-    pub universal_eval_material_queue: Ptr<MaterialEvalQueue>,
-    pub shadow_ray_queue: Ptr<ShadowRayQueue>,
-    pub bssrdf_eval_queue: PTr<GetBSSRDFAndProbeRayQueue>,
-    pub subsurface_scatter_queue: Ptr<SubsurfaceScatterQueue>,
-    pub display_rgb: Ptr<RGB>,
-    pub display_rgb_host: Ptr<RGB,
+impl WavefrontRenderer for WavefrontPathIntegrator<CpuAggregate> {
+    pub fn render(&mut self) {
+        let film = self.camera.get_film();
+        let filter = film.get_filter();
+        let pixel_bounds = film.pixel_bounds();
+        let resolution = pixel_bounds.diagonal();
 
+        let total_work = (pixel_bounds.area() as u64) * (self.samples_per_pixel as u64);
+        let options = get_options();
+        let progress = PbrtProgress::new(total_work, "Rendering", options.quiet);
+
+        for sample_index in 0..self.samples_per_pixel {
+            // Process image in scanline batches
+            let mut y0 = pixel_bounds.p_min.y();
+            while y0 < pixel_bounds.p_max.y() {
+                let y1 = (y0 + self.scanlines_per_pass as i32).min(pixel_bounds.p_max.y());
+
+                // Reset the primary ray queue for this set
+                self.ray_queues[0].reset();
+
+                self.generate_camera_rays(y0, y1, sample_index, &pixel_bounds);
+
+                for depth in 0..=self.max_depth {
+                    let current = (depth % 2) as usize;
+                    let next = ((depth + 1) % 2) as usize;
+
+                    // Reset output queues before intersection
+                    self.ray_queues[next].reset();
+                    self.escaped_ray_queue.reset();
+                    self.hit_area_light_queue.reset();
+                    self.basic_eval_material_queue.reset();
+                    self.universal_eval_material_queue.reset();
+                    self.shadow_ray_queue.reset();
+
+                    // Skip if no rays to trace
+                    if self.ray_queues[current].size() == 0 {
+                        break;
+                    }
+
+                    // Sorting of rays into output queues
+                    self.aggregate.intersect_closest(
+                        self.max_queue_size as usize,
+                        &self.ray_queues[current],
+                        &self.escaped_ray_queue,
+                        &self.hit_area_light_queue,
+                        &self.basic_eval_material_queue,
+                        &self.universal_eval_material_queue,
+                        &self.ray_queues[next],
+                        &self.pixel_sample_state,
+                    );
+
+                    // Infinite light contributions
+                    self.handle_escaped_rays();
+
+                    // Area light contributions
+                    self.handle_emissive_intersections();
+
+                    // Last depth — don't evaluate materials or sample lights
+                    if depth == self.max_depth {
+                        break;
+                    }
+
+                    // Evaluate materials, sample BSDFs, sample direct lighting
+                    // This pushes to shadow_ray_queue and ray_queues[next]
+                    self.evaluate_materials_and_bsdfs(depth);
+
+                    // Add direct lighting to pixels
+                    self.aggregate.intersect_shadow(
+                        self.max_queue_size as usize,
+                        &self.shadow_ray_queue,
+                        &self.pixel_sample_state,
+                    );
+                }
+
+                self.update_film(y0, y1, &pixel_bounds);
+                let batch_pixels =
+                    ((y1 - y0) * (pixel_bounds.p_max.x() - pixel_bounds.p_min.x())) as u64;
+                progress.inc(batch_pixels);
+
+                y0 = y1;
+            }
+        }
+    }
+
+    /// Stage 1: Generate camera rays for scanlines [y0, y1).
+    fn generate_camera_rays(
+        &mut self,
+        y0: i32,
+        y1: i32,
+        sample_index: u32,
+        pixel_bounds: &Bounds2i,
+    ) {
+        // For each pixel in the scanline range, generate a camera ray
+        // and push it to the ray queue. Also initialize the PixelSampleState.
+        for y in y0..y1 {
+            for x in pixel_bounds.p_min.x()..pixel_bounds.p_max.x() {
+                let p_pixel = Point2i::new(x, y);
+
+                // TODO: proper sampler state per pixel/sample
+                // For now, use a simple approach
+                self.sampler
+                    .start_pixel_sample(p_pixel, sample_index as i32, Some(0));
+
+                let lambda = SampledWavelengths::sample_visible(self.sampler.get1d());
+
+                let camera_sample = crate::core::sampler::get_camera_sample(
+                    &mut self.sampler,
+                    p_pixel,
+                    &self.filter,
+                );
+
+                let Some(camera_ray) = self.camera.generate_ray(camera_sample, &lambda) else {
+                    continue;
+                };
+
+                // Compute pixel index for this sample
+                let pixel_index = self.ray_queues[0].size();
+
+                // Initialize persistent pixel state
+                let pi = pixel_index as usize;
+                self.pixel_sample_state.l.set(pi, SampledSpectrum::new(0.0));
+                self.pixel_sample_state.beta.set(pi, camera_ray.weight);
+                self.pixel_sample_state.lambda.set(pi, lambda);
+                self.pixel_sample_state
+                    .r_u
+                    .set(pi, SampledSpectrum::new(1.0));
+                self.pixel_sample_state
+                    .r_l
+                    .set(pi, SampledSpectrum::new(1.0));
+                self.pixel_sample_state.depth.set(pi, 0);
+                self.pixel_sample_state.specular_bounce.set(pi, 1);
+                self.pixel_sample_state.any_non_specular_bounces.set(pi, 0);
+                self.pixel_sample_state.eta_scale.set(pi, 1.0);
+                self.pixel_sample_state.p_film.set(pi, camera_sample.p_film);
+                self.pixel_sample_state
+                    .filter_weight
+                    .set(pi, camera_sample.filter_weight);
+                self.pixel_sample_state
+                    .prev_intr_ctx
+                    .set(pi, LightSampleContext::default());
+
+                // Push ray to queue
+                self.ray_queues[0].push(RayWorkItem {
+                    ray_o: camera_ray.ray.o,
+                    ray_d: camera_ray.ray.d,
+                    ray_time: camera_ray.ray.time,
+                    ray_medium: camera_ray.ray.medium,
+                    pixel_index: pixel_index,
+                    has_differentials: true,
+                    differential: RayDifferential::default(),
+                });
+            }
+        }
+    }
+
+    /// Handle escaped rays — evaluate infinite lights.
+    fn handle_escaped_rays(&self) {
+        let n = self.escaped_ray_queue.size();
+        for i in 0..n as usize {
+            let w = unsafe { self.escaped_ray_queue.storage.get(i) };
+
+            let mut l_contrib = SampledSpectrum::new(0.0);
+
+            // Evaluate all infinite lights
+            for light_ptr in &self.infinite_lights {
+                let light = light_ptr.get().unwrap();
+                let ray = crate::core::geometry::Ray::new(w.ray_o, w.ray_d, None, Ptr::null());
+                let le = light.le(&ray, &w.lambda);
+                if le.is_black() {
+                    continue;
+                }
+
+                if w.depth == 0 || w.specular_bounce {
+                    // No MIS for direct camera rays or specular bounces
+                    l_contrib += w.beta * le / w.r_u.average();
+                } else {
+                    // MIS with light sampling
+                    // TODO: compute light PDF for MIS weight
+                    // For now, use unidirectional weight only
+                    l_contrib += w.beta * le / w.r_u.average();
+                }
+            }
+
+            if !l_contrib.is_black() {
+                let pi = w.pixel_index as usize;
+                let mut l = self.pixel_sample_state.l.get(pi);
+                l += l_contrib;
+                self.pixel_sample_state.l.set(pi, l);
+            }
+        }
+    }
+
+    /// Handle emissive intersections — area light contribution with MIS.
+    fn handle_emissive_intersections(&self) {
+        let n = self.hit_area_light_queue.size();
+        for i in 0..n as usize {
+            let w = unsafe { self.hit_area_light_queue.storage.get(i) };
+
+            let light = w.area_light.get().unwrap();
+            let le = light.l(w.p, w.n, w.uv, w.wo, &w.lambda);
+            if le.is_black() {
+                continue;
+            }
+
+            let l_contrib = if w.depth == 0 || w.specular_bounce {
+                w.beta * le / w.r_u.average()
+            } else {
+                // MIS: combine BSDF and light sampling weights
+                // TODO: full MIS with light sampler PDF
+                w.beta * le / w.r_u.average()
+            };
+
+            if !l_contrib.is_black() {
+                let pi = w.pixel_index as usize;
+                let mut l = self.pixel_sample_state.l.get(pi);
+                l += l_contrib;
+                self.pixel_sample_state.l.set(pi, l);
+            }
+        }
+    }
+
+    fn evaluate_materials_and_bsdfs(&mut self, depth: u32) {
+        self.evaluate_material_queue_impl(depth, false);
+        self.evaluate_material_queue_impl(depth, true);
+    }
+
+    fn evaluate_material_queue_impl(&mut self, depth: u32, use_universal: bool) {
+        let queue = if use_universal {
+            &self.universal_eval_material_queue
+        } else {
+            &self.basic_eval_material_queue
+        };
+
+        let n = queue.size();
+        let next = ((depth + 1) % 2) as usize;
+
+        for i in 0..n as usize {
+            let w = unsafe { queue.storage.get(i) };
+            let pi = w.pixel_index as usize;
+
+            let lambda = self.pixel_sample_state.lambda.get(pi);
+            let beta = self.pixel_sample_state.beta.get(pi);
+            let any_non_specular = self.pixel_sample_state.any_non_specular_bounces.get(pi) != 0;
+            let eta_scale = self.pixel_sample_state.eta_scale.get(pi);
+
+            let Some(material) = w.material.get() else {
+                continue;
+            };
+
+            let tex_eval = UniversalTextureEvaluator;
+            let ctx = MaterialEvalContext {
+                texture: TextureEvalContext {
+                    p: w.p,
+                    dpdx: Vector3f::zero(),
+                    dpdy: Vector3f::zero(),
+                    n: w.n,
+                    uv: w.uv,
+                    dudx: 0.0,
+                    dudy: 0.0,
+                    dvdx: 0.0,
+                    dvdy: 0.0,
+                    face_index: w.face_index,
+                },
+                wo: w.wo,
+                ns: w.ns,
+                dpdus: w.dpdu,
+            };
+            let mut bsdf = material.get_bsdf(&tex_eval, &ctx, &lambda);
+
+            if bsdf.flags().is_empty() {
+                continue;
+            }
+
+            if self.regularize && any_non_specular {
+                bsdf.regularize();
+            }
+
+            if depth >= self.max_depth {
+                continue;
+            }
+
+            // Sample a light, compute contribution,
+            // push shadow ray with deferred visibility
+            if bsdf.flags().is_non_specular() {
+                let light_ctx = LightSampleContext {
+                    pi: Point3fi::new_from_point(w.p),
+                    n: w.n,
+                    ns: w.ns,
+                };
+
+                if let Some(sampled_light) = self
+                    .light_sampler
+                    .sample_with_context(&light_ctx, self.sampler.get1d())
+                {
+                    if let Some(ls) = sampled_light.light.sample_li(
+                        &light_ctx,
+                        self.sampler.get2d(),
+                        &lambda,
+                        true,
+                    ) {
+                        if !ls.l.is_black() && ls.pdf > 0.0 {
+                            let wi = ls.wi;
+                            if let Some(f_val) = bsdf.f(w.wo, wi, TransportMode::Radiance) {
+                                let f_cos = f_val * wi.abs_dot(w.ns.into());
+                                if !f_cos.is_black() {
+                                    let p_l = sampled_light.p * ls.pdf;
+                                    let l_d = if sampled_light.light.light_type().is_delta_light() {
+                                        beta * ls.l * f_cos / p_l
+                                    } else {
+                                        let p_b = bsdf.pdf(w.wo, wi, FArgs::default());
+                                        let w_l = power_heuristic(1, p_l, 1, p_b);
+                                        beta * w_l * ls.l * f_cos / p_l
+                                    };
+
+                                    if !l_d.is_black() {
+                                        let ray_o = Ray::offset_origin(
+                                            &Point3fi::new_from_point(w.p),
+                                            &w.n,
+                                            &wi,
+                                        );
+                                        let t_max = (1.0 - 1e-4)
+                                            * (Point3f::from(ls.p_light.p()) - ray_o).norm()
+                                            / wi.norm();
+
+                                        self.shadow_ray_queue.push(ShadowRayWorkItem {
+                                            ray_o,
+                                            ray_d: wi,
+                                            ray_time: w.time,
+                                            t_max,
+                                            lambda,
+                                            l_d,
+                                            pixel_index: w.pixel_index,
+                                        });
+                                    }
+                                }
+                            }
+                        }
+                    }
+                }
+            }
+
+            // Sample BSDF for next bounce
+            let wo = w.wo;
+            let Some(bs) = bsdf.sample_f(
+                wo,
+                self.sampler.get1d(),
+                self.sampler.get2d(),
+                FArgs::default(),
+            ) else {
+                continue;
+            };
+
+            let f_cos = bs.f * bs.wi.abs_dot(w.ns.into());
+            if f_cos.is_black() || bs.pdf == 0.0 {
+                continue;
+            }
+            let new_beta = beta * f_cos / bs.pdf;
+
+            let new_depth = depth + 1;
+
+            // Russian roulette
+            if new_depth > 3 {
+                let rr_beta = new_beta.max_component_value();
+                if rr_beta < 0.25 {
+                    let q = (1.0 - rr_beta).max(0.0_f32);
+                    if self.sampler.get1d() < q {
+                        continue;
+                    }
+                }
+            }
+
+            let ray_o = Ray::offset_origin(&Point3fi::new_from_point(w.p), &w.n, &bs.wi);
+
+            // Update PixelSampleState
+            self.pixel_sample_state.beta.set(pi, new_beta);
+            self.pixel_sample_state.depth.set(pi, new_depth);
+            self.pixel_sample_state
+                .specular_bounce
+                .set(pi, bs.is_specular() as u8);
+            self.pixel_sample_state
+                .any_non_specular_bounces
+                .set(pi, (any_non_specular || !bs.is_specular()) as u8);
+            self.pixel_sample_state.eta_scale.set(
+                pi,
+                if bs.is_transmissive() {
+                    eta_scale * square(bs.eta)
+                } else {
+                    eta_scale
+                },
+            );
+            self.pixel_sample_state.prev_intr_ctx.set(
+                pi,
+                LightSampleContext {
+                    pi: Point3fi::new_from_point(w.p),
+                    n: w.n,
+                    ns: w.ns,
+                },
+            );
+
+            // Push next bounce ray
+            self.ray_queues[next].push(RayWorkItem {
+                ray_o,
+                ray_d: bs.wi,
+                ray_time: w.time,
+                ray_medium: Ptr::null(),
+                pixel_index: w.pixel_index,
+                has_differentials: true,
+                differential: RayDifferential::default(),
+            });
+        }
+    }
+
+    fn update_film(&self, y0: i32, y1: i32, pixel_bounds: &Bounds2i) {
+        // The pixel_sample_state indices map to rays generated in
+        // generate_camera_rays. We need to walk the same pixel order
+        // and read back the accumulated L values.
+        let mut pi = 0usize;
+        for y in y0..y1 {
+            for x in pixel_bounds.p_min.x()..pixel_bounds.p_max.x() {
+                let l = self.pixel_sample_state.l.get(pi);
+                let lambda = self.pixel_sample_state.lambda.get(pi);
+                let filter_weight = self.pixel_sample_state.filter_weight.get(pi);
+                let p_film = self.pixel_sample_state.p_film.get(pi);
+
+                // Add sample to film
+                self.film.add_sample(
+                    Point2i::new(p_film.x() as i32, p_film.y() as i32),
+                    l,
+                    &lambda,
+                    Some(&crate::core::film::VisibleSurface::default()),
+                    filter_weight,
+                );
+
+                pi += 1;
+            }
+        }
+    }
 }

src/wavefront/mod.rs

@@ -1,4 +1,4 @@
-pub mod integrator;
 pub mod aggregate;
+pub mod integrator;
 
-pub use aggregate::WavefrontAggregate;
+pub use aggregate::CpuAggregate;