Finally fixed import errors due to refactoring, gotta get a better IDE

.gitignore

@@ -6,3 +6,6 @@ flip.rs
 .vscode
 rust-analyzer.json
 data/
+src/gpu/
+src/tests/
+tests/

Cargo.toml

@@ -8,7 +8,7 @@ default = []
 use_f64 = []
 use_gpu = []
 use_nvtx = []
-cuda = ["cust", "cuda_builder", "shared/cuda", ]
+cuda = ["dep:cudarc"]
 
 [dependencies]
 anyhow = "1.0.100"
@@ -31,12 +31,13 @@ unicode-normalization = "0.1.25"
 wgpu = "27.0.1"
 
 shared = { path = "shared" }
-kernels = { path = "kernels" }
+ptex-filter = { path = "crates/ptex-filter" }
+# kernels = { path = "kernels" }
 
 cuda_std = { git = "https://github.com/Rust-GPU/Rust-CUDA", branch = "main", default-features = false, optional = true }
 cust = { git = "https://github.com/Rust-GPU/Rust-CUDA", branch = "main", default-features = false, features = ["glam"], optional = true }
 ptex = "0.3.0"
-ptex-sys = "0.3.0"
+# ptex-sys = "0.3.0"
 slice = "0.0.4"
 crossbeam-channel = "0.5.15"
 num_cpus = "1.17.0"
@@ -46,13 +47,15 @@ bytemuck = "1.24.0"
 once_cell = "1.21.3"
 smallvec = "1.15.1"
 cuda-runtime-sys = "0.3.0-alpha.1"
+cudarc = { version = "0.18.2", features = ["cuda-13000"], optional = true }
 
 [build-dependencies]
 spirv-builder = { git = "https://github.com/rust-gpu/rust-gpu", branch = "main", optional = true }
 cuda_builder = { git = "https://github.com/Rust-GPU/Rust-CUDA", branch = "main", optional = true }
+cc = "1.2.53"
 
 [workspace]
-members = ["kernels", "shared"]
+members = ["shared", "crates/ptex-filter"]
 
 [lints.clippy]
 excessive_precision = "allow"

build.rs

@@ -0,0 +1,43 @@
+use std::process::Command;
+
+fn main() {
+    println!("cargo:rerun-if-changed=kernels/");
+
+    if std::env::var("CARGO_FEATURE_CUDA").is_ok() {
+        compile_kernels();
+    }
+}
+
+fn compile_kernels() {
+    let out_dir = std::env::var("OUT_DIR").unwrap();
+
+    let kernels = ["test_kernels"];
+
+    for name in kernels {
+        let src = format!("kernels/{}.cu", name);
+        let dst = format!("{}/{}.ptx", out_dir, name);
+
+        println!("cargo:rerun-if-changed={}", src);
+
+        let status = Command::new("nvcc")
+            .args([
+                "-ptx",
+                "-o",
+                &dst,
+                &src,
+                "--gpu-architecture=sm_75", // Adjust for your GPU
+                "-O3",
+                "--use_fast_math",
+            ])
+            .status()
+            .expect("Failed to run nvcc");
+
+        if !status.success() {
+            panic!("nvcc failed on {}", src);
+        }
+
+        println!("cargo:warning=Compiled {} -> {}", src, dst);
+    }
+
+    println!("cargo:rustc-env=KERNEL_PTX_DIR={}", out_dir);
+}

crates/ptex-filter/Cargo.toml

@@ -0,0 +1,11 @@
+[package]
+name = "ptex-filter"
+version = "0.1.0"
+edition = "2021"
+
+[build-dependencies]
+cc = "1.0"
+pkg-config = "0.3"
+
+[dependencies]
+ptex = "0.3.0"

crates/ptex-filter/build.rs

@@ -0,0 +1,17 @@
+fn main() {
+    println!("cargo:rerun-if-changed=cpp/ptex_filter_wrapper.cpp");
+
+    let home = std::env::var("HOME").unwrap();
+    let ptex_include = format!("{}/.local/include", home);
+    let ptex_lib = format!("{}/.local/lib", home);
+
+    cc::Build::new()
+        .cpp(true)
+        .file("cpp/ptex_filter_wrapper.cpp")
+        .include(&ptex_include)
+        .flag("-std=c++17")
+        .compile("ptex_filter_wrapper");
+
+    println!("cargo:rustc-link-search=native={}", ptex_lib);
+    println!("cargo:rustc-link-lib=Ptex");
+}

crates/ptex-filter/cpp/ptex_filter_wrapper.cpp

@@ -0,0 +1,42 @@
+#include "ptex_filter_wrapper.h"
+#include <Ptexture.h>
+
+extern "C" {
+
+PtexFilterHandle ptex_filter_create(PtexTextureHandle texture, const PtexFilterOptions* opts) {
+    Ptex::PtexTexture* tex = static_cast<Ptex::PtexTexture*>(texture);
+    if (!tex || !opts) return nullptr;
+
+    Ptex::PtexFilter::Options ptex_opts;
+    ptex_opts.filter = static_cast<Ptex::PtexFilter::FilterType>(opts->filter);
+    ptex_opts.lerp = opts->lerp;
+    ptex_opts.sharpness = opts->sharpness;
+    ptex_opts.noedgeblend = opts->noedgeblend != 0;
+
+    return Ptex::PtexFilter::getFilter(tex, ptex_opts);
+}
+
+void ptex_filter_eval(
+    PtexFilterHandle filter,
+    float* result,
+    int32_t first_channel,
+    int32_t num_channels,
+    int32_t face_id,
+    float u, float v,
+    float dudx, float dvdx,
+    float dudy, float dvdy
+) {
+    Ptex::PtexFilter* f = static_cast<Ptex::PtexFilter*>(filter);
+    if (f && result) {
+        f->eval(result, first_channel, num_channels, face_id, u, v, dudx, dvdx, dudy, dvdy);
+    }
+}
+
+void ptex_filter_release(PtexFilterHandle filter) {
+    Ptex::PtexFilter* f = static_cast<Ptex::PtexFilter*>(filter);
+    if (f) {
+        f->release();
+    }
+}
+
+}

crates/ptex-filter/cpp/ptex_filter_wrapper.h

@@ -0,0 +1,47 @@
+#pragma once
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#include <stdint.h>
+
+typedef void* PtexTextureHandle;
+typedef void* PtexFilterHandle;
+
+typedef enum {
+    PTEX_FILTER_POINT = 0,
+    PTEX_FILTER_BILINEAR = 1,
+    PTEX_FILTER_BOX = 2,
+    PTEX_FILTER_GAUSSIAN = 3,
+    PTEX_FILTER_BICUBIC = 4,
+    PTEX_FILTER_BSPLINE = 5,
+    PTEX_FILTER_CATMULLROM = 6,
+    PTEX_FILTER_MITCHELL = 7
+} PtexFilterType;
+
+typedef struct {
+    PtexFilterType filter;
+    int32_t lerp;
+    float sharpness;
+    int32_t noedgeblend;
+} PtexFilterOptions;
+
+PtexFilterHandle ptex_filter_create(PtexTextureHandle texture, const PtexFilterOptions* opts);
+
+void ptex_filter_eval(
+    PtexFilterHandle filter,
+    float* result,
+    int32_t first_channel,
+    int32_t num_channels,
+    int32_t face_id,
+    float u, float v,
+    float dudx, float dvdx,
+    float dudy, float dvdy
+);
+
+void ptex_filter_release(PtexFilterHandle filter);
+
+#ifdef __cplusplus
+}
+#endif

crates/ptex-filter/src/ffi.rs

@@ -0,0 +1,54 @@
+use std::ffi::c_void;
+
+#[repr(C)]
+#[derive(Debug, Clone, Copy, PartialEq, Eq)]
+pub enum PtexFilterType {
+    Point = 0,
+    Bilinear = 1,
+    Box = 2,
+    Gaussian = 3,
+    Bicubic = 4,
+    BSpline = 5,
+    CatmullRom = 6,
+    Mitchell = 7,
+}
+
+#[repr(C)]
+#[derive(Debug, Clone)]
+pub struct PtexFilterOptions {
+    pub filter: PtexFilterType,
+    pub lerp: i32,
+    pub sharpness: f32,
+    pub noedgeblend: i32,
+}
+
+impl Default for PtexFilterOptions {
+    fn default() -> Self {
+        Self {
+            filter: PtexFilterType::BSpline,
+            lerp: 1,
+            sharpness: 0.0,
+            noedgeblend: 0,
+        }
+    }
+}
+
+extern "C" {
+    pub fn ptex_filter_create(texture: *mut c_void, opts: *const PtexFilterOptions) -> *mut c_void;
+
+    pub fn ptex_filter_eval(
+        filter: *mut c_void,
+        result: *mut f32,
+        first_channel: i32,
+        num_channels: i32,
+        face_id: i32,
+        u: f32,
+        v: f32,
+        dudx: f32,
+        dvdx: f32,
+        dudy: f32,
+        dvdy: f32,
+    );
+
+    pub fn ptex_filter_release(filter: *mut c_void);
+}

crates/ptex-filter/src/lib.rs

@@ -0,0 +1,60 @@
+mod ffi;
+
+pub use ffi::{ptex_filter_create, PtexFilterOptions, PtexFilterType};
+
+use std::ffi::c_void;
+use std::marker::PhantomData;
+use std::ptr::NonNull;
+
+pub struct PtexFilter {
+    handle: NonNull<c_void>,
+    _marker: PhantomData<*mut ()>,
+}
+
+impl PtexFilter {
+    pub unsafe fn new(texture_ptr: *mut c_void, opts: &PtexFilterOptions) -> Option<Self> {
+        let handle = ffi::ptex_filter_create(texture_ptr, opts);
+        NonNull::new(handle).map(|h| Self {
+            handle: h,
+            _marker: PhantomData,
+        })
+    }
+
+    pub fn eval(
+        &self,
+        face_id: i32,
+        u: f32,
+        v: f32,
+        dudx: f32,
+        dvdx: f32,
+        dudy: f32,
+        dvdy: f32,
+        num_channels: i32,
+    ) -> [f32; 4] {
+        let mut result = [0.0f32; 4];
+        unsafe {
+            ffi::ptex_filter_eval(
+                self.handle.as_ptr(),
+                result.as_mut_ptr(),
+                0,
+                num_channels.min(4),
+                face_id,
+                u,
+                v,
+                dudx,
+                dvdx,
+                dudy,
+                dvdy,
+            );
+        }
+        result
+    }
+}
+
+impl Drop for PtexFilter {
+    fn drop(&mut self) {
+        unsafe {
+            ffi::ptex_filter_release(self.handle.as_ptr());
+        }
+    }
+}

kernels/src/lib.rs

@@ -1,16 +1,30 @@
-#![no_std]
+#![cfg_attr(target_arch = "nvptx64", no_std)]
+#![cfg_attr(target_arch = "nvptx64", feature(abi_ptx))]
 
 use cuda_std::prelude::*;
-use shared::Float;
 
+/// Scales each element: data[i] *= scale
 #[kernel]
-pub unsafe fn scale_array_kernel(data: *mut Float, n: u32, scale: Float) {
-    let i = thread::index_1d() as usize;
-
-    if i < n as usize {
-        let ptr = unsafe { data.add(i) };
-        unsafe {
-            *ptr = *ptr * scale;
-        }
+#[allow(improper_ctypes_definitions)]
+pub unsafe fn scale_array(data: *mut f32, len: u32, scale: f32) {
+    let idx = thread::index_1d() as u32;
+    if idx >= len {
+        return;
     }
+
+    let ptr = unsafe { data.add(idx as usize) };
+    *ptr = *ptr * scale;
+}
+
+/// Adds two arrays: c[i] = a[i] + b[i]
+#[kernel]
+#[allow(improper_ctypes_definitions)]
+pub unsafe fn add_arrays(a: *const f32, b: *const f32, c: *mut f32, len: u32) {
+    let idx = thread::index_1d() as u32;
+    if idx >= len {
+        return;
+    }
+
+    let i = idx as usize;
+    *c.add(i) = *a.add(i) + *b.add(i);
 }

kernels/test_kernels.cu

@@ -0,0 +1,27 @@
+extern "C" __global__ void scale_array(float* data, unsigned int len, float scale) {
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= len) return;
+    data[idx] *= scale;
+}
+
+extern "C" __global__ void add_arrays(
+    const float* __restrict__ a,
+    const float* __restrict__ b,
+    float* __restrict__ c,
+    unsigned int len
+) {
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= len) return;
+    c[idx] = a[idx] + b[idx];
+}
+
+extern "C" __global__ void saxpy(
+    float a,
+    const float* __restrict__ x,
+    float* __restrict__ y,
+    unsigned int len
+) {
+    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= len) return;
+    y[idx] = a * x[idx] + y[idx];
+}

shared/src/bxdfs/complex.rs

@@ -138,11 +138,15 @@ impl HairBxDF {
         std::array::from_fn(|i| ap[i].average() / sum_y)
     }
 
-    pub fn sigma_a_from_concentration(&self, ce: Float, cp: Float) -> RGBUnboundedSpectrum {
+    pub fn sigma_a_from_concentration(
+        ce: Float,
+        cp: Float,
+        stdcs: StandardColorSpaces,
+    ) -> RGBUnboundedSpectrum {
         let eumelanin_sigma_a = RGB::new(0.419, 0.697, 1.37);
         let pheomelanin_sigma_a = RGB::new(0.187, 0.4, 1.05);
         let sigma_a = ce * eumelanin_sigma_a + cp * pheomelanin_sigma_a;
-        RGBUnboundedSpectrum::new(&self.colorspaces.srgb, sigma_a)
+        RGBUnboundedSpectrum::new(&stdcs.srgb, sigma_a)
     }
 }
 

shared/src/core/bssrdf.rs

@@ -94,8 +94,8 @@ impl From<&SubsurfaceInteraction> for SurfaceInteraction {
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]
 pub struct BSSRDFTable {
-    pub n_rho_samples: u32,
-    pub n_radius_samples: u32,
+    pub n_rho: u32,
+    pub n_radius: u32,
     pub rho_samples: Ptr<Float>,
     pub radius_samples: Ptr<Float>,
     pub profile: Ptr<Float>,
@@ -105,34 +105,27 @@ pub struct BSSRDFTable {
 
 impl BSSRDFTable {
     pub fn get_rho(&self) -> &[Float] {
-        unsafe {
-            core::slice::from_raw_parts(self.rho_samples.as_ref(), self.n_rho_samples as usize)
-        }
+        unsafe { core::slice::from_raw_parts(self.rho_samples.as_ref(), self.n_rho as usize) }
     }
 
     pub fn get_radius(&self) -> &[Float] {
-        unsafe {
-            core::slice::from_raw_parts(
-                self.radius_samples.as_ref(),
-                self.n_radius_samples as usize,
-            )
-        }
+        unsafe { core::slice::from_raw_parts(self.radius_samples.as_ref(), self.n_radius as usize) }
     }
 
     pub fn get_profile(&self) -> &[Float] {
-        let n_profile = (self.n_rho_samples * self.n_radius_samples) as usize;
+        let n_profile = (self.n_rho * self.n_radius) as usize;
         unsafe { core::slice::from_raw_parts(self.profile.as_ref(), n_profile) }
     }
 
     pub fn get_cdf(&self) -> &[Float] {
-        let n_profile = (self.n_rho_samples * self.n_radius_samples) as usize;
+        let n_profile = (self.n_rho * self.n_radius) as usize;
         unsafe { core::slice::from_raw_parts(self.profile_cdf.as_ref(), n_profile) }
     }
 
     pub fn eval_profile(&self, rho_index: u32, radius_index: u32) -> Float {
-        debug_assert!(rho_index < self.n_rho_samples);
-        debug_assert!(radius_index < self.n_radius_samples);
-        let idx = (rho_index * self.n_radius_samples + radius_index) as usize;
+        debug_assert!(rho_index < self.n_rho);
+        debug_assert!(radius_index < self.n_radius);
+        let idx = (rho_index * self.n_radius + radius_index) as usize;
         unsafe { *self.profile.add(idx) }
     }
 }

shared/src/core/spectrum.rs

@@ -51,7 +51,7 @@ impl Spectrum {
         XYZ::new(x, y, z) / CIE_Y_INTEGRAL
     }
 
-    fn to_rgb(&self, cs: &RGBColorSpace, std: &StandardSpectra) -> RGB {
+    pub fn to_rgb(&self, cs: &RGBColorSpace, std: &StandardSpectra) -> RGB {
         let xyz = self.to_xyz(std);
         cs.to_rgb(xyz)
     }

shared/src/materials/dielectric.rs

@@ -16,12 +16,12 @@ use crate::utils::math::clamp;
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]
 pub struct DielectricMaterial {
-    normal_map: Ptr<DeviceImage>,
-    displacement: Ptr<GPUFloatTexture>,
-    u_roughness: Ptr<GPUFloatTexture>,
-    v_roughness: Ptr<GPUFloatTexture>,
-    eta: Ptr<Spectrum>,
-    remap_roughness: bool,
+    pub normal_map: Ptr<DeviceImage>,
+    pub displacement: Ptr<GPUFloatTexture>,
+    pub u_roughness: Ptr<GPUFloatTexture>,
+    pub v_roughness: Ptr<GPUFloatTexture>,
+    pub eta: Ptr<Spectrum>,
+    pub remap_roughness: bool,
 }
 
 impl MaterialTrait for DielectricMaterial {
@@ -87,6 +87,7 @@ pub struct ThinDielectricMaterial {
     pub normal_map: Ptr<DeviceImage>,
     pub eta: Ptr<Spectrum>,
 }
+
 impl MaterialTrait for ThinDielectricMaterial {
     fn get_bsdf<T: TextureEvaluator>(
         &self,

shared/src/shapes/bilinear.rs

@@ -7,7 +7,7 @@ use crate::core::pbrt::{Float, gamma};
 use crate::core::shape::{Shape, ShapeIntersection, ShapeSample, ShapeSampleContext, ShapeTrait};
 use crate::utils::Transform;
 use crate::utils::math::{SquareMatrix, clamp, difference_of_products, lerp, quadratic};
-use crate::utils::mesh::BilinearPatchMesh;
+use crate::utils::mesh::DeviceBilinearPatchMesh;
 use crate::utils::sampling::{
     bilinear_pdf, invert_spherical_rectangle_sample, sample_bilinear, sample_spherical_rectangle,
 };
@@ -46,7 +46,7 @@ impl BilinearIntersection {
 #[repr(C)]
 #[derive(Debug, Clone, Copy)]
 pub struct BilinearPatchShape {
-    pub mesh: BilinearPatchMesh,
+    pub mesh: DeviceBilinearPatchMesh,
     pub blp_index: u32,
     pub area: Float,
     pub rectangle: bool,
@@ -54,7 +54,7 @@ pub struct BilinearPatchShape {
 
 impl BilinearPatchShape {
     pub const MIN_SPHERICAL_SAMPLE_AREA: Float = 1e-4;
-    fn mesh(&self) -> BilinearPatchMesh {
+    fn mesh(&self) -> DeviceBilinearPatchMesh {
         self.mesh
     }
 
@@ -117,7 +117,7 @@ impl BilinearPatchShape {
     }
 
     #[cfg(not(target_os = "cuda"))]
-    pub fn new(mesh: BilinearPatchMesh, blp_index: u32) -> Self {
+    pub fn new(mesh: DeviceBilinearPatchMesh, blp_index: u32) -> Self {
         let mut bp = BilinearPatchShape {
             mesh,
             blp_index,

shared/src/shapes/triangle.rs

@@ -10,7 +10,7 @@ use crate::core::interaction::{
 use crate::core::pbrt::gamma;
 use crate::core::shape::{ShapeIntersection, ShapeSample, ShapeSampleContext, ShapeTrait};
 use crate::utils::math::{difference_of_products, square};
-use crate::utils::mesh::TriangleMesh;
+use crate::utils::mesh::DeviceTriangleMesh;
 use crate::utils::sampling::{
     bilinear_pdf, invert_spherical_triangle_sample, sample_bilinear, sample_spherical_triangle,
     sample_uniform_triangle,
@@ -34,7 +34,7 @@ impl TriangleIntersection {
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]
 pub struct TriangleShape {
-    pub mesh: TriangleMesh,
+    pub mesh: DeviceTriangleMesh,
     pub tri_index: u32,
 }
 
@@ -111,11 +111,11 @@ impl TriangleShape {
         }
     }
 
-    pub fn new(mesh: TriangleMesh, tri_index: u32) -> Self {
+    pub fn new(mesh: DeviceTriangleMesh, tri_index: u32) -> Self {
         Self { mesh, tri_index }
     }
 
-    pub fn get_mesh(&self) -> TriangleMesh {
+    pub fn get_mesh(&self) -> DeviceTriangleMesh {
         self.mesh
     }
 

shared/src/spectra/colorspace.rs

@@ -16,6 +16,50 @@ pub struct StandardColorSpaces {
     pub aces2065_1: Ptr<RGBColorSpace>,
 }
 
+impl StandardColorSpaces {
+    #[cfg(not(target_arch = "nvptx64"))]
+    pub fn get_named(&self, name: &str) -> Option<Ptr<RGBColorSpace>> {
+        match name.to_lowercase().as_str() {
+            "srgb" => Some(self.srgb),
+            "dci-p3" => Some(self.dci_p3),
+            "rec2020" => Some(self.rec2020),
+            "aces2065-1" => Some(self.aces2065_1),
+            _ => None,
+        }
+    }
+
+    pub fn get_by_id(&self, id: ColorSpaceId) -> Ptr<RGBColorSpace> {
+        match id {
+            ColorSpaceId::SRGB => self.srgb,
+            ColorSpaceId::DciP3 => self.dci_p3,
+            ColorSpaceId::Rec2020 => self.rec2020,
+            ColorSpaceId::Aces2065_1 => self.aces2065_1,
+        }
+    }
+}
+
+#[repr(u8)]
+#[derive(Copy, Clone, Debug, PartialEq, Eq)]
+pub enum ColorSpaceId {
+    SRGB = 0,
+    DciP3 = 1,
+    Rec2020 = 2,
+    Aces2065_1 = 3,
+}
+
+impl ColorSpaceId {
+    #[cfg(not(target_arch = "nvptx64"))]
+    pub fn from_name(name: &str) -> Option<Self> {
+        match name.to_lowercase().as_str() {
+            "srgb" => Some(Self::SRGB),
+            "dci-p3" => Some(Self::DciP3),
+            "rec2020" => Some(Self::Rec2020),
+            "aces2065-1" => Some(Self::Aces2065_1),
+            _ => None,
+        }
+    }
+}
+
 #[repr(C)]
 #[derive(Debug, Clone, Copy)]
 pub struct RGBColorSpace {

shared/src/spectra/sampled.rs

@@ -32,6 +32,10 @@ impl SampledSpectrum {
         }
     }
 
+    pub fn zero() -> Self {
+        Self::default()
+    }
+
     #[inline(always)]
     pub fn from_fn<F>(cb: F) -> Self
     where

shared/src/utils/mesh.rs

@@ -6,7 +6,7 @@ use crate::utils::sampling::DevicePiecewiseConstant2D;
 
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
-pub struct TriangleMesh {
+pub struct DeviceTriangleMesh {
     pub n_triangles: u32,
     pub n_vertices: u32,
     pub vertex_indices: Ptr<u32>,
@@ -21,14 +21,19 @@ pub struct TriangleMesh {
 
 #[repr(C)]
 #[derive(Debug, Clone, Copy)]
-pub struct BilinearPatchMesh {
-    pub reverse_orientation: bool,
-    pub transform_swaps_handedness: bool,
+pub struct DeviceBilinearPatchMesh {
     pub n_patches: u32,
     pub n_vertices: u32,
     pub vertex_indices: Ptr<u32>,
     pub p: Ptr<Point3f>,
     pub n: Ptr<Normal3f>,
     pub uv: Ptr<Point2f>,
+    pub reverse_orientation: bool,
+    pub transform_swaps_handedness: bool,
     pub image_distribution: Ptr<DevicePiecewiseConstant2D>,
 }
+
+unsafe impl Send for DeviceTriangleMesh {}
+unsafe impl Sync for DeviceTriangleMesh {}
+unsafe impl Send for DeviceBilinearPatchMesh {}
+unsafe impl Sync for DeviceBilinearPatchMesh {}

shared/src/utils/sampling.rs

@@ -772,7 +772,7 @@ impl DevicePiecewiseConstant1D {
 #[repr(C)]
 #[derive(Debug, Copy, Clone)]
 pub struct DevicePiecewiseConstant2D {
-    pub conditional: *const DevicePiecewiseConstant1D, // Array of n_v conditionals
+    pub conditionals: *const DevicePiecewiseConstant1D, // Array of n_v conditionals
     pub marginal: DevicePiecewiseConstant1D,
     pub n_u: u32,
     pub n_v: u32,
@@ -792,7 +792,7 @@ impl DevicePiecewiseConstant2D {
 
     pub fn sample(&self, u: Point2f) -> (Point2f, f32, Point2i) {
         let (d1, pdf1, off_y) = self.marginal.sample(u.y());
-        let (d0, pdf0, off_x) = (unsafe { &*self.conditional.add(off_y) }).sample(u.x());
+        let (d0, pdf0, off_x) = (unsafe { &*self.conditionals.add(off_y) }).sample(u.x());
         let pdf = pdf0 * pdf1;
         let offset = Point2i::new(off_x as i32, off_y as i32);
         (Point2f::new(d0, d1), pdf, offset)
@@ -803,7 +803,7 @@ impl DevicePiecewiseConstant2D {
         // let delta_v = 1.0 / self.n_v as Float;
         let v_offset = ((p.y() * self.n_v as Float) as usize).min(self.n_v as usize - 1);
 
-        let conditional = unsafe { &*self.conditional.add(v_offset) };
+        let conditional = unsafe { &*self.conditionals.add(v_offset) };
 
         // Find which column
         // let delta_u = 1.0 / self.n_u as Float;

src/core/aggregates.rs

@@ -1,7 +1,7 @@
-use crate::core::primitive::PrimitiveTrait;
 use rayon::prelude::*;
 use shared::Float;
 use shared::core::geometry::{Bounds3f, Point3f, Ray, Vector3f};
+use shared::core::primitive::PrimitiveTrait;
 use shared::core::shape::ShapeIntersection;
 use shared::utils::math::encode_morton_3;
 use shared::utils::{find_interval, partition_slice};

src/core/bssrdf.rs

@@ -1,3 +1,6 @@
+use shared::Float;
+use shared::core::bssrdf::BSSRDFTable;
+
 pub struct BSSRDFTableData {
     pub rho_samples: Vec<Float>,
     pub radius_samples: Vec<Float>,
@@ -22,13 +25,15 @@ impl BSSRDFTableData {
         }
     }
 
-    pub fn view(&self, rho_ptr: *const f32, radius_ptr: *const f32) -> BSSRDFTableView {
+    pub fn view(&self, rho_ptr: *const Float, radius_ptr: *const Float) -> BSSRDFTable {
         BSSRDFTable {
             rho_samples: rho_ptr,
             n_rho: self.rho_samples.len() as u32,
             radius_samples: radius_ptr,
             n_radius: self.radius_samples.len() as u32,
-            // ...
+            profile: self.profile,
+            profile_cdf: self.profile_cdf,
+            rho_eff: self.rho_eff,
         }
     }
 }

src/core/camera.rs

@@ -1,11 +1,18 @@
+use crate::core::image::Image;
 use crate::core::image::ImageMetadata;
+use crate::utils::read_float_file;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
-use shared::Float;
 use shared::cameras::*;
 use shared::core::camera::{Camera, CameraBase, CameraTrait, CameraTransform};
+use shared::core::color::ColorEncoding::SRGB;
 use shared::core::film::Film;
+use shared::core::geometry::{Bounds2f, Point2f, Point2i, Vector2f, Vector3f};
+use shared::core::image::PixelFormat;
 use shared::core::medium::Medium;
 use shared::core::options::get_options;
+use shared::utils::math::square;
+use shared::{Float, PI};
+use std::path::Path;
 use std::sync::Arc;
 
 #[repr(C)]

src/core/color.rs

@@ -1,6 +1,8 @@
+use crate::utils::read_float_file;
 use shared::Float;
 use shared::core::color::{Coeffs, RES, RGBToSpectrumTable};
-use shared::spectra::RGBSigmoidPolynomial;
+use std::ops::Deref;
+use std::path::Path;
 
 pub struct RGBToSpectrumTableData {
     _z_nodes: Vec<Float>,
@@ -33,7 +35,7 @@ impl RGBToSpectrumTableData {
         }
     }
 
-    pub fn load(base_dir: &Path, name: &str) -> io::Result<Self> {
+    pub fn load(base_dir: &Path, name: &str) -> Result<Self> {
         let z_path = base_dir.join(format!("{}_znodes.dat", name));
         let c_path = base_dir.join(format!("{}_coeffs.dat", name));
 

src/core/film.rs

@@ -1,15 +1,33 @@
+use crate::core::image::{Image, ImageChannelDesc, ImageChannelValues, ImageMetadata};
+use crate::spectra::{SRGB, data::get_named_spectrum};
 use shared::Float;
-use shared::core::spectrum::Spectrum;
+use shared::core::camera::CameraTransform;
+use shared::core::color::{RGB, XYZ, white_balance};
+use shared::core::film::SpectralPixel;
+use shared::core::filter::Filter;
+use shared::core::geometry::{Bounds2f, Bounds2i, Point2f, Point2i};
+use shared::core::image::PixelFormat;
+use shared::core::spectrum::{Spectrum, StandardSpectra};
 use shared::film::{Film, FilmBase, GBufferFilm, PixelSensor, PixelSensor, RGBFilm, SpectralFilm};
-use shared::spectra::{PiecewiseLinearSpectrum, RGBColorSpace};
+use shared::spectra::cie::SWATCHES_RAW;
+use shared::spectra::{
+    DenselySampledSpectrum, LAMBDA_MAX, LAMBDA_MIN, PiecewiseLinearSpectrum, RGBColorSpace,
+};
+use shared::utils::containers::Array2D;
+use shared::utils::math::{SquareMatrix, linear_least_squares};
+use shared::utils::{AnimatedTransform, AtomicFloat};
+use std::cmp::Ordering;
+use std::path::Path;
+use std::sync::atomic::AtomicUsize;
+use std::sync::{Arc, LazyLock};
 
 use crate::spectra::DenselySampledSpectrumBuffer;
 use crate::utils::{FileLoc, ParameterDictionary};
 
 const N_SWATCH_REFLECTANCES: usize = 24;
-const SWATCH_REFLECTANCES: Lazy<[Spectrum; N_SWATCH_REFLECTANCES]> = Lazy::new(|| {
+const SWATCH_REFLECTANCES: LazyLock<[Spectrum; N_SWATCH_REFLECTANCES]> = LazyLock::new(|| {
     std::array::from_fn(|i| {
-        let raw_data = crate::core::cie::SWATCHES_RAW[i];
+        let raw_data = SWATCHES_RAW[i];
         let pls = PiecewiseLinearSpectrum::from_interleaved(raw_data, false);
         Spectrum::PiecewiseLinear(pls)
     })
@@ -248,10 +266,10 @@ impl SpectralFilmHost {
             filter_integral: base.filter.integral(),
             output_rgbf_from_sensor_rgb: SquareMatrix::identity(), // Logic omitted
 
-            pixels: Array2DView {
-                data: storage.pixels.as_mut_ptr(),
+            pixels: Array2D {
+                values: storage.pixels.as_mut_ptr(),
                 extent: base.pixel_bounds,
-                x_stride: base.pixel_bounds.max.x - base.pixel_bounds.min.x,
+                stride: base.pixel_bounds.max.x - base.pixel_bounds.min.x,
             },
 
             bucket_sums: storage.bucket_sums.as_ptr() as *mut f64,
@@ -266,6 +284,10 @@ impl SpectralFilmHost {
     }
 }
 
+pub struct GBufferFilmHost {
+    pub device: GBufferFilm,
+}
+
 impl GBufferFilmHost {
     pub fn new(
         base: &FilmBase,
@@ -285,7 +307,7 @@ impl GBufferFilmHost {
         let filter_integral = base.filter.integral();
         let pixels = Array2D::new(base.pixel_bounds);
 
-        Self {
+        let device = GBufferFilm {
             base: base.clone(),
             output_from_render: output_from_render.clone(),
             apply_inverse,
@@ -295,7 +317,9 @@ impl GBufferFilmHost {
             write_fp16,
             filter_integral,
             output_rgbf_from_sensor_rgb,
-        }
+        };
+
+        Self { device }
     }
 }
 
@@ -495,7 +519,7 @@ impl FilmFactory for Film {
                 let film_base = FilmBase::create(params, filter, Some(sensor), loc);
 
                 let filename = params.get_one_string("filename", "pbrt.exr");
-                if Path::new(&ilename).extension() != Some("exr".as_ref()) {
+                if Path::new(&filename).extension() != Some("exr".as_ref()) {
                     return Err(format!(
                         "{}: EXR is the only format supported by GBufferFilm",
                         loc

src/core/filter.rs

@@ -1,14 +1,17 @@
-use shared::core::filter::FilterSampler;
-use shared::core::geometry::Point2f;
-use shared::filter::Filter;
+use crate::utils::sampling::PiecewiseConstant2D;
+use crate::utils::{FileLoc, ParameterDictionary};
+use shared::Float;
+use shared::core::filter::{Filter, FilterSampler};
+use shared::core::geometry::{Bounds2f, Point2f, Vector2f};
 use shared::filters::*;
+use shared::utils::containers::Array2D;
 
 pub trait FilterFactory {
     fn create(name: &str, params: &ParameterDictionary, loc: &FileLoc) -> Result<Filter, String>;
 }
 
 impl FilterFactory for Filter {
-    fn create(name: &str, params: ParameterDictionary, loc: &FileLoc) -> Result<Self, String> {
+    fn create(name: &str, params: &ParameterDictionary, loc: &FileLoc) -> Result<Self, String> {
         match name {
             "box" => {
                 let xw = params.get_one_float("xradius", 0.5);
@@ -78,7 +81,7 @@ impl CreateFilterSampler for FilterSampler {
                 f[(x as i32, y as i32)] = func(p);
             }
         }
-        let distrib = DevicePiecewiseConstant2D::new_with_bounds(&f, domain);
+        let distrib = PiecewiseConstant2D::new_with_bounds(&f, domain);
         Self { domain, f, distrib }
     }
 }

src/core/image/io.rs

@@ -1,11 +1,13 @@
-use super::{Image, ImageAndMetadata};
-use crate::core::image::PixelStorage;
+use super::{Image, ImageAndMetadata, ImageMetadata};
+use crate::core::image::{PixelStorage, WrapMode};
 use crate::utils::error::ImageError;
 use anyhow::{Context, Result, bail};
 use exr::prelude::{read_first_rgba_layer_from_file, write_rgba_file};
 use image_rs::{DynamicImage, ImageReader};
 use shared::Float;
-use shared::core::color::{ColorEncoding, LINEAR, SRGB};
+use shared::core::color::{ColorEncoding, LINEAR};
+use shared::core::geometry::Point2i;
+use shared::core::image::PixelFormat;
 use std::fs::File;
 use std::io::{BufRead, BufReader, BufWriter, Read, Write};
 use std::path::Path;

src/core/image/mod.rs

@@ -1,8 +1,15 @@
 use half::f16;
-use shared::core::geometry::Point2i;
-use shared::core::image::{DeviceImage, ImageAccess, ImageBase, PixelFormat, WrapMode};
-use smallvec::smallvec;
-use std::ops::Deref;
+use shared::Float;
+use shared::core::color::ColorEncoding;
+use shared::core::color::LINEAR;
+use shared::core::geometry::{Bounds2f, Point2f, Point2i};
+use shared::core::image::{
+    DeviceImage, ImageAccess, ImageBase, PixelFormat, Pixels, WrapMode, WrapMode2D,
+};
+use shared::utils::containers::Array2D;
+use shared::utils::math::square;
+use smallvec::{SmallVec, smallvec};
+use std::ops::{Deref, DerefMut};
 
 pub mod io;
 pub mod metadata;
@@ -542,7 +549,7 @@ impl Image {
     }
 
     pub fn has_any_infinite_pixels(&self) -> bool {
-        if format == PixelFormat::Float {
+        if self.format() == PixelFormat::Float {
             return false;
         }
 

src/core/image/ops.rs

@@ -1,7 +1,8 @@
 use super::Image;
-use crate::core::image::PixelStorage;
+use crate::core::image::pixel::PixelStorage;
 use rayon::prelude::*;
 use shared::Float;
+use shared::core::color::ColorEncoding;
 use shared::core::geometry::{Bounds2i, Point2i};
 use shared::core::image::{PixelFormat, WrapMode, WrapMode2D};
 use shared::utils::math::windowed_sinc;
@@ -242,7 +243,7 @@ fn copy_rect_in_kernel<T: PixelStorage>(
     dst: &mut [T],
     res: Point2i,
     channels: usize,
-    enc: crate::spectra::color::ColorEncoding,
+    enc: ColorEncoding,
     extent: Bounds2i,
     buf: &[Float],
 ) {

src/core/image/pixel.rs

@@ -1,5 +1,5 @@
-use crate::core::pbrt::Float;
-use crate::spectra::color::{ColorEncoding, ColorEncodingTrait};
+use shared::Float;
+use shared::core::color::ColorEncoding;
 use half::f16;
 
 // Allows writing generic algorithms that work on any image format.

src/core/light.rs

@@ -7,10 +7,12 @@ use log::error;
 use shared::core::camera::CameraTransform;
 use shared::core::light::Light;
 use shared::core::medium::Medium;
+use shared::core::shape::Shape;
 use shared::core::spectrum::Spectrum;
 use shared::lights::*;
 use shared::spectra::{DenselySampledSpectrum, RGBColorSpace};
 use shared::utils::Transform;
+use std::fmt::Error;
 
 pub fn lookup_spectrum(s: &Spectrum) -> DenselySampledSpectrum {
     let cache = SPECTRUM_CACHE.get_or_init(InternCache::new);
@@ -28,7 +30,7 @@ pub trait CreateLight {
         shape: &Shape,
         alpha_text: &FloatTexture,
         colorspace: Option<&RGBColorSpace>,
-    ) -> Light;
+    ) -> Result<Light, Error>;
 }
 
 pub trait LightFactory {
@@ -60,7 +62,8 @@ impl LightFactory for Light {
         camera_transform: CameraTransform,
     ) -> Result<Self, Error> {
         match name {
-            "diffuse" => lights::diffuse::create(
+            "diffuse" => DiffuseAreaLight::create(
+                name,
                 arena,
                 render_from_light,
                 medium,
@@ -68,7 +71,6 @@ impl LightFactory for Light {
                 loc,
                 shape,
                 alpha_tex,
-                colorspace,
             )?,
             "point" => PointLight::create(
                 arena,

src/core/material.rs

@@ -5,41 +5,26 @@ use crate::utils::error::FileLoc;
 use shared::core::material::Material;
 use shared::materials::*;
 use std::collections::HashMap;
+use std::fmt::Error;
+use std::sync::Arc;
 
 pub trait CreateMaterial: Sized {
     fn create(
         parameters: &TextureParameterDictionary,
-        normal_map: Option<Ptr<Image>>,
+        normal_map: Option<Arc<Image>>,
         named_materials: &HashMap<String, Material>,
         loc: &FileLoc,
         arena: &mut Arena,
-    ) -> Result<Self, Error>;
-}
-
-macro_rules! make_material_factory {
-    (
-        $name:ident, $params:ident, $nmap:ident, $mats:ident, $loc:ident;
-        $($key:literal => $variant:ident($concrete:ty)),+ $(,)?
-    ) => {
-        match $name {
-            $(
-                $key => {
-                    let mat = <$concrete>::create($params, $nmap, $mats, $loc);
-                    Ok(Material::$variant(mat))
-                }
-            )+
-            _ => Err(format!("Material type '{}' unknown at {}", $name, $loc)),
-        }
-    };
+    ) -> Result<Material, Error>;
 }
 
 pub trait MaterialFactory {
     fn create(
         name: &str,
         params: &TextureParameterDictionary,
-        normal_map: Ptr<Image>,
+        normal_map: Option<Arc<Image>>,
         named_materials: HashMap<String, Material>,
-        loc: &FileLoc,
+        loc: FileLoc,
         arena: &mut Arena,
     ) -> Result<Self, Error>;
 }
@@ -47,26 +32,52 @@ pub trait MaterialFactory {
 impl MaterialFactory for Material {
     fn create(
         name: &str,
-        params: &TextureParameterDictionary,
+        parameters: &TextureParameterDictionary,
         normal_map: Option<Arc<Image>>,
-        named_materials: &HashMap<String, Material>,
-        loc: &FileLoc,
+        named_materials: HashMap<String, Material>,
+        loc: FileLoc,
         arena: &mut Arena,
-    ) -> Result<Self, Error> {
-        make_material_factory!(
-            name, params, normal_map, named_materials, loc;
+    ) -> Result<Material, Error> {
+        match name {
+            "diffuse" => {
+                DiffuseMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "coateddiffuse" => {
+                CoatedDiffuseMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "coatedconductor" => CoatedConductorMaterial::create(
+                parameters,
+                normal_map,
+                named_materials,
+                loc,
+                arena,
+            )?,
+            "diffusetransmission" => DiffuseTransmissionMaterial::create(
+                parameters,
+                normal_map,
+                named_materials,
+                loc,
+                arena,
+            )?,
+            "dielectric" => {
+                DielectricMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "thindielectric" => {
+                ThinDielectricMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "hair" => HairMaterial::create(parameters, normal_map, named_materials, loc, arena)?,
+            "conductor" => {
+                ConductorMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "measured" => {
+                MeasuredMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "subsurface" => {
+                SubsurfaceMaterial::create(parameters, normal_map, named_materials, loc, arena)?
+            }
+            "mix" => MixMaterial::create(parameters, normal_map, named_materials, loc, arena)?,
 
-            "diffuse"             => Diffuse(DiffuseMaterial),
-            "coateddiffuse"       => CoatedDiffuse(CoatedDiffuseMaterial),
-            "coatedconductor"     => Conductor(CoatedConductorMaterial),
-            "diffusetransmission" => DiffuseTransmission(DiffuseTransmissionMaterial),
-            "dielectric"          => Dielectric(DielectricMaterial),
-            "thindielectric"      => ThinDielectric(ThinDielectricMaterial),
-            "hair"                => Hair(HairMaterial),
-            "conductor"           => Conductor(ConductorMaterial),
-            "measured"            => Measured(MeasuredMaterial),
-            "subsurface"          => Subsurface(SubsurfaceMaterial),
-            "mix"                 => Mix(MixMaterial)
-        )
+            _ => Err(format!("Material type '{}' unknown at {}", $name, $loc)),
+        }
     }
 }

src/core/medium.rs

@@ -1,7 +1,8 @@
 use crate::spectra::dense::DenselySampledSpectrumBuffer;
-use shared::core::geometry::Bounds3f;
-use shared::core::medium::{GridMedium, HomogeneousMedium, RGBGridMedium};
-use shared::spectra::{RGBIlluminantSpectrum, RGBUnboundedSpectrum};
+use shared::core::geometry::{Bounds3f, Point3i};
+use shared::core::medium::{GridMedium, HGPhaseFunction, HomogeneousMedium, RGBGridMedium};
+use shared::core::spectrum::Spectrum;
+use shared::spectra::{DenselySampledSpectrum, RGBIlluminantSpectrum, RGBUnboundedSpectrum};
 use shared::utils::Transform;
 use shared::utils::containers::SampledGrid;
 use shared::{Float, core::medium::MajorantGrid};
@@ -80,7 +81,7 @@ impl RGBGridMediumCreator for RGBGridMedium {
 }
 
 pub trait GridMediumCreator {
-    pub fn new(
+    fn new(
         bounds: &Bounds3f,
         render_from_medium: &Transform,
         sigma_a: &Spectrum,

src/core/sampler.rs

@@ -1,6 +1,11 @@
-use shared::core::sampler::Sampler;
-
 use crate::Arena;
+use crate::utils::{FileLoc, ParameterDictionary};
+use shared::core::geometry::Point2i;
+use shared::core::sampler::{
+    HaltonSampler, IndependentSampler, PaddedSobolSampler, Sampler, SobolSampler,
+    StratifiedSampler, ZSobolSampler,
+};
+use std::fmt::Error;
 
 pub trait CreateSampler {
     fn create(

src/core/scene.rs

@@ -1,6 +1,6 @@
 use crate::core::filter::FilterFactory;
 use crate::core::image::{Image, io::ImageIO};
-use crate::core::texture::SpectrumTexture;
+use crate::core::material::MaterialFactory;
 use crate::utils::parallel::{AsyncJob, run_async};
 use crate::utils::parameters::{
     NamedTextures, ParameterDictionary, ParsedParameterVector, TextureParameterDictionary,
@@ -12,22 +12,23 @@ use image_rs::Primitive;
 use parking_lot::Mutex;
 use shared::Float;
 use shared::core::camera::{Camera, CameraTransform};
+use shared::core::color::ColorEncoding;
+use shared::textures::*;
 // use shared::core::color::LINEAR;
-use shared::core::film::{Film, FilmTrait};
+use crate::core::texture::{FloatTexture, SpectrumTexture};
+use crate::utils::error::FileLoc;
+use shared::core::film::Film;
 use shared::core::filter::Filter;
 use shared::core::geometry::Vector3f;
-use shared::core::lights::Light;
+use shared::core::light::Light;
 use shared::core::material::Material;
 use shared::core::medium::{Medium, MediumInterface};
 use shared::core::options::RenderingCoordinateSystem;
-use shared::core::primitive::PrimitiveTrait;
+use shared::core::primitive::{GeometricPrimitive, PrimitiveTrait, SimplePrimitive};
 use shared::core::sampler::Sampler;
-use shared::core::spectrum::SpectrumType;
-use shared::core::texture::{FloatTexture, SpectrumTexture};
-use shared::images::Image;
+use shared::core::shape::Shape;
+use shared::core::texture::SpectrumType;
 use shared::spectra::RGBColorSpace;
-use shared::utils::error::FileLoc;
-// use shared::utils::math::SquareMatrix;
 use shared::utils::transform::{AnimatedTransform, Transform, look_at};
 use std::collections::{HashMap, HashSet};
 use std::ops::{Index as IndexTrait, IndexMut as IndexMutTrait};
@@ -106,7 +107,7 @@ pub struct LightSceneEntity {
 #[derive(Clone, Debug)]
 pub enum InstanceTransform {
     Animated(AnimatedTransform),
-    Static(Arc<TransformGeneric<Float>>),
+    Static(Arc<Transform>),
 }
 
 #[derive(Clone, Debug)]
@@ -195,9 +196,9 @@ impl<'a> SceneLookup<'a> {
     fn resolve_material(&self, name: &str, loc: &FileLoc) -> Material {
         if !name.is_empty() {
             *self.named_materials.get(name).expect("Material not found")
-        } else {
-            self.materials[index]
-        }
+        } // else {
+        //     self.materials[index]
+        // }
     }
 }
 
@@ -415,7 +416,7 @@ impl BasicScene {
         let camera_transform = self.get_camera().camera_transform;
 
         let light_clone = light.clone();
-        let mut light_state = self.light_data.lock().unwrap();
+        let mut light_state = self.light_state.lock().unwrap();
         let job = run_async(move || {
             let render_from_light = light_clone
                 .transformed_base
@@ -428,10 +429,10 @@ impl BasicScene {
                 render_from_light,
                 camera_transform,
                 medium,
-                &light_cmd.transformed_base.base.loc,
+                &light_clone.transformed_base.base.loc,
             )
         });
-        light_state.light_jobs.insert(name.to_string(), job);
+        light_state.light_jobs.insert(job);
     }
 
     pub fn add_area_light(&self, light: SceneEntity) -> usize {
@@ -465,6 +466,7 @@ impl BasicScene {
     pub fn create_materials(
         &self,
         textures: &NamedTextures,
+        arena: &mut Arena,
     ) -> (HashMap<String, Material>, Vec<Material>) {
         let mut state = self.material_state.lock().unwrap();
         log::info!(
@@ -482,7 +484,7 @@ impl BasicScene {
             state.normal_map_jobs.insert(filename, image);
         }
 
-        let mut named_materials_out: HashMap<String, Materal> = HashMap::new();
+        let mut named_materials_out: HashMap<String, Material> = HashMap::new();
 
         for (name, entity) in &state.named_materials {
             if named_materials_out.contains_key(name) {
@@ -520,11 +522,12 @@ impl BasicScene {
             let tex_dict = TextureParameterDictionary::new(&entity.parameters, textures);
 
             let mat = Material::create(
-                &mat_type,
+                name,
                 &tex_dict,
                 normal_map_img,
                 &named_materials_out,
                 &entity.loc,
+                arena,
             );
 
             named_materials_out.insert(name.clone(), mat);
@@ -549,11 +552,12 @@ impl BasicScene {
             let tex_dict = TextureParameterDictionary::new(&entity.parameters, textures);
 
             let mat = Material::create(
-                &mat_type,
+                entity.materials.name,
                 &tex_dict,
                 normal_map_img,
                 &named_materials_out,
                 &entity.loc,
+                arena,
             );
         }
         (named_materials_out, materials_out)
@@ -903,7 +907,7 @@ impl BasicScene {
 
         let filename_clone = filename.clone();
 
-        let job = crate::parallel::run_async(move || {
+        let job = run_async(move || {
             let path = std::path::Path::new(&filename_clone);
 
             let immeta = Image::read(path, Some(ColorEncoding::Linear))
@@ -954,7 +958,7 @@ const MAX_TRANSFORMS: usize = 2;
 
 #[derive(Debug, Default, Clone, Copy)]
 struct TransformSet {
-    t: [TransformGeneric<Float>; MAX_TRANSFORMS],
+    t: [Transform; MAX_TRANSFORMS],
 }
 
 impl TransformSet {
@@ -970,7 +974,7 @@ impl TransformSet {
 
     pub fn map<F>(&mut self, bits: u32, f: F)
     where
-        F: Fn(&TransformGeneric<Float>) -> TransformGeneric<Float>,
+        F: Fn(&Transform) -> Transform,
     {
         if (bits & 1) != 0 {
             self.t[0] = f(&self.t[0]);
@@ -982,7 +986,7 @@ impl TransformSet {
 }
 
 impl IndexTrait<usize> for TransformSet {
-    type Output = TransformGeneric<Float>;
+    type Output = Transform;
 
     fn index(&self, index: usize) -> &Self::Output {
         &self.t[index]
@@ -1032,7 +1036,7 @@ pub struct BasicSceneBuilder {
     graphics_state: GraphicsState,
     pushed_graphics_states: Vec<GraphicsState>,
     push_stack: Vec<(char, FileLoc)>,
-    render_from_world: TransformGeneric<Float>,
+    render_from_world: Transform,
     named_coordinate_systems: HashMap<String, TransformSet>,
     active_instance_definition: Option<InstanceDefinitionSceneEntity>,
 
@@ -1061,7 +1065,7 @@ impl BasicSceneBuilder {
             graphics_state: GraphicsState::default(),
             pushed_graphics_states: Vec::new(),
             push_stack: Vec::new(),
-            render_from_world: TransformGeneric::identity(),
+            render_from_world: Transform::identity(),
             named_coordinate_systems: HashMap::new(),
             active_instance_definition: None,
             shapes: Vec::new(),
@@ -1105,7 +1109,7 @@ impl BasicSceneBuilder {
 
     fn for_active_transforms<F>(&mut self, f: F)
     where
-        F: Fn(&TransformGeneric<Float>) -> TransformGeneric<Float>,
+        F: Fn(&Transform) -> Transform,
     {
         let bits = self.graphics_state.active_transform_bits;
 
@@ -1148,21 +1152,21 @@ impl ParserTarget for BasicSceneBuilder {
     }
 
     fn identity(&mut self, _loc: FileLoc) {
-        self.for_active_transforms(|_| TransformGeneric::identity());
+        self.for_active_transforms(|_| Transform::identity());
     }
 
     fn translate(&mut self, dx: Float, dy: Float, dz: Float, _loc: FileLoc) {
-        let t = TransformGeneric::translate(Vector3f::new(dx, dy, dz));
+        let t = Transform::translate(Vector3f::new(dx, dy, dz));
         self.for_active_transforms(|cur| cur * &t);
     }
 
     fn rotate(&mut self, angle: Float, ax: Float, ay: Float, az: Float, _loc: FileLoc) {
-        let t = TransformGeneric::rotate_around_axis(angle, Vector3f::new(ax, ay, az));
+        let t = Transform::rotate_around_axis(angle, Vector3f::new(ax, ay, az));
         self.for_active_transforms(|cur| cur * &t);
     }
 
     fn scale(&mut self, sx: Float, sy: Float, sz: Float, _loc: FileLoc) {
-        let t = TransformGeneric::scale(sx, sy, sz);
+        let t = Transform::scale(sx, sy, sz);
         self.for_active_transforms(|cur| cur * &t);
     }
 
@@ -1191,7 +1195,7 @@ impl ParserTarget for BasicSceneBuilder {
     }
 
     fn concat_transform(&mut self, m: &[Float; 16], loc: FileLoc) {
-        let result = TransformGeneric::from_flat(m);
+        let result = Transform::from_flat(m);
         match result {
             Ok(t) => {
                 self.for_active_transforms(|cur| cur * &t);
@@ -1203,7 +1207,7 @@ impl ParserTarget for BasicSceneBuilder {
     }
 
     fn transform(&mut self, m: &[Float; 16], loc: FileLoc) {
-        let result = TransformGeneric::from_flat(m);
+        let result = Transform::from_flat(m);
         match result {
             Ok(t) => {
                 self.for_active_transforms(|_| t);
@@ -1379,7 +1383,7 @@ impl ParserTarget for BasicSceneBuilder {
         self.verify_options("WorldBegin", &loc);
         self.current_block = BlockState::WorldBlock;
         for i in 0..MAX_TRANSFORMS {
-            self.graphics_state.ctm[i] = TransformGeneric::<Float>::default();
+            self.graphics_state.ctm[i] = Transform::default();
         }
         self.graphics_state.active_transform_bits = Self::ALL_TRANSFORM_BITS;
         self.named_coordinate_systems
@@ -1491,7 +1495,7 @@ impl ParserTarget for BasicSceneBuilder {
                 &loc,
                 &format!(
                     "{}: texture type unknown. Must be \"float\" or \"spectrum\".",
-                    tex_type
+                    tex_name
                 ),
             );
             return;
@@ -1528,9 +1532,14 @@ impl ParserTarget for BasicSceneBuilder {
         }
     }
 
-    fn material(&mut self, _name: &str, _params: &ParsedParameterVector, loc: FileLoc) {
+    fn material(&mut self, name: &str, params: &ParsedParameterVector, loc: FileLoc) {
         self.verify_world("material", loc);
-        self.graphics_state.current_material_name = self.scene.add_material(name, material)
+        let entity = SceneEntity {
+            name,
+            loc,
+            parameters: params,
+        };
+        self.graphics_state.current_material_name = self.scene.add_material(name, entity);
     }
     fn make_named_material(&mut self, _name: &str, _params: &ParsedParameterVector, _loc: FileLoc) {
         todo!()

src/core/shape.rs

@@ -1,19 +1,20 @@
 use crate::core::texture::FloatTexture;
-use crate::shapes::{BilinearPatchMeshHost, TriQuadMesh, TriangleMeshHost};
+use crate::shapes::{BilinearPatchMesh, TriQuadMesh, TriangleMesh};
 use crate::utils::{Arena, FileLoc, ParameterDictionary, resolve_filename};
 use shared::core::options::get_options;
 use shared::core::shape::*;
 use shared::shapes::*;
 // use shared::spectra::*;
+use parking_lot::Mutex;
 use shared::utils::Transform;
-use std::sync::{Arc, Mutex};
+use std::collections::HashMap;
+use std::sync::Arc;
 
-pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<TriangleMeshHost>>> = Mutex::new(Vec::new());
-pub static ALL_BILINEAR_MESHES: Mutex<Vec<Arc<BilinearPatchMeshHost>>> = Mutex::new(Vec::new());
+pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<TriangleMesh>>> = Mutex::new(Vec::new());
+pub static ALL_BILINEAR_MESHES: Mutex<Vec<Arc<BilinearPatchMesh>>> = Mutex::new(Vec::new());
 
 pub trait CreateShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         object_from_render: Transform,
         reverse_orientation: bool,
@@ -21,7 +22,7 @@ pub trait CreateShape {
         float_textures: HashMap<String, FloatTexture>,
         loc: FileLoc,
         arena: &mut Arena,
-    ) -> Vec<Shape>;
+    ) -> Result<Vec<Shape>, String>;
 }
 
 pub trait ShapeFactory {
@@ -34,7 +35,7 @@ pub trait ShapeFactory {
         float_textures: HashMap<String, FloatTexture>,
         loc: FileLoc,
         arena: &mut Arena,
-    ) -> Vec<Shape>;
+    ) -> Result<Vec<Shape>, String>;
 }
 
 impl ShapeFactory for Shape {
@@ -47,10 +48,9 @@ impl ShapeFactory for Shape {
         float_textures: HashMap<String, FloatTexture>,
         loc: FileLoc,
         arena: &mut Arena,
-    ) -> Vec<Shape> {
+    ) -> Result<Vec<Shape>, String> {
         match name {
             "sphere" => SphereShape::create(
-                name,
                 render_from_object,
                 object_from_render,
                 reverse_orientation,
@@ -58,9 +58,8 @@ impl ShapeFactory for Shape {
                 float_textures,
                 loc,
                 arena,
-            ),
+            )?,
             "cylinder" => CylinderShape::create(
-                name,
                 render_from_object,
                 object_from_render,
                 reverse_orientation,
@@ -68,9 +67,8 @@ impl ShapeFactory for Shape {
                 float_textures,
                 loc,
                 arena,
-            ),
+            )?,
             "disk" => DiskShape::create(
-                name,
                 render_from_object,
                 object_from_render,
                 reverse_orientation,
@@ -78,9 +76,8 @@ impl ShapeFactory for Shape {
                 float_textures,
                 loc,
                 arena,
-            ),
+            )?,
             "bilinearmesh" => BilinearPatchShape::create(
-                name,
                 render_from_object,
                 object_from_render,
                 reverse_orientation,
@@ -88,9 +85,8 @@ impl ShapeFactory for Shape {
                 float_textures,
                 loc,
                 arena,
-            ),
+            )?,
             "trianglemesh" => TriangleShape::create(
-                name,
                 render_from_object,
                 object_from_render,
                 reverse_orientation,
@@ -98,7 +94,7 @@ impl ShapeFactory for Shape {
                 float_textures,
                 loc,
                 arena,
-            ),
+            )?,
             "plymesh" => {
                 let filename = resolve_filename(parameters.get_one_string("filename", ""));
                 let ply_mesh = TriQuadMesh::read_ply(filename);

src/core/spectrum.rs

@@ -1,11 +1,14 @@
-use crate::core::light::LightBaseTrait;
 use crate::spectra::cie_y;
 use crate::utils::containers::InternCache;
+use parking_lot::Mutex;
 use shared::Float;
 use shared::core::spectrum::Spectrum;
+use shared::spectra::DenselySampledSpectrum;
+use std::collections::HashMap;
+use std::sync::LazyLock;
 
-pub static SPECTRUM_CACHE: Lazy<Mutex<HashMap<String, Spectrum>>> =
-    Lazy::new(|| Mutex::new(HashMap::new()));
+pub static SPECTRUM_CACHE: LazyLock<Mutex<HashMap<String, Spectrum>>> =
+    LazyLock::new(|| Mutex::new(HashMap::new()));
 
 fn get_spectrum_cache() -> &'static InternCache<DenselySampledSpectrum> {
     SPECTRUM_CACHE.get_or_init(InternCache::new)

src/core/texture.rs

@@ -1,11 +1,30 @@
 use crate::textures::*;
+use crate::utils::mipmap::MIPMap;
 use crate::utils::mipmap::MIPMapFilterOptions;
 use crate::utils::{Arena, FileLoc, TextureParameterDictionary};
 use enum_dispatch::enum_dispatch;
+use shared::Float;
+use shared::core::color::ColorEncoding;
+use shared::core::geometry::Vector3f;
+use shared::core::image::WrapMode;
+use shared::core::texture::{
+    CylindricalMapping, PlanarMapping, SphericalMapping, TextureEvalContext, TextureMapping2D,
+    UVMapping,
+};
+use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use shared::textures::*;
 use shared::utils::Transform;
+use std::collections::HashMap;
 use std::sync::{Arc, Mutex, OnceLock};
 
+pub trait FloatTextureTrait {
+    fn evaluate(&self, ctx: &TextureEvalContext) -> Float;
+}
+
+pub trait SpectrumTextureTrait {
+    fn evaluate(&self, _ctx: &TextureEvalContext, _lambda: &SampledWavelengths) -> SampledSpectrum;
+}
+
 #[enum_dispatch(FloatTextureTrait)]
 #[derive(Debug, Clone)]
 pub enum FloatTexture {
@@ -90,8 +109,8 @@ impl FloatTexture {
 #[derive(Clone, Debug)]
 #[enum_dispatch(SpectrumTextureTrait)]
 pub enum SpectrumTexture {
-    RGBConstant(RGBConstantTexture),
-    RGBReflectanceConstant(RGBReflectanceConstantTexture),
+    // RGBConstant(RGBConstantTexture),
+    // RGBReflectanceConstant(RGBReflectanceConstantTexture),
     Constant(SpectrumConstantTexture),
     Bilerp(SpectrumBilerpTexture),
     Checkerboard(SpectrumCheckerboardTexture),
@@ -104,8 +123,16 @@ pub enum SpectrumTexture {
     Scaled(SpectrumScaledTexture),
 }
 
-impl TextureMapping2D {
-    pub fn create(
+pub trait CreateTextureMapping {
+    fn create(
+        params: &TextureParameterDictionary,
+        render_from_texture: &Transform,
+        loc: &FileLoc,
+    ) -> Self;
+}
+
+impl CreateTextureMapping for TextureMapping2D {
+    fn create(
         params: &TextureParameterDictionary,
         render_from_texture: &Transform,
         loc: &FileLoc,

src/filters/gaussian.rs

@@ -1,7 +1,7 @@
 use crate::core::filter::CreateFilterSampler;
 use shared::Float;
 use shared::core::filter::FilterSampler;
-use shared::core::geometry::Vector2f;
+use shared::core::geometry::{Point2f, Vector2f};
 use shared::filters::GaussianFilter;
 use shared::utils::math::gaussian;
 

src/filters/lanczos.rs

@@ -1,4 +1,5 @@
 use shared::Float;
+use shared::core::filter::FilterSampler;
 use shared::core::geometry::{Point2f, Vector2f};
 use shared::filters::LanczosSincFilter;
 use shared::utils::math::windowed_sinc;

src/filters/mitchell.rs

@@ -1,4 +1,5 @@
 use shared::Float;
+use shared::core::filter::FilterSampler;
 use shared::core::geometry::{Point2f, Vector2f};
 use shared::filters::MitchellFilter;
 

src/globals.rs

@@ -1,7 +1,7 @@
-use crate::Float;
 use crate::core::color::RGBToSpectrumTableData;
 use bytemuck::cast_slice;
 use once_cell::sync::Lazy;
+use shared::Float;
 
 static SRGB_SCALE_BYTES: &[u8] = include_bytes!("../../data/srgb_scale.dat");
 static SRGB_COEFFS_BYTES: &[u8] = include_bytes!("../../data/srgb_coeffs.dat");

src/gpu/memory.rs

@@ -1,534 +1,57 @@
-#![allow(clippy::too_many_arguments)]
-use super::Float4;
-use crate::Float;
-use crate::core::geometry::{Normal3f, Point2f, Point2i, Point3f, Point3fi, Ray, Vector3f};
-use crate::lights::LightSampleContext;
-use crate::spectra::{SampledSpectrum, SampledWavelengths};
-// use cust::memory::{CopyDestination, DeviceMemory};
-// use cust::prelude::*;
+use cudarc::driver::{CudaDevice, CudaSlice, DeviceSlice};
 
-#[macro_export]
-macro_rules! soa_struct {
-    (
-    $(#[$outer:meta])*
-    pub struct $name:ident {
-    $(
-        pub $field:ident : $type:ty
-    ),* $(,)?
-    }
-) => {
-    #[cfg(feature = "use_gpu")]
-    $(#[$outer])*
-    pub struct $name {
-        capacity: u32,
-        pub count: cust::memory::DeviceBuffer<u32>,
-        $(
-            pub $field: cust::memory::DeviceBuffer<$type>,
-        )*
+use super::{GpuError, gpu_unwrap};
+
+/// Device-only buffer (faster, but not CPU-accessible)
+pub struct DeviceBuffer<T: cudarc::driver::DeviceRepr> {
+    inner: CudaSlice<T>,
+}
+
+impl<T: cudarc::driver::DeviceRepr + Clone> DeviceBuffer<T> {
+    /// Allocate uninitialized
+    pub fn new(len: usize) -> Result<Self, GpuError> {
+        let ctx = gpu_unwrap();
+        let inner = unsafe { ctx.device.alloc::<T>(len)? };
+        Ok(Self { inner })
     }
 
-    #[cfg(feature = "use_gpu")]
-    impl $name {
-        pub fn new(capacity: usize) -> cust::error::CudaResult<Self> {
-            use cust::memory::DeviceBuffer;
-            Ok(Self {
-                capacity: capacity as u32,
-                count: DeviceBuffer::zeroed(1)?,
-                $(
-                    $field: DeviceBuffer::zeroed(capacity)?,
-                )*
-            })
-        }
-
-        pub fn len(&self) -> cust::error::CudaResult<u32> {
-            let mut host_count = [0u32; 1];
-            self.count.copy_to(&mut host_count)?;
-            Ok(host_count[0])
-        }
-
-        pub fn reset(&mut self) -> cust::error::CudaResult<()> {
-            self.count.copy_from(&[0])
-        }
-
-        // Generate the View name
-        pub fn as_view(&mut self) -> paste::paste! { [<$name View>] } {
-            paste::paste! {
-                [<$name View>] {
-                    capacity: self.capacity,
-                    count: self.count.as_device_ptr().as_mut_ptr(),
-                    $(
-                        $field: self.$field.as_device_ptr().as_raw() as *mut $type,
-                    )*
-                }
-            }
-        }
+    /// Allocate and copy from host slice
+    pub fn from_slice(data: &[T]) -> Result<Self, GpuError> {
+        let ctx = gpu_unwrap();
+        let inner = ctx.device.htod_sync_copy(data)?;
+        Ok(Self { inner })
     }
 
-    paste::paste! {
-            #[repr(C)]
-            #[derive(Clone, Copy)]
-            pub struct [<$name View>] {
-                pub capacity: u32,
-                pub count: *mut u32,
-                $(
-                    pub $field: *mut $type,
-                )*
-            }
+    /// Copy back to host
+    pub fn to_vec(&self) -> Result<Vec<T>, GpuError> {
+        let ctx = gpu_unwrap();
+        Ok(ctx.device.dtoh_sync_copy(&self.inner)?)
+    }
 
-            unsafe impl cust::memory::DeviceCopy for [<$name View>] {}
+    /// Raw device pointer (for kernel params)
+    pub fn as_ptr(&self) -> *mut T {
+        *self.inner.device_ptr() as *mut T
+    }
 
-            impl [<$name View>] {
-                // The raw push that fills every field
-                #[cfg(feature = "use_gpu")]
-                pub unsafe fn push(&self, $( $field : $type ),* ) -> Option<u32> {
-                    use core::sync::atomic::{AtomicU32, Ordering};
+    pub fn len(&self) -> usize {
+        self.inner.len()
+    }
 
-                    let index = unsafe {
-                        let counter_ptr = self.count as *mut AtomicU32;
-                        (*counter_ptr).fetch_add(1, Ordering::Relaxed)
-                    };
+    pub fn is_empty(&self) -> bool {
+        self.inner.len() == 0
+    }
 
-                    if index >= self.capacity {
-                        return None;
-                    }
+    /// Get the underlying CudaSlice (for cudarc APIs)
+    pub fn as_cuda_slice(&self) -> &CudaSlice<T> {
+        &self.inner
+    }
 
-                    unsafe {
-                        $(
-                            *self.$field.add(index as usize) = $field;
-                        )*
-                    }
-
-                    Some(index)
-                }
-
-                #[cfg(feature = "use_gpu")]
-                pub unsafe fn size(&self) -> u32 {
-                    use core::sync::atomic::{AtomicU32, Ordering};
-                    unsafe {
-                        (*(self.count as *const AtomicU32)).load(Ordering::Relaxed)
-                    }
-                }
-
-                $(
-                    #[cfg(feature = "use_gpu")]
-                    pub fn [<$field _ptr>](&self) -> *mut $type {
-                        self.$field
-                    }
-                )*
-            }
-        }
-    };
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Default)]
-pub struct RaySamplesDirect {
-    pub u: Point2f,
-    pub uc: Float,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Default)]
-pub struct RaySamplesIndirect {
-    pub uc: Float,
-    pub rr: Float,
-    pub u: Point2f,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Default)]
-pub struct RaySamplesSubsurface {
-    pub uc: Float,
-    pub u: Point2f,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Default)]
-pub struct RaySamples {
-    pub direct: RaySamplesDirect,
-    pub indirect: RaySamplesIndirect,
-    pub have_subsurface: bool,
-    pub subsurface: RaySamplesSubsurface,
-}
-
-soa_struct! {
-    pub struct RayQueue {
-        pub ray_o: Point3f,
-        pub ray_d: Vector3f,
-
-        pub depth: i32,
-        pub lambda: SampledWavelengths,
-        pub pixel_index: u32,
-
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub r_l: SampledSpectrum,
-
-        pub ctx_pi: Point3f,
-        pub ctx_n: Normal3f,
-        pub ctx_ns: Normal3f,
-
-        pub eta_scale: Float,
-        pub specular_bounce: u32,
-        pub any_non_specular_bounces: u32,
+    pub fn as_cuda_slice_mut(&mut self) -> &mut CudaSlice<T> {
+        &mut self.inner
     }
 }
 
-soa_struct! {
-    pub struct PixelSampleStateStorage {
-        pub p_pixel: Point2i,
-        pub l: SampledSpectrum,
-        pub lambda: SampledWavelengths,
-        pub filter_weight: Float,
-        pub visible_surface: u32,
-        pub camera_ray_weight: SampledSpectrum,
-
-        pub rs_direct_packed: Float4,
-        pub rs_indirect_packed: Float4,
-        pub rs_subsurface_packed: Float4,
-    }
-}
-
-soa_struct! {
-    pub struct EscapedRayQueue {
-        pub ray_o: Point3f,
-        pub ray_d: Vector3f,
-        pub depth: i32,
-        pub lambda: SampledWavelengths,
-        pub pixel_index: u32,
-        pub beta: SampledSpectrum,
-        pub specular_bounce: u32,
-        pub r_u: SampledSpectrum,
-        pub r_l: SampledSpectrum,
-        pub ctx_pi: Point3f,
-        pub ctx_n: Normal3f,
-        pub ctx_ns: Normal3f,
-    }
-}
-
-soa_struct! {
-    pub struct HitAreaLightQueue {
-        pub area_light_id: u32, // Light ID
-        pub p: Point3f,
-        pub n: Normal3f,
-        pub uv: Point2f,
-        pub wo: Vector3f,
-        pub lambda: SampledWavelengths,
-        pub depth: i32,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub r_l: SampledSpectrum,
-        pub ctx_pi: Point3f,
-        pub ctx_n: Normal3f,
-        pub ctx_ns: Normal3f,
-        pub specular_bounce: u32,
-        pub pixel_index: u32,
-    }
-}
-
-soa_struct! {
-    pub struct ShadowRayQueue {
-        pub ray_o: Point3f,
-        pub ray_d: Vector3f,
-        pub t_max: Float,
-        pub lambda: SampledWavelengths,
-        pub ld: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub r_l: SampledSpectrum,
-        pub pixel_index: u32,
-    }
-}
-
-soa_struct! {
-    pub struct GetBSSRDFAndProbeRayQueue {
-        pub material_id: u32,
-        pub lambda: SampledWavelengths,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub p: Point3f,
-        pub wo: Vector3f,
-        pub n: Normal3f,
-        pub ns: Normal3f,
-        pub dpdus: Vector3f,
-        pub uv: Point2f,
-        pub depth: i32,
-        pub mi_inside: u32,
-        pub mi_outside: u32,
-        pub eta_scale: Float,
-        pub pixel_index: u32,
-    }
-}
-
-soa_struct! {
-    pub struct SubsurfaceScatterQueue {
-        pub p0: Point3f,
-        pub p1: Point3f,
-        pub depth: i32,
-        pub material_id: u32,
-        pub lambda: SampledWavelengths,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub mi_inside: u32,
-        pub mi_outside: u32,
-        pub eta_scale: Float,
-        pub pixel_index: u32,
-    }
-}
-
-soa_struct! {
-    pub struct MediumSampleQueue {
-        pub ray_o: Point3f,
-        pub ray_d: Vector3f,
-        pub t_max: Float,
-        pub lambda: SampledWavelengths,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub r_l: SampledSpectrum,
-        pub pixel_index: u32,
-
-        pub ctx_pi: Point3f,
-        pub ctx_n: Normal3f,
-        pub ctx_ns: Normal3f,
-
-        pub specular_bounce: u32,
-        pub any_non_specular_bounces: u32,
-        pub eta_scale: Float,
-
-        pub area_light_id: u32,
-        pub pi: Point3fi,
-        pub n: Normal3f,
-        pub dpdu: Vector3f,
-        pub dpdv: Vector3f,
-        pub wo: Vector3f,
-        pub uv: Point2f,
-        pub material_id: u32,
-        pub ns: Normal3f,
-        pub dpdus: Vector3f,
-        pub dpdvs: Vector3f,
-        pub dndus: Normal3f,
-        pub dndvs: Normal3f,
-        pub face_index: i32,
-        pub mi_inside: u32,
-        pub mi_outside: u32,
-    }
-}
-
-soa_struct! {
-    pub struct MaterialEvalQueue {
-        pub material_id: u32,
-        pub pi: Point3fi,
-        pub n: Normal3f,
-        pub dpdu: Vector3f,
-        pub dpdv: Vector3f,
-        pub time: Float,
-        pub depth: i32,
-        pub ns: Normal3f,
-        pub dpdus: Vector3f,
-        pub dpdvs: Vector3f,
-        pub dndus: Normal3f,
-        pub dndvs: Normal3f,
-        pub uv: Point2f,
-        pub face_index: i32,
-        pub lambda: SampledWavelengths,
-        pub pixel_index: u32,
-        pub any_non_specular_bounces: u32,
-        pub wo: Vector3f,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub eta_scale: Float,
-        pub mi_inside: u32,
-        pub mi_outside: u32,
-    }
-}
-
-soa_struct! {
-    pub struct MediumScatterQueue {
-        pub p: Point3f,
-        pub depth: usize,
-        pub lambda: SampledWavelengths,
-        pub beta: SampledSpectrum,
-        pub r_u: SampledSpectrum,
-        pub wo: Vector3f,
-        pub time: Float,
-        pub eta_scale: Float,
-        pub pixel_index: usize,
-
-        // ID
-        pub phase_function: u32,
-        pub medium: u32,
-    }
-}
-
-#[repr(C)]
-#[derive(Clone, Copy)]
-pub struct RayWorkItem {
-    pub ray: Ray,
-    pub depth: i32,
-    pub lambda: SampledWavelengths,
-    pub pixel_index: u32,
-    pub beta: SampledSpectrum,
-    pub r_u: SampledSpectrum,
-    pub r_l: SampledSpectrum,
-    pub prev_intr_ctx: LightSampleContext,
-    pub eta_scale: Float,
-    pub specular_bounce: bool,
-    pub any_non_specular_bounces: bool,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy)]
-pub struct EscapedRayWorkItem {
-    pub ray_o: Point3f,
-    pub ray_d: Vector3f,
-    pub depth: i32,
-    pub lambda: SampledWavelengths,
-    pub pixel_index: u32,
-    pub beta: SampledSpectrum,
-    pub specular_bounce: bool,
-    pub r_u: SampledSpectrum,
-    pub r_l: SampledSpectrum,
-    pub prev_intr_ctx: LightSampleContext,
-}
-
-#[repr(C)]
-#[derive(Clone, Copy)]
-pub struct ShadowRayWorkItem {
-    pub ray: Ray,
-    pub t_max: Float,
-    pub lambda: SampledWavelengths,
-    pub ld: SampledSpectrum,
-    pub r_u: SampledSpectrum,
-    pub r_l: SampledSpectrum,
-    pub pixel_index: u32,
-}
-
-impl RayQueueView {
-    #[cfg(feature = "use_gpu")]
-    pub unsafe fn push_work_item(&self, item: RayWorkItem) -> Option<u32> {
-        unsafe {
-            self.push(
-                item.ray.o,
-                item.ray.d,
-                item.depth,
-                item.lambda,
-                item.pixel_index,
-                item.beta,
-                item.r_u,
-                item.r_l,
-                item.prev_intr_ctx.pi.into(),
-                item.prev_intr_ctx.n,
-                item.prev_intr_ctx.ns,
-                item.eta_scale,
-                if item.specular_bounce { 1 } else { 0 },
-                if item.any_non_specular_bounces { 1 } else { 0 },
-            )
-        }
-    }
-}
-
-impl EscapedRayQueueView {
-    #[cfg(feature = "use_gpu")]
-    pub unsafe fn push_work_item(&self, r: &RayWorkItem) -> Option<u32> {
-        unsafe {
-            self.push(
-                r.ray.o,
-                r.ray.d,
-                r.depth,
-                r.lambda,
-                r.pixel_index,
-                r.beta,
-                if r.specular_bounce { 1 } else { 0 },
-                r.r_u,
-                r.r_l,
-                r.prev_intr_ctx.pi.into(),
-                r.prev_intr_ctx.n,
-                r.prev_intr_ctx.ns,
-            )
-        }
-    }
-}
-
-impl PixelSampleStateStorageView {
-    #[cfg(feature = "use_gpu")]
-    pub unsafe fn get_samples(&self, index: u32) -> RaySamples {
-        let i = index as usize;
-
-        let (dir, ind, ss) = unsafe {
-            (
-                *self.rs_direct_packed.add(i),
-                *self.rs_indirect_packed.add(i),
-                *self.rs_subsurface_packed.add(i),
-            )
-        };
-
-        let direct_u = Point2f::new(dir.v[0], dir.v[1]);
-        let direct_uc = dir.v[2];
-        let flags = dir.v[3] as i32;
-        let have_subsurface = (flags & 1) != 0;
-
-        let indirect_uc = ind.v[0];
-        let indirect_rr = ind.v[1];
-        let indirect_u = Point2f::new(ind.v[2], ind.v[3]);
-
-        let subsurface_uc = ss.v[0];
-        let subsurface_u = Point2f::new(ss.v[1], ss.v[2]);
-
-        RaySamples {
-            direct: RaySamplesDirect {
-                u: direct_u,
-                uc: direct_uc,
-            },
-            indirect: RaySamplesIndirect {
-                uc: indirect_uc,
-                rr: indirect_rr,
-                u: indirect_u,
-            },
-            have_subsurface,
-            subsurface: RaySamplesSubsurface {
-                uc: subsurface_uc,
-                u: subsurface_u,
-            },
-        }
-    }
-
-    #[cfg(feature = "use_gpu")]
-    pub unsafe fn set_samples(&self, index: u32, rs: RaySamples) {
-        if index >= self.capacity {
-            return;
-        }
-        let i = index as usize;
-
-        let flags = if rs.have_subsurface { 1.0 } else { 0.0 };
-        let dir = Float4 {
-            v: [rs.direct.u.0[0], rs.direct.u.0[1], rs.direct.uc, flags],
-        };
-
-        let ind = Float4 {
-            v: [
-                rs.indirect.uc,
-                rs.indirect.rr,
-                rs.indirect.u.0[0],
-                rs.indirect.u.0[1],
-            ],
-        };
-
-        unsafe {
-            *self.rs_direct_packed.add(i) = dir;
-            *self.rs_indirect_packed.add(i) = ind;
-        }
-
-        if rs.have_subsurface {
-            let ss = Float4 {
-                v: [
-                    rs.subsurface.uc,
-                    rs.subsurface.u.0[0],
-                    rs.subsurface.u.0[1],
-                    0.0,
-                ],
-            };
-            unsafe {
-                *self.rs_subsurface_packed.add(i) = ss;
-            }
-        }
-    }
-}
+/// Unified memory buffer (CPU + GPU accessible)
+/// Note: cudarc doesn't have built-in unified memory,
+/// so we use raw CUDA calls or just use DeviceBuffer + explicit copies
+pub type UnifiedBuffer<T> = DeviceBuffer<T>;

src/gpu/mod.rs

@@ -1,5 +1,11 @@
-pub mod driver;
-pub mod memory;
+#[cfg(feature = "use_gpu")]
+mod context;
+#[cfg(feature = "use_gpu")]
+mod memory;
+#[cfg(feature = "use_gpu")]
 pub mod wavefront;
 
-pub use driver::launch_scale_kernel;
+#[cfg(feature = "use_gpu")]
+pub use context::{GpuContext, GpuError, gpu, gpu_init, gpu_unwrap};
+#[cfg(feature = "use_gpu")]
+pub use memory::UnifiedBuffer;

src/gpu/wavefront/mod.rs

@@ -1,41 +1,41 @@
 // use crate::core::scene::BasicScene;
-use crate::{
-    EscapedRayQueue, GetBSSRDFAndProbeRayQueue, HitAreaLightQueue, MaterialEvalQueue,
-    MediumSampleQueue, MediumScatterQueue, PixelSampleStateStorage, RayQueue, ShadowRayQueue,
-    SubsurfaceScatterQueue,
-};
-use shared::core::camera::Camera;
-use shared::core::film::Film;
-use shared::core::filter::Filter;
-use shared::core::light::Light;
-use shared::core::sampler::Sampler;
-use shared::lights::LightSampler;
-use std::sync::Arc;
-
-pub struct WavefrontPathIntegrator {
-    pub film: Film,
-    pub filter: Filter,
-    pub sampler: Sampler,
-    pub camera: Arc<Camera>,
-    pub light_sampler: LightSampler,
-    pub infinite_lights: Option<Vec<Arc<Light>>>,
-    pub max_depth: i32,
-    pub samples_per_pixel: i32,
-    pub regularize: bool,
-    pub scanlines_per_pixel: i32,
-    pub max_queue_size: i32,
-    pub pixel_sample_state: PixelSampleStateStorage,
-    pub ray_queue: [RayQueue; 2],
-    pub hit_area_light_queue: HitAreaLightQueue,
-    pub shadow_ray_queue: ShadowRayQueue,
-    pub escaped_ray_queue: Option<EscapedRayQueue>,
-    pub basic_material_queue: Option<MaterialEvalQueue>,
-    pub universal_material_queue: Option<MaterialEvalQueue>,
-    pub medium_sample_queue: Option<MediumSampleQueue>,
-    pub medium_scatter_queue: Option<MediumScatterQueue>,
-    pub bssrf_queue: Option<GetBSSRDFAndProbeRayQueue>,
-    pub subsurface_queue: Option<SubsurfaceScatterQueue>,
-}
+// use crate::{
+//     EscapedRayQueue, GetBSSRDFAndProbeRayQueue, HitAreaLightQueue, MaterialEvalQueue,
+//     MediumSampleQueue, MediumScatterQueue, PixelSampleStateStorage, RayQueue, ShadowRayQueue,
+//     SubsurfaceScatterQueue,
+// };
+// use shared::core::camera::Camera;
+// use shared::core::film::Film;
+// use shared::core::filter::Filter;
+// use shared::core::light::Light;
+// use shared::core::sampler::Sampler;
+// use shared::lights::sampler::LightSampler;
+// use std::sync::Arc;
+//
+// pub struct WavefrontPathIntegrator {
+//     pub film: Film,
+//     pub filter: Filter,
+//     pub sampler: Sampler,
+//     pub camera: Arc<Camera>,
+//     pub light_sampler: LightSampler,
+//     pub infinite_lights: Option<Vec<Arc<Light>>>,
+//     pub max_depth: i32,
+//     pub samples_per_pixel: i32,
+//     pub regularize: bool,
+//     pub scanlines_per_pixel: i32,
+//     pub max_queue_size: i32,
+//     pub pixel_sample_state: PixelSampleStateStorage,
+//     pub ray_queue: [RayQueue; 2],
+//     pub hit_area_light_queue: HitAreaLightQueue,
+//     pub shadow_ray_queue: ShadowRayQueue,
+//     pub escaped_ray_queue: Option<EscapedRayQueue>,
+//     pub basic_material_queue: Option<MaterialEvalQueue>,
+//     pub universal_material_queue: Option<MaterialEvalQueue>,
+//     pub medium_sample_queue: Option<MediumSampleQueue>,
+//     pub medium_scatter_queue: Option<MediumScatterQueue>,
+//     pub bssrf_queue: Option<GetBSSRDFAndProbeRayQueue>,
+//     pub subsurface_queue: Option<SubsurfaceScatterQueue>,
+// }
 
 #[cfg(feature = "use_gpu")]
 impl WavefrontPathIntegrator {

src/integrators/mod.rs

@@ -1,7 +1,8 @@
 mod pipeline;
 
+use shared::core::bsdf::BSDF;
 use shared::core::bssrdf::{BSSRDFTrait, SubsurfaceInteraction};
-use shared::core::bxdf::{BSDF, BxDFFlags, BxDFTrait, FArgs, TransportMode};
+use shared::core::bxdf::{BxDFFlags, BxDFTrait, FArgs, TransportMode};
 use shared::core::camera::Camera;
 use shared::core::film::VisibleSurface;
 use shared::core::geometry::{Bounds2i, Point2f, Point2i, Point3fi, Ray, Vector3f, VectorLike};
@@ -14,9 +15,9 @@ use shared::core::options::get_options;
 use shared::core::pbrt::{Float, SHADOW_EPSILON};
 use shared::core::primitive::{Primitive, PrimitiveTrait};
 use shared::core::sampler::{Sampler, SamplerTrait};
+use shared::core::shape::ShapeIntersection;
 use shared::lights::sampler::LightSamplerTrait;
 use shared::lights::sampler::{LightSampler, UniformLightSampler};
-use shared::shapes::ShapeIntersection;
 use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use shared::utils::hash::{hash_buffer, mix_bits};
 use shared::utils::math::{float_to_bits, sample_discrete, square};
@@ -29,6 +30,7 @@ use shared::utils::sampling::{
 use std::sync::Arc;
 
 use crate::Arena;
+use crate::spectra::get_spectra_context;
 
 #[derive(Clone, Debug)]
 pub struct IntegratorBase {
@@ -282,8 +284,9 @@ impl RayIntegratorTrait for SimplePathIntegrator {
             specular_bounce = bs.is_specular;
             ray = isect.spawn_ray(bs.wi);
         }
-        assert!(beta.y(lambda) > 0.);
-        debug_assert!(beta.y(lambda).is_finite());
+        let stdspec = get_spectra_context();
+        assert!(beta.y(lambda, stdspec) > 0.);
+        debug_assert!(beta.y(lambda, stdspec).is_finite());
         (l, None)
     }
 }
@@ -815,7 +818,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
             sampler,
             p_pixel,
             sample_ind,
-            scratch,
+            arena,
         );
     }
 
@@ -898,8 +901,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
                             r_u *= t_maj * mp.sigma_s / pdf;
 
                             if !beta.is_black() && !r_u.is_black() {
-                                let ray_medium =
-                                    ray.medium.as_ref().expect("Void scattering").clone();
+                                let ray_medium = ray.medium.expect("Void scattering").clone();
                                 let intr = MediumInteraction::new(
                                     p, -ray.d, ray.time, ray_medium, mp.phase,
                                 );
@@ -958,6 +960,8 @@ impl RayIntegratorTrait for VolPathIntegrator {
 
             // Handle surviving unscattered rays
             // Add emitted light at volume path vertex or from the environment
+
+            let stdspec = get_spectra_context();
             let Some(mut hit) = si else {
                 for light in &self.base.infinite_lights {
                     let le = light.le(&ray, lambda);
@@ -1058,7 +1062,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
                     beta,
                     r_u,
                 );
-                debug_assert!(l.y(lambda).is_finite());
+                debug_assert!(l.y(lambda, stdspec).is_finite());
             }
             let wo = isect.wo();
             let n = isect.shading.n;
@@ -1073,7 +1077,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
             } else {
                 r_l = r_u / bs.pdf;
             }
-            debug_assert!(beta.y(lambda).is_finite());
+            debug_assert!(beta.y(lambda, stdspec).is_finite());
             // Update volumetric integrator path state after surface scattering
             specular_bounce = bs.is_specular();
             if bs.is_transmissive() {
@@ -1081,7 +1085,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
             }
             ray = isect.spawn_ray_with_differentials(&ray, bs.wi, bs.flags, bs.eta);
 
-            if let Some(bssrdf) = (*isect).get_bssrdf(&ray, lambda, self.camera.as_ref(), scratch)
+            if let Some(bssrdf) = (*isect).get_bssrdf(&ray, lambda, self.camera.as_ref())
                 && bs.is_transmissive()
             {
                 let uc = sampler.get1d();
@@ -1092,8 +1096,7 @@ impl RayIntegratorTrait for VolPathIntegrator {
                 let seed = mix_bits(float_to_bits(sampler.get1d()).into());
                 let mut interaction_sampler =
                     WeightedReservoirSampler::<SubsurfaceInteraction>::new(seed);
-                let base =
-                    SimpleInteraction::new(Point3fi::new_from_point(probe_seg.p0), ray.time, None);
+                let base = SimpleInteraction::new(Point3fi::new_from_point(probe_seg.p0), ray.time);
                 loop {
                     let r = base.spawn_ray_to_point(probe_seg.p1);
                     if r.d == Vector3f::zero() {

src/integrators/pipeline.rs

@@ -1,7 +1,7 @@
-use crate::core::image::Image;
-use crate::core::{options::PBRTOptions, sampler::get_camera_sample};
+use crate::core::image::{Image, ImageMetadata};
 use crate::spectra::get_spectra_context;
 use indicatif::{ProgressBar, ProgressStyle};
+use shared::core::sampler::get_camera_sample;
 use std::io::Write;
 use std::path::Path;
 
@@ -237,8 +237,13 @@ pub fn evaluate_pixel_sample<T: RayIntegratorTrait>(
         }
 
         let initialize_visible_surface = film.uses_visible_surface();
-        let (mut l, visible_surface) =
-            integrator.li(camera_ray.ray, &lambda, sampler, initialize_visible_surface, arena);
+        let (mut l, visible_surface) = integrator.li(
+            camera_ray.ray,
+            &lambda,
+            sampler,
+            initialize_visible_surface,
+            arena,
+        );
         l *= camera_ray.weight;
 
         let std_spectra = get_spectra_context();

src/lib.rs

@@ -1,7 +1,6 @@
 pub mod core;
 pub mod filters;
 pub mod globals;
-pub mod gpu;
 pub mod integrators;
 pub mod lights;
 pub mod materials;
@@ -11,4 +10,7 @@ pub mod spectra;
 pub mod textures;
 pub mod utils;
 
+#[cfg(feature = "cuda")]
+pub mod gpu;
+
 pub use utils::arena::Arena;

src/lights/diffuse.rs

@@ -1,19 +1,20 @@
 use crate::core::image::{Image, ImageIO};
 use crate::core::light::{CreateLight, lookup_spectrum};
 use crate::core::spectrum::spectrum_to_photometric;
-use crate::core::texture::FloatTexture;
+use crate::core::texture::{FloatTexture, FloatTextureTrait};
 use crate::utils::{Arena, FileLoc, ParameterDictionary, Upload, resolve_filename};
 use log::error;
-use shared::Float;
-use shared::core::ligh::{Light, LightBase, LightType};
-use shared::core::medium::MediumInterface;
+use shared::core::geometry::Point2i;
+use shared::core::light::{Light, LightBase, LightType};
+use shared::core::medium::{Medium, MediumInterface};
 use shared::core::shape::Shape;
 use shared::core::spectrum::{Spectrum, SpectrumTrait};
-use shared::core::texture::SpectrumType;
-use shared::core::texture::{FloatTextureTrait, TextureEvalContext};
+use shared::core::texture::{SpectrumType, TextureEvalContext};
 use shared::lights::DiffuseAreaLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::{Ptr, Transform};
+use shared::{Float, PI};
+use std::fmt::Error;
 
 pub trait CreateDiffuseLight {
     fn new(
@@ -88,7 +89,7 @@ impl CreateDiffuseLight for DiffuseAreaLight {
             area: shape.area(),
             image,
             colorspace,
-            shape: Ptr::from(&*storage.shape),
+            shape,
             alpha: stored_alpha,
             lemit,
             two_sided,
@@ -105,7 +106,7 @@ impl CreateLight for DiffuseAreaLight {
         params: &ParameterDictionary,
         loc: &FileLoc,
         shape: &Shape,
-        alpha_text: &FloatTexture,
+        alpha: &FloatTexture,
         colorspace: Option<&RGBColorSpace>,
     ) -> Result<Light, Error> {
         let mut l = params.get_one_spectrum("l", None, SpectrumType::Illuminant);
@@ -146,7 +147,7 @@ impl CreateLight for DiffuseAreaLight {
         let l_for_scale = l.as_ref().unwrap_or(&colorspace.illuminant);
         scale /= spectrum_to_photometric(l_for_scale);
 
-        let phi_v = parameters.get_one_float("power", -1.0);
+        let phi_v = params.get_one_float("power", -1.0);
         if phi_v > 0.0 {
             // k_e is the emissive power of the light as defined by the spectral
             // distribution and texture and is used to normalize the emitted
@@ -155,7 +156,7 @@ impl CreateLight for DiffuseAreaLight {
 
             let mut k_e: Float = 1.0;
 
-            if let Some(ref img) = image_host {
+            if let Some(ref img) = image {
                 // Get the appropriate luminance vector from the image colour space
                 let lum_vec = image_color_space.luminance_vector();
 
@@ -175,7 +176,7 @@ impl CreateLight for DiffuseAreaLight {
             }
 
             let side_factor = if two_sided { 2.0 } else { 1.0 };
-            k_e *= side_factor * shape_data.area() * PI;
+            k_e *= side_factor * shape.area() * PI;
 
             // now multiply up scale to hit the target power
             scale *= phi_v / k_e;

src/lights/distant.rs

@@ -2,13 +2,17 @@ use crate::core::light::{CreateLight, lookup_spectrum};
 use crate::core::spectrum::spectrum_to_photometric;
 use crate::core::texture::FloatTexture;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
+use shared::Float;
 use shared::core::geometry::{Point3f, Vector3f, VectorLike};
-use shared::core::light::{Light, LightBase};
+use shared::core::light::{Light, LightBase, LightType};
 use shared::core::medium::{Medium, MediumInterface};
 use shared::core::shape::Shape;
+use shared::core::spectrum::Spectrum;
+use shared::core::texture::SpectrumType;
 use shared::lights::DistantLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::Transform;
+use std::fmt::Error;
 
 pub trait CreateDistantLight {
     fn new(render_from_light: Transform, le: Spectrum, scale: Float) -> Self;
@@ -82,7 +86,7 @@ impl CreateLight for DistantLight {
         // patch.
         let e_v = parameters.get_one_float("illuminance", -1.);
         if e_v > 0. {
-            sc *= e_v;
+            scale *= e_v;
         }
 
         let specific = DistantLight::new(final_render, l, scale);

src/lights/goniometric.rs

@@ -6,14 +6,19 @@ use crate::lights::distant::CreateDistantLight;
 use crate::utils::sampling::PiecewiseConstant2D;
 use crate::utils::{Arena, FileLoc, ParameterDictionary, resolve_filename};
 use log::error;
-use shared::Float;
+use shared::core::color::ColorEncoding;
+use shared::core::geometry::Point2i;
+use shared::core::image::PixelFormat;
 use shared::core::light::{Light, LightBase, LightType};
 use shared::core::medium::{Medium, MediumInterface};
+use shared::core::shape::Shape;
 use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
 use shared::lights::GoniometricLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::{Ptr, Transform};
+use shared::{Float, PI};
+use std::fmt::Error;
 
 pub trait CreateGoniometricLight {
     fn new(
@@ -68,7 +73,7 @@ impl CreateLight for GoniometricLight {
             colorspace.unwrap().illuminant,
             SpectrumType::Illuminant,
         );
-        let sc = params.get_one_float("scale", 1.);
+        let mut scale = params.get_one_float("scale", 1.);
         let filename = resolve_filename(params.get_one_string("filename", ""));
         let mut image: Option<Image> = None;
         let image = if filename.is_empty() {
@@ -99,13 +104,13 @@ impl CreateLight for GoniometricLight {
             Some(convert_to_luminance_image(&loaded, &filename, loc)?)
         };
 
-        scale /= spectrum_to_photometric(&lemit_data);
+        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) = image {
                 let k_e = compute_emissive_power(image);
-                scale *= phi_v / phi_e;
+                scale *= phi_v / k_e;
             }
         }
 
@@ -116,7 +121,7 @@ impl CreateLight for GoniometricLight {
         let final_render_from_light = render_from_light * t;
 
         let specific =
-            GoniometricLight::new(final_render_from_light, medium.into(), le, scale, image);
+            GoniometricLight::new(final_render_from_light, medium.into(), i, scale, image);
 
         Ok(Light::Goniometric(specific))
     }

src/lights/infinite.rs

@@ -1,16 +1,26 @@
+use crate::Arena;
+use crate::core::image::Image;
+use crate::core::spectrum::spectrum_to_photometric;
+use crate::utils::{FileLoc, ParameterDictionary, resolve_filename};
 use log::error;
-use shared::Float;
-use shared::core::geometry::Point2f;
-use shared::core::light::{CreateLight, Light, LightBase, LightType};
+use shared::core::camera::CameraTransform;
+use shared::core::geometry::{Bounds2f, Frame, Point2f, Point2i, Point3f, cos_theta};
+use shared::core::image::{PixelFormat, WrapMode};
+use shared::core::light::{Light, LightBase, LightType};
 use shared::core::medium::MediumInterface;
 use shared::core::spectrum::Spectrum;
+use shared::core::texture::SpectrumType;
 use shared::lights::{ImageInfiniteLight, PortalInfiniteLight, UniformInfiniteLight};
 use shared::spectra::RGBColorSpace;
-use shared::utils::sampling::DevicePiecewiseConstant2D;
+use shared::utils::hash::hash_float;
+use shared::utils::math::{equal_area_sphere_to_square, equal_area_square_to_sphere};
+use shared::utils::sampling::{DevicePiecewiseConstant2D, WindowedPiecewiseConstant2D};
 use shared::utils::{Ptr, Transform};
+use shared::{Float, PI};
+use std::fmt::Error;
 use std::sync::Arc;
 
-use crate::core::light::{LightBaseTrait, lookup_spectrum};
+use crate::core::light::lookup_spectrum;
 
 pub trait CreateImageInfiniteLight {
     fn new(
@@ -43,10 +53,9 @@ impl CreateImageInfiniteLight for ImageInfiniteLight {
         assert_eq!(3, desc.size());
         assert!(desc.is_identity());
         if image.resolution().x() != image.resolution().y() {
-            hash(hashee, into);
+            let into = hash_float(hashee);
             panic!(
-                "{}: image resolution ({}, {}) is non-square. It's unlikely this is an equal area environment map.",
-                filename,
+                "Image resolution ({}, {}) is non-square. It's unlikely this is an equal area environment map.",
                 image.resolution.x(),
                 image.resolution.y()
             );
@@ -75,7 +84,7 @@ impl CreateImageInfiniteLight for ImageInfiniteLight {
         ImageInfiniteLight {
             base,
             image: &image,
-            image_color_space: &storage.image_color_space,
+            image_color_space,
             scene_center: Point3f::default(),
             scene_radius: 0.,
             scale,
@@ -126,18 +135,14 @@ impl CreatePortalInfiniteLight for PortalInfiniteLight {
         let desc = image
             .get_channel_desc(&["R", "G", "B"])
             .unwrap_or_else(|_| {
-                panic!(
-                    "{}: image used for PortalImageInfiniteLight doesn't have R, G, B channels.",
-                    filename
-                )
+                panic!("Image used for PortalImageInfiniteLight doesn't have R, G, B channels.",)
             });
 
         assert_eq!(3, desc.offset.len());
-        let src_res = equal_area_image.resolution;
+        let src_res = image.resolution;
         if src_res.x() != src_res.y() {
             panic!(
-                "{}: image resolution ({}, {}) is non-square. It's unlikely this is an equal area environment map.",
-                filename,
+                "Image resolution ({}, {}) is non-square. It's unlikely this is an equal area environment map.",
                 src_res.x(),
                 src_res.y()
             );
@@ -195,7 +200,7 @@ impl CreatePortalInfiniteLight for PortalInfiniteLight {
                     let pixel_idx = (x * 3) as usize;
 
                     for c in 0..3 {
-                        let val = equal_area_image.bilerp_channel_with_wrap(
+                        let val = image.bilerp_channel_with_wrap(
                             uv_equi,
                             c,
                             WrapMode::OctahedralSphere.into(),
@@ -205,19 +210,14 @@ impl CreatePortalInfiniteLight for PortalInfiniteLight {
                 }
             });
 
-        let image = Image::new(
-            PixelFormat::F32,
-            src_res,
-            &["R", "G", "B"],
-            equal_area_image.encoding,
-        );
+        let img = Image::new(PixelFormat::F32, src_res, &["R", "G", "B"], image.encoding);
 
         let duv_dw_closure = |p: Point2f| -> Float {
             let (_, jacobian) = Self::render_from_image(portal_frame, p);
             jacobian
         };
 
-        let d = image.get_sampling_distribution(
+        let d = img.get_sampling_distribution(
             duv_dw_closure,
             Bounds2f::from_points(Point2f::new(0., 0.), Point2f::new(1., 1.)),
         );
@@ -226,7 +226,7 @@ impl CreatePortalInfiniteLight for PortalInfiniteLight {
 
         PortalInfiniteLight {
             base,
-            image,
+            image: img,
             image_color_space: &image_color_space,
             scale,
             scene_center: Point3f::default(),
@@ -340,7 +340,7 @@ fn load_image_or_constant(
     loc: &FileLoc,
 ) -> Result<(Image, RGBColorSpace), Error> {
     if filename.is_empty() {
-        let rgb = spectrum_to_rgb(&l[0], colorspace);
+        let rgb = &l[0].to_rgb();
         let image = Image::new_constant(Point2i::new(1, 1), &["R", "G", "B"], &rgb);
         Ok((image, colorspace.clone()))
     } else {

src/lights/point.rs

@@ -1,15 +1,18 @@
-use crate::core::light::{CreateLight, LightBaseTrait, lookup_spectrum};
+use crate::core::light::{CreateLight, lookup_spectrum};
 use crate::core::spectrum::spectrum_to_photometric;
 use crate::core::texture::FloatTexture;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
+use shared::core::geometry::Point3f;
 use shared::core::light::{Light, LightBase, LightType};
 use shared::core::medium::{Medium, MediumInterface};
 use shared::core::shape::Shape;
 use shared::core::spectrum::Spectrum;
+use shared::core::texture::SpectrumType;
 use shared::lights::PointLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::Transform;
 use shared::{Float, PI};
+use std::fmt::Error;
 
 pub trait CreatePointLight {
     fn new(
@@ -61,7 +64,7 @@ impl CreateLight for PointLight {
         let phi_v = parameters.get_one_float("power", 1.);
         if phi_v > 0. {
             let k_e = 4. * PI;
-            sc *= phi_v / k_e;
+            scale *= phi_v / k_e;
         }
 
         let from = parameters.get_one_point3f("from", Point3f::zero());

src/lights/projection.rs

@@ -1,25 +1,28 @@
 use crate::core::image::{Image, ImageIO};
 use crate::core::light::CreateLight;
 use crate::core::spectrum::spectrum_to_photometric;
-use crate::spectra::colorspace::new;
-use crate::utils::{Arena, ParameterDictionary, Upload, resolve_filename};
+use crate::core::texture::FloatTexture;
+use crate::utils::sampling::PiecewiseConstant2D;
+use crate::utils::{Arena, FileLoc, ParameterDictionary, Upload, resolve_filename};
 use log::error;
 use shared::Float;
-use shared::core::geometry::{Bounds2f, VectorLike};
+use shared::core::geometry::{
+    Bounds2f, Point2f, Point2i, Point3f, Vector3f, VectorLike, cos_theta,
+};
 use shared::core::image::ImageAccess;
-use shared::core::light::{Light, LightBase};
-use shared::core::medium::MediumInterface;
-use shared::core::spectrum::Spectrum;
+use shared::core::light::{Light, LightBase, LightType};
+use shared::core::medium::{Medium, MediumInterface};
+use shared::core::shape::Shape;
 use shared::lights::ProjectionLight;
 use shared::spectra::RGBColorSpace;
 use shared::utils::math::{radians, square};
 use shared::utils::{Ptr, Transform};
+use std::fmt::Error;
 
 pub trait CreateProjectionLight {
     fn new(
         render_from_light: Transform,
         medium_interface: MediumInterface,
-        le: Spectrum,
         scale: Float,
         image: Ptr<Image>,
         image_color_space: Ptr<RGBColorSpace>,
@@ -31,7 +34,6 @@ impl CreateProjectionLight for ProjectionLight {
     fn new(
         render_from_light: Transform,
         medium_interface: MediumInterface,
-        le: Spectrum,
         scale: Float,
         image: Ptr<Image>,
         image_color_space: Ptr<RGBColorSpace>,
@@ -140,8 +142,7 @@ impl CreateLight for ProjectionLight {
 
         let specific = ProjectionLight::new(
             render_from_light_flip,
-            medium_interface,
-            le,
+            medium.into(),
             scale,
             image.upload(arena),
             colorspace.upload(arena),

src/lights/sampler.rs

@@ -1,5 +1,9 @@
 use crate::utils::sampling::AliasTableHost;
+use shared::core::light::Light;
+use shared::spectra::{SampledSpectrum, SampledWavelengths};
+use shared::utils::sampling::AliasTable;
 use std::collections::HashMap;
+use std::sync::Arc;
 
 pub struct PowerSamplerHost {
     pub lights: Vec<Light>,

src/lights/spot.rs

@@ -1,17 +1,20 @@
 // use crate::core::image::{Image, ImageIO, ImageMetadata};
-use crate::core::light::CreateLight;
+use crate::core::light::{CreateLight, lookup_spectrum};
 use crate::core::spectrum::spectrum_to_photometric;
-use crate::utils::{Arena, ParameterDictionary};
-use shared::core::geometry::{Frame, VectorLike};
+use crate::core::texture::FloatTexture;
+use crate::utils::{Arena, FileLoc, ParameterDictionary};
+use shared::core::geometry::{Frame, Point3f, VectorLike};
 use shared::core::light::{Light, LightBase, LightType};
-use shared::core::medium::MediumInterface;
+use shared::core::medium::{Medium, MediumInterface};
+use shared::core::shape::Shape;
 use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
 use shared::lights::SpotLight;
 use shared::spectra::RGBColorSpace;
-use shared::utils::Transformk
+use shared::utils::Transform;
 use shared::utils::math::radians;
 use shared::{Float, PI};
+use std::fmt::Error;
 
 pub trait CreateSpotLight {
     fn new(
@@ -69,8 +72,8 @@ impl CreateLight for SpotLight {
             )
             .expect("No spectrum");
         let mut scale = parameters.get_one_float("scale", 1.);
-        let coneangle = parameters.get_one_float("coneangle", def);
-        let conedelta = parameters.get_one_float("conedelta", def);
+        let coneangle = parameters.get_one_float("coneangle", 30.);
+        let conedelta = parameters.get_one_float("conedelta", 5.);
         let from = parameters.get_one_point3f("from", Point3f::zero());
         let to = parameters.get_one_point3f("to", Point3f::new(0., 0., 1.));
         let dir_to_z = Transform::from(Frame::from_z((to - from).normalize()));
@@ -90,7 +93,7 @@ impl CreateLight for SpotLight {
         let specific = SpotLight::new(
             final_render,
             medium.into(),
-            le,
+            i,
             scale,
             coneangle,
             coneangle - conedelta,

src/materials/coated.rs

@@ -1,7 +1,8 @@
 use crate::core::image::Image;
 use crate::core::material::CreateMaterial;
 use crate::core::texture::SpectrumTexture;
-use crate::utils::{Arena, FileLoc, TextureParameterDictionary, Upload};
+use crate::spectra::data::get_named_spectrum;
+use crate::utils::{Arena, FileLoc, TextureParameterDictionary, Upload, parameters::error_exit};
 use shared::core::material::Material;
 use shared::core::spectrum::Spectrum;
 use shared::core::texture::SpectrumType;
@@ -9,6 +10,8 @@ use shared::materials::coated::*;
 use shared::spectra::ConstantSpectrum;
 use shared::textures::SpectrumConstantTexture;
 use std::collections::HashMap;
+use std::fmt::Error;
+use std::sync::Arc;
 
 impl CreateMaterial for CoatedDiffuseMaterial {
     fn create(
@@ -65,7 +68,7 @@ impl CreateMaterial for CoatedDiffuseMaterial {
             n_samples,
         );
 
-        Ok(Material::CoatedDiffuse(specific));
+        Ok(Material::CoatedDiffuse(specific))
     }
 }
 
@@ -76,7 +79,7 @@ impl CreateMaterial for CoatedConductorMaterial {
         named_materials: &HashMap<String, Material>,
         loc: &FileLoc,
         arena: &mut Arena,
-    ) -> Result<Self, String> {
+    ) -> Result<Material, Error> {
         let interface_u_roughness = parameters
             .get_float_texture_or_null("interface.uroughness")
             .r_else(|| parameters.get_float_texture("interface.roughness", 0.));
@@ -147,6 +150,6 @@ impl CreateMaterial for CoatedConductorMaterial {
             n_samples,
         );
         arena.alloc(material);
-        return material;
+        Ok(Material::CoatedConductor(material))
     }
 }

src/materials/complex.rs

@@ -1,15 +1,26 @@
+use crate::core::image::Image;
 use crate::core::material::CreateMaterial;
+use crate::spectra::get_colorspace_context;
 use crate::utils::{Arena, FileLoc, TextureParameterDictionary, Upload};
 use shared::bxdfs::HairBxDF;
-use shared::core::material::Material;
+use shared::core::bsdf::BSDF;
+use shared::core::bssrdf::BSSRDF;
+use shared::core::material::{Material, MaterialEvalContext, MaterialTrait};
 use shared::core::spectrum::Spectrum;
-use shared::core::texture::GPUFloatTexture;
+use shared::core::texture::{GPUFloatTexture, SpectrumType, TextureEvaluator};
 use shared::materials::complex::*;
+use shared::spectra::RGBUnboundedSpectrum;
+use shared::spectra::SampledWavelengths;
+use shared::textures::SpectrumConstantTexture;
+use shared::utils::Ptr;
+use std::collections::HashMap;
+use std::fmt::Error;
+use std::sync::Arc;
 
 impl CreateMaterial for HairMaterial {
     fn create(
         parameters: &TextureParameterDictionary,
-        normal_map: Option<Arc<ImageBuffer>>,
+        normal_map: Option<Arc<Image>>,
         _named_materials: &HashMap<String, Material>,
         loc: &FileLoc,
         arena: &mut Arena,
@@ -23,10 +34,14 @@ impl CreateMaterial for HairMaterial {
         let has_melanin = eumelanin.is_some() || pheomelanin.is_some();
 
         // Default distribution if nothing is spceified
-        let sigma_a = if sigma_a.is_none() && color.is_none() && !has_melanin {
-            let default_rgb = HairBxDF::sigma_a_from_concentration(1.3, 0.0);
+        let sigma_a = if sigma_a.is_none() && !reflectance.is_none() && !has_melanin {
+            let stdcs = get_colorspace_context();
+            let default_rgb = HairBxDF::sigma_a_from_concentration(1.3, 0.0, stdcs);
             Some(Arc::new(SpectrumConstantTexture::new(
-                Spectrum::RGBUnbounded(RGBUnboundedSpectrum::new(default_rgb)),
+                Spectrum::RGBUnbounded(RGBUnboundedSpectrum::new(
+                    reflectance.to_rgb(),
+                    default_rgb,
+                )),
             )))
         } else {
             sigma_a
@@ -47,7 +62,7 @@ impl CreateMaterial for HairMaterial {
             alpha.upload(arena),
         );
 
-        return material;
+        Ok(Material::Hair(material))
     }
 }
 
@@ -89,34 +104,14 @@ impl MaterialTrait for MeasuredMaterial {
 #[repr(C)]
 #[derive(Clone, Copy, Debug)]
 pub struct SubsurfaceMaterial;
-impl MaterialTrait for SubsurfaceMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> BSDF {
-        todo!()
-    }
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        todo!()
-    }
-
-    fn can_evaluate_textures(&self, _tex_eval: &dyn TextureEvaluator) -> bool {
-        todo!()
-    }
-    fn get_normal_map(&self) -> *const Image {
-        todo!()
-    }
-    fn get_displacement(&self) -> Ptr<GPUFloatTexture> {
-        todo!()
-    }
-    fn has_subsurface_scattering(&self) -> bool {
+impl CreateMaterial for SubsurfaceMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &HashMap<String, Material>,
+        loc: &FileLoc,
+        arena: &mut Arena,
+    ) -> Result<Material, Error> {
         todo!()
     }
 }

src/materials/conductor.rs

@@ -1,52 +1,19 @@
 use crate::core::image::Image;
-use crate::core::scattering::TrowbridgeReitzDistribution;
-use shared::core::bsdf::BSDF;
-use shared::core::bssrdf::BSSRDF;
-use shared::core::material::{MaterialEvalContext, MaterialTrait};
-use shared::core::texture::{GPUFloatTexture, GPUSpectrumTexture, TextureEvaluator};
-use shared::spectra::SampledWavelengths;
-use shared::utils::Ptr;
+use crate::core::material::CreateMaterial;
+// use crate::core::scattering::TrowbridgeReitzDistribution;
+use crate::utils::TextureParameterDictionary;
+use shared::core::material::Material;
+use shared::materials::ConductorMaterial;
+use std::sync::Arc;
 
-#[repr(C)]
-#[derive(Clone, Copy, Debug)]
-pub struct ConductorMaterial {
-    pub displacement: Ptr<GPUFloatTexture>,
-    pub eta: Ptr<GPUSpectrumTexture>,
-    pub k: Ptr<GPUSpectrumTexture>,
-    pub reflectance: Ptr<GPUSpectrumTexture>,
-    pub u_roughness: Ptr<GPUFloatTexture>,
-    pub v_roughness: Ptr<GPUFloatTexture>,
-    pub remap_roughness: bool,
-    pub normal_map: Ptr<Image>,
-}
-
-impl MaterialTrait for ConductorMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> BSDF {
-        todo!()
-    }
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        todo!()
-    }
-    fn can_evaluate_textures(&self, _tex_eval: &dyn TextureEvaluator) -> bool {
-        todo!()
-    }
-    fn get_normal_map(&self) -> *const Image {
-        todo!()
-    }
-    fn get_displacement(&self) -> Ptr<GPUFloatTexture> {
-        todo!()
-    }
-    fn has_subsurface_scattering(&self) -> bool {
+impl CreateMaterial for ConductorMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &std::collections::HashMap<String, shared::core::material::Material>,
+        loc: &crate::utils::FileLoc,
+        arena: &mut crate::Arena,
+    ) -> Result<Material, std::fmt::Error> {
         todo!()
     }
 }

src/materials/dielectric.rs

@@ -1,111 +1,33 @@
+use crate::Arena;
 use crate::core::image::Image;
-use crate::core::scattering::TrowbridgeReitzDistribution;
-use shared::core::bsdf::BSDF;
-use shared::core::bssrdf::BSSRDF;
-use shared::core::bxdf::BxDF;
-use shared::core::material::{MaterialEvalContext, MaterialTrait};
-use shared::core::spectrum::{Spectrum, SpectrumTrait};
-use shared::core::texture::{GPUFloatTexture, TextureEvaluator};
-use shared::spectra::SampledWavelengths;
-use shared::utils::Ptr;
+use crate::core::material::CreateMaterial;
+use crate::utils::{FileLoc, TextureParameterDictionary};
+use shared::core::material::Material;
+use shared::materials::{DielectricMaterial, ThinDielectricMaterial};
+use std::collections::HashMap;
+use std::fmt::Error;
+use std::sync::Arc;
 
-#[repr(C)]
-#[derive(Clone, Copy, Debug)]
-pub struct DielectricMaterial {
-    normal_map: *const Image,
-    displacement: Ptr<GPUFloatTexture>,
-    u_roughness: Ptr<GPUFloatTexture>,
-    v_roughness: Ptr<GPUFloatTexture>,
-    remap_roughness: bool,
-    eta: Ptr<Spectrum>,
-}
-
-impl MaterialTrait for DielectricMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        tex_eval: &T,
-        ctx: &MaterialEvalContext,
-        lambda: &SampledWavelengths,
-    ) -> BSDF {
-        let mut sampled_eta = self.eta.evaluate(lambda[0]);
-        if !self.eta.is_constant() {
-            lambda.terminate_secondary();
-        }
-
-        if sampled_eta == 0.0 {
-            sampled_eta = 1.0;
-        }
-
-        let mut u_rough = tex_eval.evaluate_float(&self.u_roughness, ctx);
-        let mut v_rough = tex_eval.evaluate_float(&self.v_roughness, ctx);
-
-        if self.remap_roughness {
-            u_rough = TrowbridgeReitzDistribution::roughness_to_alpha(u_rough);
-            v_rough = TrowbridgeReitzDistribution::roughness_to_alpha(v_rough);
-        }
-
-        let distrib = TrowbridgeReitzDistribution::new(u_rough, v_rough);
-        let bxdf = BxDF::Dielectric(DielectricBxDF::new(sampled_eta, distrib));
-
-        BSDF::new(ctx.ns, ctx.dpdus, Some(bxdf))
-    }
-
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        None
-    }
-
-    fn can_evaluate_textures(&self, tex_eval: &dyn TextureEvaluator) -> bool {
-        tex_eval.can_evaluate(&[&self.u_roughness, &self.v_roughness], &[])
-    }
-
-    fn get_normal_map(&self) -> *const Image {
-        self.normal_map
-    }
-
-    fn get_displacement(&self) -> Ptr<GPUFloatTexture> {
-        self.displacement
-    }
-
-    fn has_subsurface_scattering(&self) -> bool {
-        false
-    }
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug)]
-pub struct ThinDielectricMaterial;
-impl MaterialTrait for ThinDielectricMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> BSDF {
-        todo!()
-    }
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        todo!()
-    }
-    fn can_evaluate_textures(&self, _tex_eval: &dyn TextureEvaluator) -> bool {
-        todo!()
-    }
-    fn get_normal_map(&self) -> *const Image {
-        todo!()
-    }
-    fn get_displacement(&self) -> Ptr<GPUFloatTexture> {
-        todo!()
-    }
-    fn has_subsurface_scattering(&self) -> bool {
+impl CreateMaterial for DielectricMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &HashMap<String, Material>,
+        loc: &FileLoc,
+        arena: &mut Arena,
+    ) -> Result<Material, Error> {
+        todo!()
+    }
+}
+
+impl CreateMaterial for ThinDielectricMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &HashMap<String, Material>,
+        loc: &FileLoc,
+        arena: &mut Arena,
+    ) -> Result<Material, Error> {
         todo!()
     }
 }

src/materials/diffuse.rs

@@ -1,94 +1,33 @@
+use crate::Arena;
 use crate::core::image::Image;
-use crate::core::scattering::TrowbridgeReitzDistribution;
-use shared::core::bsdf::BSDF;
-use shared::core::bssrdf::BSSRDF;
-use shared::core::bxdf::BxDF;
-use shared::core::material::{MaterialEvalContext, MaterialTrait};
-use shared::core::texture::{GPUFloatTexture, GPUSpectrumTexture, TextureEvaluator};
-use shared::spectra::SampledWavelengths;
-use shared::utils::Ptr;
+use crate::core::material::CreateMaterial;
+use crate::utils::{FileLoc, TextureParameterDictionary};
+use shared::core::material::Material;
+use shared::materials::{DiffuseMaterial, DiffuseTransmissionMaterial};
+use std::collections::HashMap;
+use std::fmt::Error;
+use std::sync::Arc;
 
-#[repr(C)]
-#[derive(Clone, Copy, Debug)]
-pub struct DiffuseMaterial {
-    pub normal_map: *const Image,
-    pub displacement: Ptr<GPUFloatTexture>,
-    pub reflectance: Ptr<GPUSpectrumTexture>,
-}
-
-impl MaterialTrait for DiffuseMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        tex_eval: &T,
-        ctx: &MaterialEvalContext,
-        lambda: &SampledWavelengths,
-    ) -> BSDF {
-        let r = tex_eval.evaluate_spectrum(&self.reflectance, ctx, lambda);
-        let bxdf = BxDF::Diffuse(DiffuseBxDF::new(r));
-        BSDF::new(ctx.ns, ctx.dpdus, Some(bxdf))
-    }
-
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        todo!()
-    }
-
-    fn can_evaluate_textures(&self, tex_eval: &dyn TextureEvaluator) -> bool {
-        tex_eval.can_evaluate(&[], &[self.reflectance])
-    }
-
-    fn get_normal_map(&self) -> *const Image {
-        self.normal_map
-    }
-
-    fn get_displacement(&self) -> Ptr<GPUFloatTexture> {
-        self.displacement
-    }
-
-    fn has_subsurface_scattering(&self) -> bool {
-        false
-    }
-}
-
-#[repr(C)]
-#[derive(Clone, Copy, Debug)]
-pub struct DiffuseTransmissionMaterial;
-
-impl MaterialTrait for DiffuseTransmissionMaterial {
-    fn get_bsdf<T: TextureEvaluator>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> BSDF {
-        todo!()
-    }
-    fn get_bssrdf<T>(
-        &self,
-        _tex_eval: &T,
-        _ctx: &MaterialEvalContext,
-        _lambda: &SampledWavelengths,
-    ) -> Option<BSSRDF> {
-        todo!()
-    }
-
-    fn can_evaluate_textures(&self, _tex_eval: &dyn TextureEvaluator) -> bool {
-        todo!()
-    }
-
-    fn get_normal_map(&self) -> Option<Arc<Image>> {
-        todo!()
-    }
-
-    fn get_displacement(&self) -> Option<FloatTexture> {
-        todo!()
-    }
-
-    fn has_subsurface_scattering(&self) -> bool {
+impl CreateMaterial for DiffuseMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &HashMap<String, Material>,
+        loc: &FileLoc,
+        arena: &mut Arena,
+    ) -> Result<Material, Error> {
+        todo!()
+    }
+}
+
+impl CreateMaterial for DiffuseTransmissionMaterial {
+    fn create(
+        parameters: &TextureParameterDictionary,
+        normal_map: Option<Arc<Image>>,
+        named_materials: &HashMap<String, Material>,
+        loc: &FileLoc,
+        arena: &mut Arena,
+    ) -> Result<Material, Error> {
         todo!()
     }
 }

src/materials/mix.rs

@@ -1,5 +1,4 @@
 use crate::core::image::Image;
-use crate::core::scattering::TrowbridgeReitzDistribution;
 use shared::core::bsdf::BSDF;
 use shared::core::bssrdf::BSSRDF;
 use shared::core::material::{Material, MaterialEvalContext, MaterialTrait};

src/samplers/halton.rs

@@ -1,7 +1,10 @@
+use crate::Arena;
 use crate::core::sampler::CreateSampler;
 use crate::utils::{FileLoc, ParameterDictionary};
 use shared::core::geometry::Point2i;
+use shared::core::options::get_options;
 use shared::core::sampler::{HaltonSampler, RandomizeStrategy};
+use std::fmt::Error;
 
 impl CreateSampler for HaltonSampler {
     fn create(

src/samplers/independent.rs

@@ -1,7 +1,10 @@
+use crate::Arena;
 use crate::core::sampler::CreateSampler;
 use crate::utils::{FileLoc, ParameterDictionary};
 use shared::core::geometry::Point2i;
+use shared::core::options::get_options;
 use shared::core::sampler::IndependentSampler;
+use std::fmt::Error;
 
 impl CreateSampler for IndependentSampler {
     fn create(

src/samplers/sobol.rs

@@ -1,7 +1,10 @@
+use crate::Arena;
 use crate::core::sampler::CreateSampler;
 use crate::utils::{FileLoc, ParameterDictionary};
 use shared::core::geometry::Point2i;
+use shared::core::options::get_options;
 use shared::core::sampler::{PaddedSobolSampler, RandomizeStrategy, SobolSampler, ZSobolSampler};
+use std::fmt::Error;
 
 impl CreateSampler for SobolSampler {
     fn create(

src/samplers/stratified.rs

@@ -1,8 +1,10 @@
 use crate::Arena;
 use crate::core::sampler::CreateSampler;
 use crate::utils::{FileLoc, ParameterDictionary};
-use shared::core::geometry::{FileLoc, ParameterDictionary};
+use shared::core::geometry::Point2i;
+use shared::core::options::get_options;
 use shared::core::sampler::StratifiedSampler;
+use std::fmt::Error;
 
 impl CreateSampler for StratifiedSampler {
     fn create(

src/shapes/bilinear.rs

@@ -1,22 +1,27 @@
-use crate::core::shape::{ALL_BILINEAR_MESHES, CreateMesh};
+use crate::core::image::Image;
+use crate::core::shape::{ALL_BILINEAR_MESHES, CreateShape};
 use crate::core::texture::FloatTexture;
+use crate::shapes::mesh::BilinearPatchMesh;
+use crate::utils::sampling::PiecewiseConstant2D;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
 use log::warn;
+use shared::core::geometry::{Bounds2f, Point2f};
+use shared::core::shape::Shape;
 use shared::shapes::BilinearPatchShape;
-use shared::shapes::mesh::Mesh;
 use shared::utils::Transform;
 use std::collections::HashMap;
+use std::sync::Arc;
 
-impl CreateMesh for BilinearPatchShape {
+impl CreateShape for BilinearPatchShape {
     fn create(
-        render_from_object: &Transform,
-        _object_from_render: &Transform,
+        render_from_object: Transform,
+        _object_from_render: Transform,
         reverse_orientation: bool,
-        parameters: &ParameterDictionary,
-        _float_textures: &HashMap<String, FloatTexture>,
-        _loc: &FileLoc,
+        parameters: ParameterDictionary,
+        _float_textures: HashMap<String, FloatTexture>,
+        _loc: FileLoc,
         arena: &mut Arena,
-    ) -> Result<Mesh, String> {
+    ) -> Result<Vec<Shape>, String> {
         let mut vertex_indices = parameters.get_int_array("indices");
         let mut p = parameters.get_point3f_array("P");
         let mut uv = parameters.get_point2f_array("uv");
@@ -105,7 +110,7 @@ impl CreateMesh for BilinearPatchShape {
             }
         }
 
-        let host = BilinearPatchMeshHost::new(
+        let host = BilinearPatchMesh::new(
             render_from_object,
             reverse_orientation,
             vertex_indices,

src/shapes/curves.rs

@@ -2,12 +2,16 @@ use crate::Arena;
 use crate::core::shape::CreateShape;
 use crate::core::texture::FloatTexture;
 use crate::utils::{FileLoc, ParameterDictionary};
-use shared::core::geometry::Normal3f;
+use shared::Float;
+use shared::core::geometry::{Normal3f, Point3f};
+use shared::core::shape::Shape;
 use shared::shapes::{CurveCommon, CurveShape, CurveType};
 use shared::utils::Transform;
+use shared::utils::math::lerp;
 use shared::utils::splines::{
     cubic_bspline_to_bezier, elevate_quadratic_bezier_to_cubic, quadratic_bspline_to_bezier,
 };
+use std::sync::Arc;
 
 use log::warn;
 use std::collections::HashMap;
@@ -40,8 +44,8 @@ pub fn create_curve(
         let u_min = i as Float / n_segments as Float;
         let u_max = (i + 1) as Float / n_segments as Float;
 
-        let curve = Curve {
-            common: common.clone(),
+        let curve = CurveShape {
+            common: curve_common.clone(),
             u_min,
             u_max,
         };
@@ -53,7 +57,6 @@ pub fn create_curve(
 
 impl CreateShape for CurveShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         object_from_render: Transform,
         reverse_orientation: bool,
@@ -146,7 +149,7 @@ impl CreateShape for CurveShape {
             parameters.get_one_int("splitdepth", 3)
         };
 
-        let mut curves: Vec<Arc<dyn Shape>> = Vec::new();
+        let mut curves: Vec<Arc<Shape>> = Vec::new();
         let mut cp_offset = 0;
 
         for seg in 0..n_segments {

src/shapes/cylinder.rs

@@ -8,7 +8,6 @@ use std::collections::HashMap;
 
 impl CreateShape for CylinderShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         object_from_render: Transform,
         reverse_orientation: bool,

src/shapes/disk.rs

@@ -8,7 +8,6 @@ use std::collections::HashMap;
 
 impl CreateShape for DiskShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         object_from_render: Transform,
         reverse_orientation: bool,

src/shapes/mesh.rs

@@ -1,22 +1,261 @@
+// use crate::Arena;
+use crate::utils::sampling::PiecewiseConstant2D;
 use anyhow::{Context, Result as AnyResult, bail};
 use ply_rs::parser::Parser;
 use ply_rs::ply::{DefaultElement, Property};
+use shared::core::geometry::{Normal3f, Point2f, Point3f, Vector3f};
 use shared::utils::Transform;
-use shared::utils::mesh::{BilinearPatchMesh, TriangleMesh};
-use shared::utils::sampling::DevicePiecewiseConstant2D;
+use shared::utils::mesh::{DeviceBilinearPatchMesh, DeviceTriangleMesh};
 use std::fs::File;
+use std::ops::Deref;
 use std::path::Path;
 
-#[derive(Debug, Clone, Copy, Default)]
+/// Intermediate mesh from PLY
+#[derive(Debug, Clone, Default)]
 pub struct TriQuadMesh {
     pub p: Vec<Point3f>,
     pub n: Vec<Normal3f>,
     pub uv: Vec<Point2f>,
-    pub face_indices: Vec<u32>,
+    pub face_indices: Vec<i32>,
     pub tri_indices: Vec<u32>,
     pub quad_indices: Vec<u32>,
 }
 
+#[derive(Debug)]
+struct TriangleMeshStorage {
+    vertex_indices: Vec<u32>,
+    p: Vec<Point3f>,
+    n: Vec<Normal3f>,
+    s: Vec<Vector3f>,
+    uv: Vec<Point2f>,
+    face_indices: Vec<i32>,
+}
+
+#[derive(Debug)]
+struct BilinearMeshStorage {
+    vertex_indices: Vec<u32>,
+    p: Vec<Point3f>,
+    n: Vec<Normal3f>,
+    uv: Vec<Point2f>,
+    image_distribution: Option<PiecewiseConstant2D>,
+}
+
+#[derive(Debug)]
+pub struct TriangleMesh {
+    _storage: Box<TriangleMeshStorage>,
+    pub device: DeviceTriangleMesh,
+}
+
+#[derive(Debug)]
+pub struct BilinearPatchMesh {
+    _storage: Box<BilinearMeshStorage>,
+    pub device: DeviceBilinearPatchMesh,
+}
+
+impl Deref for TriangleMesh {
+    type Target = DeviceTriangleMesh;
+    #[inline]
+    fn deref(&self) -> &Self::Target {
+        &self.device
+    }
+}
+
+impl Deref for BilinearPatchMesh {
+    type Target = DeviceBilinearPatchMesh;
+    #[inline]
+    fn deref(&self) -> &Self::Target {
+        &self.device
+    }
+}
+
+impl TriangleMesh {
+    pub fn new(
+        render_from_object: &Transform,
+        reverse_orientation: bool,
+        vertex_indices: Vec<u32>,
+        mut p: Vec<Point3f>,
+        mut n: Vec<Normal3f>,
+        mut s: Vec<Vector3f>,
+        uv: Vec<Point2f>,
+        face_indices: Vec<i32>,
+    ) -> Self {
+        let n_triangles = vertex_indices.len() / 3;
+        let n_vertices = p.len();
+
+        // Transform positions to render space
+        for pt in p.iter_mut() {
+            *pt = render_from_object.apply_point(*pt);
+        }
+
+        // Transform and optionally flip normals
+        if !n.is_empty() {
+            assert_eq!(n_vertices, n.len(), "Normal count mismatch");
+            for nn in n.iter_mut() {
+                *nn = render_from_object.apply_normal(*nn);
+                if reverse_orientation {
+                    *nn = -*nn;
+                }
+            }
+        }
+
+        // Transform tangents
+        if !s.is_empty() {
+            assert_eq!(n_vertices, s.len(), "Tangent count mismatch");
+            for ss in s.iter_mut() {
+                *ss = render_from_object.apply_vector(*ss);
+            }
+        }
+
+        // Validate UVs
+        if !uv.is_empty() {
+            assert_eq!(n_vertices, uv.len(), "UV count mismatch");
+        }
+
+        // Validate face indices
+        if !face_indices.is_empty() {
+            assert_eq!(n_triangles, face_indices.len(), "Face index count mismatch");
+        }
+
+        let transform_swaps_handedness = render_from_object.swaps_handedness();
+
+        let storage = Box::new(TriangleMeshStorage {
+            vertex_indices,
+            p,
+            n,
+            s,
+            uv,
+            face_indices,
+        });
+
+        // Build device struct with pointers into storage
+        let device = DeviceTriangleMesh {
+            n_triangles: n_triangles as u32,
+            n_vertices: n_vertices as u32,
+            vertex_indices: storage.vertex_indices.as_ptr(),
+            p: storage.p.as_ptr(),
+            n: if storage.n.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.n.as_ptr()
+            },
+            s: if storage.s.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.s.as_ptr()
+            },
+            uv: if storage.uv.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.uv.as_ptr()
+            },
+            face_indices: if storage.face_indices.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.face_indices.as_ptr()
+            },
+            reverse_orientation,
+            transform_swaps_handedness,
+        };
+
+        Self {
+            _storage: storage,
+            device,
+        }
+    }
+
+    pub fn positions(&self) -> &[Point3f] {
+        &self._storage.p
+    }
+
+    pub fn indices(&self) -> &[u32] {
+        &self._storage.vertex_indices
+    }
+
+    pub fn normals(&self) -> &[Normal3f] {
+        &self._storage.n
+    }
+
+    pub fn uvs(&self) -> &[Point2f] {
+        &self._storage.uv
+    }
+}
+
+impl BilinearPatchMesh {
+    pub fn new(
+        render_from_object: &Transform,
+        reverse_orientation: bool,
+        vertex_indices: Vec<u32>,
+        mut p: Vec<Point3f>,
+        mut n: Vec<Normal3f>,
+        uv: Vec<Point2f>,
+        image_distribution: Option<PiecewiseConstant2D>,
+    ) -> Self {
+        let n_patches = vertex_indices.len() / 4;
+        let n_vertices = p.len();
+
+        for pt in p.iter_mut() {
+            *pt = render_from_object.apply_point(*pt);
+        }
+
+        if !n.is_empty() {
+            assert_eq!(n_vertices, n.len(), "Normal count mismatch");
+            for nn in n.iter_mut() {
+                *nn = render_from_object.apply_normal(*nn);
+                if reverse_orientation {
+                    *nn = -*nn;
+                }
+            }
+        }
+
+        if !uv.is_empty() {
+            assert_eq!(n_vertices, uv.len(), "UV count mismatch");
+        }
+
+        let transform_swaps_handedness = render_from_object.swaps_handedness();
+
+        let storage = Box::new(BilinearMeshStorage {
+            vertex_indices,
+            p,
+            n,
+            uv,
+            image_distribution,
+        });
+
+        let device = DeviceBilinearPatchMesh {
+            n_patches: n_patches as u32,
+            n_vertices: n_vertices as u32,
+            vertex_indices: storage.vertex_indices.as_ptr(),
+            p: storage.p.as_ptr(),
+            n: if storage.n.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.n.as_ptr()
+            },
+            uv: if storage.uv.is_empty() {
+                std::ptr::null()
+            } else {
+                storage.uv.as_ptr()
+            },
+            reverse_orientation,
+            transform_swaps_handedness,
+            image_distribution: storage
+                .image_distribution
+                .as_ref()
+                .map(|d| &d.device as *const _)
+                .unwrap_or(std::ptr::null()),
+        };
+
+        Self {
+            _storage: storage,
+            device,
+        }
+    }
+}
+
+// ============================================================================
+// PLY Helper Functions
+// ============================================================================
+
 fn get_float(elem: &DefaultElement, key: &str) -> AnyResult<f32> {
     match elem.get(key) {
         Some(Property::Float(v)) => Ok(*v),
@@ -49,30 +288,39 @@ fn get_float_any(elem: &DefaultElement, keys: &[&str]) -> Option<f32> {
 
 fn get_list_uint(elem: &DefaultElement, key: &str) -> AnyResult<Vec<u32>> {
     match elem.get(key) {
-        Some(Property::List_Int(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
-        Some(Property::List_UInt(vec)) => Ok(vec.clone()),
-        Some(Property::List_UChar(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
-        Some(Property::List_Char(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
+        Some(Property::ListInt(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
+        Some(Property::ListUInt(vec)) => Ok(vec.clone()),
+        Some(Property::ListUChar(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
+        Some(Property::ListChar(vec)) => Ok(vec.iter().map(|&x| x as u32).collect()),
         _ => bail!("Property {} is not a list", key),
     }
 }
 
+// ============================================================================
+// TriQuadMesh Implementation
+// ============================================================================
+
 impl TriQuadMesh {
     pub fn read_ply<P: AsRef<Path>>(filename: P) -> AnyResult<Self> {
         let path = filename.as_ref();
         let filename_display = path.display().to_string();
 
-        let mut f = File::open(&filename)
+        let mut f = File::open(path)
             .with_context(|| format!("Couldn't open PLY file \"{}\"", filename_display))?;
 
-        // Going to ply-rs
         let p = Parser::<DefaultElement>::new();
         let ply = p
             .read_ply(&mut f)
             .with_context(|| format!("Unable to read/parse PLY file \"{}\"", filename_display))?;
 
         let mut mesh = TriQuadMesh::default();
+
+        // Parse vertices
         if let Some(vertices) = ply.payload.get("vertex") {
+            if vertices.is_empty() {
+                bail!("{}: PLY file has no vertices", filename_display);
+            }
+
             let first = &vertices[0];
             let has_uv = (first.contains_key("u") && first.contains_key("v"))
                 || (first.contains_key("s") && first.contains_key("t"))
@@ -80,73 +328,77 @@ impl TriQuadMesh {
                 || (first.contains_key("texture_s") && first.contains_key("texture_t"));
             let has_normal =
                 first.contains_key("nx") && first.contains_key("ny") && first.contains_key("nz");
+
+            mesh.p.reserve(vertices.len());
+            if has_normal {
+                mesh.n.reserve(vertices.len());
+            }
+            if has_uv {
+                mesh.uv.reserve(vertices.len());
+            }
+
             for v_elem in vertices {
-                // Read Position
                 let x = get_float(v_elem, "x")?;
                 let y = get_float(v_elem, "y")?;
                 let z = get_float(v_elem, "z")?;
-                mesh.p.push([x, y, z]);
+                mesh.p.push(Point3f::new(x, y, z));
 
-                // Read Normal
                 if has_normal {
                     let nx = get_float(v_elem, "nx").unwrap_or(0.0);
                     let ny = get_float(v_elem, "ny").unwrap_or(0.0);
                     let nz = get_float(v_elem, "nz").unwrap_or(0.0);
-                    mesh.n.push([nx, ny, nz]);
+                    mesh.n.push(Normal3f::new(nx, ny, nz));
                 }
 
-                // Read UVs (Optional, handling variable naming convention)
                 if has_uv {
                     let u =
                         get_float_any(v_elem, &["u", "s", "texture_u", "texture_s"]).unwrap_or(0.0);
                     let v =
                         get_float_any(v_elem, &["v", "t", "texture_v", "texture_t"]).unwrap_or(0.0);
-                    mesh.uv.push([u, v]);
+                    mesh.uv.push(Point2f::new(u, v));
                 }
             }
         } else {
-            bail!(
-                "{}: PLY file is invalid! No vertex elements found!",
-                filename_display
-            );
+            bail!("{}: PLY file has no vertex elements", filename_display);
         }
 
+        // Parse faces
         if let Some(faces) = ply.payload.get("face") {
             mesh.tri_indices.reserve(faces.len() * 3);
             mesh.quad_indices.reserve(faces.len() * 4);
 
             for f_elem in faces {
                 if let Ok(fi) = get_int(f_elem, "face_indices") {
-                    mesh.faceIndices.push(fi);
+                    mesh.face_indices.push(fi);
                 }
 
                 if let Ok(indices) = get_list_uint(f_elem, "vertex_indices") {
                     match indices.len() {
-                        3 => {
-                            mesh.tri_indices.extend_from_slice(&indices);
+                        3 => mesh.tri_indices.extend_from_slice(&indices),
+                        4 => mesh.quad_indices.extend_from_slice(&indices),
+                        n => {
+                            log::warn!(
+                                "{}: Skipping face with {} vertices (only 3 or 4 supported)",
+                                filename_display,
+                                n
+                            );
                         }
-                        4 => {
-                            mesh.quad_indices.extend_from_slice(&indices);
-                        }
-                        _ => {}
                     }
                 } else {
-                    bail!("{}: vertex indices not found in PLY file", filename_display);
+                    bail!("{}: vertex_indices not found in face", filename_display);
                 }
             }
         } else {
-            bail!(
-                "{}: PLY file is invalid! No face elements found!",
-                filename_display
-            );
+            bail!("{}: PLY file has no face elements", filename_display);
         }
 
+        // Validate indices
         let vertex_count = mesh.p.len() as u32;
-
         for &idx in &mesh.tri_indices {
             if idx >= vertex_count {
                 bail!(
-                    "plymesh: Vertex index {} is out of bounds! Valid range is [0..{})",
+                    "{}: Vertex index {} out of bounds [0..{})",
+                    filename_display,
                     idx,
                     vertex_count
                 );
@@ -155,7 +407,8 @@ impl TriQuadMesh {
         for &idx in &mesh.quad_indices {
             if idx >= vertex_count {
                 bail!(
-                    "plymesh: Vertex index {} is out of bounds! Valid range is [0..{})",
+                    "{}: Vertex index {} out of bounds [0..{})",
+                    filename_display,
                     idx,
                     vertex_count
                 );
@@ -170,262 +423,110 @@ impl TriQuadMesh {
             return;
         }
 
-        for i in (0..self.quad_indices.len()).step_by(4) {
-            // First triangle: 0, 1, 3
-            self.tri_indices.push(self.quad_indices[i]);
-            self.tri_indices.push(self.quad_indices[i + 1]);
-            self.tri_indices.push(self.quad_indices[i + 3]);
+        // Each quad becomes 2 triangles
+        self.tri_indices.reserve(self.quad_indices.len() / 4 * 6);
 
-            // Second triangle: 0, 3, 2
-            self.tri_indices.push(self.quad_indices[i]);
-            self.tri_indices.push(self.quad_indices[i + 3]);
-            self.tri_indices.push(self.quad_indices[i + 2]);
+        for quad in self.quad_indices.chunks_exact(4) {
+            let (v0, v1, v2, v3) = (quad[0], quad[1], quad[2], quad[3]);
+
+            // Triangle 1: v0, v1, v2
+            self.tri_indices.push(v0);
+            self.tri_indices.push(v1);
+            self.tri_indices.push(v2);
+
+            // Triangle 2: v0, v2, v3
+            self.tri_indices.push(v0);
+            self.tri_indices.push(v2);
+            self.tri_indices.push(v3);
         }
 
         self.quad_indices.clear();
     }
 
     pub fn compute_normals(&mut self) {
-        self.n.resize(self.p.len(), Normal3f::zero());
-        for i in (0..self.tri_indices.len()).step_by(3) {
-            let v = vec![
-                self.tri_indices[i],
-                self.tri_indices[i + 1],
-                self.tri_indices[i + 2],
-            ];
-            let v10 = self.p[v[1]] - self.p[v[0]];
-            let v21 = self.p[v[2]] - self.p[v[1]];
+        // Initialize normals to zero
+        self.n = vec![Normal3f::new(0.0, 0.0, 0.0); self.p.len()];
 
-            let mut vn = v10.cross(v21);
-            if vn.norm_squared() > 0. {
-                vn = vn.normalize();
-                self.n[v[0]] += vn;
-                self.n[v[1]] += vn;
-                self.n[v[2]] += vn;
+        // Accumulate face normals for triangles
+        for tri in self.tri_indices.chunks_exact(3) {
+            let (i0, i1, i2) = (tri[0] as usize, tri[1] as usize, tri[2] as usize);
+
+            let p0 = self.p[i0];
+            let p1 = self.p[i1];
+            let p2 = self.p[i2];
+
+            let v10 = p1 - p0;
+            let v20 = p2 - p0;
+            let face_normal = v10.cross(&v20);
+
+            // Accumulate (will normalize later)
+            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);
+        }
+
+        // Accumulate face normals for quads
+        for quad in self.quad_indices.chunks_exact(4) {
+            let indices: Vec<usize> = quad.iter().map(|&i| i as usize).collect();
+
+            let p0 = self.p[indices[0]];
+            let p1 = self.p[indices[1]];
+            let p2 = self.p[indices[2]];
+
+            let v10 = p1 - p0;
+            let v20 = p2 - p0;
+            let face_normal = v10.cross(&v20);
+
+            for &idx in &indices {
+                self.n[idx] = self.n[idx] + Normal3f::from(face_normal);
             }
         }
 
-        assert!(self.quad_indices == 0.);
-        for i in 0..self.n.len() {
-            if n[i].normalize() > 0. {
-                self.n[i] = self.n[i].normalize()
+        // Normalize all normals
+        for normal in &mut self.n {
+            let len_sq = normal.length_squared();
+            if len_sq > 0.0 {
+                *normal = *normal / len_sq.sqrt();
             }
         }
     }
-}
 
-pub trait TriangleMeshFactory {
-    fn create(
-        arena: &mut Arena,
+    /// Convert to a TriangleMesh, consuming self
+    pub fn into_triangle_mesh(
+        mut self,
         render_from_object: &Transform,
         reverse_orientation: bool,
-        vertex_indices: Vec<u32>,
-        p: Vec<Point3f>,
-        n: Vec<Normal3f>,
-        s: Vec<Vector3f>,
-        uv: Vec<Point2f>,
-        face_indices: Vec<u32>,
-    ) -> ArenaPtr<TriangleMesh>;
-}
+    ) -> TriangleMesh {
+        self.convert_to_only_triangles();
 
-impl TriangleMeshFactory for TriangleMesh {
-    fn create(
-        arena: &mut Arena,
-        render_from_object: &Transform,
-        reverse_orientation: bool,
-        vertex_indices: Vec<u32>,
-        p: Vec<Point3f>,
-        n: Vec<Normal3f>,
-        s: Vec<Vector3f>,
-        uv: Vec<Point2f>,
-        face_indices: Vec<u32>,
-    ) -> ArenaPtr<TriangleMesh> {
-        let n_triangles = indices.len() / 3;
-        let n_vertices = p.len();
-        for pt in p.iter_mut() {
-            *pt = render_from_object.apply_to_point(*pt);
+        if self.n.is_empty() {
+            self.compute_normals();
         }
 
-        let transform_swaps_handedness = render_from_object.swaps_handedness();
-
-        let uv = if !uv.is_empty() {
-            assert_eq!(n_vertices, uv.len());
-            Some(uv)
-        } else {
-            None
-        };
-
-        let n = if !n.is_empty() {
-            assert_eq!(n_vertices, n.len());
-            for nn in n.iter_mut() {
-                *nn = render_from_object.apply_to_normal(*nn);
-                if reverse_orientation {
-                    *nn = -*nn;
-                }
-            }
-            Some(n)
-        } else {
-            None
-        };
-
-        let s = if !s.is_empty() {
-            assert_eq!(n_vertices, s.len());
-            for ss in s.iter_mut() {
-                *ss = render_from_object.apply_to_vector(*ss);
-            }
-            Some(s)
-        } else {
-            None
-        };
-
-        let face_indices = if !face_indices.is_empty() {
-            assert_eq!(n_triangles, face_indices.len());
-            Some(face_indices)
-        } else {
-            None
-        };
-
-        let storage = Box::new(TriangleMeshStorage {
-            p,
-            vertex_indices,
-            n,
-            s,
-            uv,
-            face_indices,
-        });
-
-        assert!(p.len() <= i32::MAX as usize);
-        assert!(indices.len() <= i32::MAX as usize);
-
-        let p_ptr = storage.p.as_ptr();
-        let idx_ptr = storage.vertex_indices.as_ptr();
-
-        let n_ptr = if storage.n.is_empty() {
-            ptr::null()
-        } else {
-            storage.n.as_ptr()
-        };
-
-        let uv_ptr = if storage.uv.is_empty() {
-            ptr::null()
-        } else {
-            storage.uv.as_ptr()
-        };
-
-        let s_ptr = if storage.s.is_empty() {
-            ptr::null()
-        } else {
-            storage.s.as_ptr()
-        };
-
-        let mesh = TriangleMeshHost::new(
+        TriangleMesh::new(
             render_from_object,
             reverse_orientation,
-            vertex_indices,
-            p,
-            s,
-            n,
-            uv,
-            face_indices,
-        );
+            self.tri_indices,
+            self.p,
+            self.n,
+            Vec::new(), // s (tangents)
+            self.uv,
+            self.face_indices,
+        )
     }
 }
 
-#[derive(Debug, Clone, Copy)]
-pub struct BilinearMeshStorage {
-    vertex_indices: Vec<u32>,
-    p: Vec<Point3f>,
-    n: Vec<Normal3f>,
-    uv: Vec<Point2f>,
-    image_distribution: Option<DevicePiecewiseConstant2D>,
-}
+// ============================================================================
+// Convenience: Create from PLY directly
+// ============================================================================
 
-#[derive(Debug, Clone, Copy)]
-pub struct BilinearPatchMeshHost {
-    pub view: BilinearPatchMesh,
-    _storage: Box<BilinearMeshStorage>,
-}
-
-#[derive(Debug, Clone, Copy)]
-pub struct TriangleMeshHost {
-    pub view: BilinearPatchMesh,
-    _storage: Box<BilinearMeshStorage>,
-}
-
-impl Deref for TriangleMeshHost {
-    type Target = TriangleMesh;
-    fn deref(&self) -> &Self::Target {
-        &self.view
-    }
-}
-
-impl BilinearPatchMeshHost {
-    pub fn new(
-        render_from_object: Transform,
+impl TriangleMesh {
+    pub fn from_ply<P: AsRef<Path>>(
+        filename: P,
+        render_from_object: &Transform,
         reverse_orientation: bool,
-        vertex_indices: Vec<usize>,
-        mut p: Vec<Point3f>,
-        mut n: Vec<Normal3f>,
-        uv: Vec<Point2f>,
-        image_distribution: Option<DevicePiecewiseConstant2D>,
-    ) -> Self {
-        let n_patches = indices.len() / 3;
-        let n_vertices = p.len();
-        for pt in p.iter_mut() {
-            *pt = render_from_object.apply_to_point(*pt);
-        }
-
-        if !n.is_empty() {
-            assert_eq!(n_vertices, n.len(), "Normal count mismatch");
-            for nn in n.iter_mut() {
-                *nn = render_from_object.apply_to_normal(*nn);
-                if reverse_orientation {
-                    *nn = -*nn;
-                }
-            }
-        }
-
-        if !uv.is_empty() {
-            assert_eq!(n_vertices, uv.len(), "UV count mismatch");
-        }
-
-        let storage = Box::new(BilinearMeshStorage {
-            vertex_indices,
-            p,
-            n,
-            uv,
-            image_distribution,
-        });
-
-        let transform_swaps_handedness = render_from_object.swaps_handedness();
-
-        let p_ptr = storage.p.as_ptr();
-        let idx_ptr = storage.vertex_indices.as_ptr();
-
-        let n_ptr = if storage.n.is_empty() {
-            ptr::null()
-        } else {
-            storage.n.as_ptr()
-        };
-
-        let uv_ptr = if storage.uv.is_empty() {
-            ptr::null()
-        } else {
-            storage.uv.as_ptr()
-        };
-
-        let view = BilinearPatchMesh {
-            n_patches: n_patches as u32,
-            n_vertices: n_vertices as u32,
-            vertex_indices: idx_ptr,
-            p: p_ptr,
-            n: n_ptr,
-            uv: uv_ptr,
-            reverse_orientation,
-            transform_swaps_handedness,
-            image_distribution: dist_ptr,
-        };
-
-        Self { view, _storage }
+    ) -> AnyResult<Self> {
+        let mesh = TriQuadMesh::read_ply(filename)?;
+        Ok(mesh.into_triangle_mesh(render_from_object, reverse_orientation))
     }
 }

src/shapes/mod.rs

@@ -12,5 +12,5 @@ pub use mesh::*;
 
 use std::sync::{Arc, Mutex};
 
-pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<TriangleMeshHost>>> = Mutex::new(Vec::new());
-pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<BilinearPatchMeshHost>>> = Mutex::new(Vec::new());
+pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<TriangleMesh>>> = Mutex::new(Vec::new());
+pub static ALL_TRIANGLE_MESHES: Mutex<Vec<Arc<BilinearPatchMesh>>> = Mutex::new(Vec::new());

src/shapes/sphere.rs

@@ -4,10 +4,10 @@ use crate::utils::{Arena, FileLoc, ParameterDictionary};
 use shared::core::shape::Shape;
 use shared::shapes::SphereShape;
 use shared::utils::Transform;
+use std::collections::HashMap;
 
 impl CreateShape for SphereShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         object_from_render: Transform,
         reverse_orientation: bool,
@@ -21,15 +21,15 @@ impl CreateShape for SphereShape {
         let zmax = parameters.get_one_float("zmax", radius);
         let phimax = parameters.get_one_float("phimax", 360.);
         let shape = SphereShape::new(
-            renderFromObject,
-            objectFromRender,
-            reverseOrientation,
+            render_from_object,
+            object_from_render,
+            reverse_orientation,
             radius,
             zmin,
             zmax,
             phimax,
         );
-        arena.alloc(vec![Shape::Sphere(SphereShape)]);
-        Ok(vec![Shape::Sphere(SphereShape)])
+        arena.alloc(vec![Shape::Sphere(shape)]);
+        Ok(vec![Shape::Sphere(shape)])
     }
 }

src/shapes/triangle.rs

@@ -1,14 +1,16 @@
 use crate::core::shape::{ALL_TRIANGLE_MESHES, CreateShape};
 use crate::core::texture::FloatTexture;
-use crate::shapes::mesh::TriangleMeshHost;
+use crate::shapes::mesh::TriangleMesh;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
 use log::warn;
+use shared::core::shape::Shape;
 use shared::shapes::TriangleShape;
 use shared::utils::Transform;
+use std::collections::HashMap;
+use std::sync::Arc;
 
 impl CreateShape for TriangleShape {
     fn create(
-        name: &str,
         render_from_object: Transform,
         _object_from_render: Transform,
         reverse_orientation: bool,
@@ -16,7 +18,7 @@ impl CreateShape for TriangleShape {
         _float_texture: HashMap<String, FloatTexture>,
         loc: FileLoc,
         arena: &mut Arena,
-    ) -> Result<Mesh, String> {
+    ) -> Result<Vec<Shape>, String> {
         let mut vertex_indices = parameters.get_int_array("indices");
         let mut p = parameters.get_point3f_array("P");
         let mut uvs = parameters.get_point2f_array("uv");
@@ -34,7 +36,7 @@ impl CreateShape for TriangleShape {
             let excess = vertex_indices.len() % 3;
             warn!(
                 "Number of vertex indices {} not a multiple of 3. Discarding {} excess.",
-                vi.len(),
+                vertex_indices.len(),
                 excess
             );
             let new_len = vertex_indices.len() - excess;
@@ -78,12 +80,12 @@ impl CreateShape for TriangleShape {
             warn!(
                 "Number of face indices {} does not match number of triangles {}. Discarding face indices.",
                 face_indices.len(),
-                num_triangles
+                n_triangles
             );
             face_indices.clear();
         }
 
-        let host = TriangleMeshHost::new(
+        let host = TriangleMesh::new(
             render_from_object,
             reverse_orientation,
             vertex_indices,

src/spectra/colorspace.rs

@@ -1,8 +1,9 @@
 use super::DenselySampledSpectrumBuffer;
-use shared::core::color::RGBToSpectrumTable;
+use shared::core::color::{RGB, RGBToSpectrumTable, XYZ};
 use shared::core::geometry::Point2f;
 use shared::spectra::RGBColorSpace;
 use shared::utils::math::SquareMatrix;
+use shared::utils::ptr::Ptr;
 
 pub struct RGBColorSpaceData {
     _illuminant: DenselySampledSpectrumBuffer,
@@ -15,7 +16,7 @@ impl RGBColorSpaceData {
         g: Point2f,
         b: Point2f,
         illuminant: DenselySampledSpectrumBuffer,
-        rgb_to_spectrum_table: *const RGBToSpectrumTable,
+        rgb_to_spectrum_table: Ptr<RGBToSpectrumTable>,
     ) -> Self {
         let w_xyz: XYZ = illuminant.to_xyz();
         let w = w_xyz.xy();
@@ -38,24 +39,14 @@ impl RGBColorSpaceData {
             g,
             b,
             w,
-            illuminant: illuminant_buffer.view, // Extract view
+            illuminant: illuminant.view,
             xyz_from_rgb,
             rgb_from_xyz,
-            rgb_to_spectrum_table: table_ptr,
+            rgb_to_spectrum_table,
         };
         Self {
-            _illuminant: illuminant_buffer,
+            _illuminant: illuminant,
             view,
         }
     }
 }
-
-pub fn get_named(name: &str) -> Result<Arc<RGBColorSpace>, String> {
-    match name.to_lowercase().as_str() {
-        "aces2065-1" => Ok(Self::aces2065_1().clone()),
-        "rec2020" => Ok(Self::rec2020().clone()),
-        "dci-p3" => Ok(Self::dci_p3().clone()),
-        "srgb" => Ok(Self::srgb().clone()),
-        _ => Err(format!("Color space '{}' not found", name)),
-    }
-}

src/spectra/data.rs

@@ -1,7 +1,10 @@
-use crate::spectra::DenselySampledSpectrumBuffer;
+use crate::spectra::{DenselySampledSpectrumBuffer, piecewise::PiecewiseLinearSpectrumBuffer};
 use shared::Float;
 use shared::core::spectrum::Spectrum;
 use shared::spectra::PiecewiseLinearSpectrum;
+use shared::spectra::cie::*;
+use std::collections::HashMap;
+use std::sync::LazyLock;
 
 pub fn create_cie_buffer(data: &[Float]) -> Spectrum {
     let buffer = PiecewiseLinearSpectrum::from_interleaved(data, false);
@@ -150,6 +153,6 @@ pub static NAMED_SPECTRA: LazyLock<HashMap<String, Spectrum>> = LazyLock::new(||
 });
 
 pub fn get_named_spectrum(name: &str) -> Option<Spectrum> {
-    let buffer = NAMED_SPECTRA_DATA.get(name)?;
+    let buffer = NAMED_SPECTRA.get(name)?;
     Some(Spectrum::PiecewiseLinear(buffer.view))
 }

src/spectra/dense.rs

@@ -1,9 +1,15 @@
 use shared::Float;
-use shared::spectra::{DenselySampledSpectrum, SampledSpectrum};
+use shared::core::spectrum::Spectrum;
+use shared::spectra::cie::{CIE_S_LAMBDA, CIE_S0, CIE_S1, CIE_S2, N_CIES};
+use shared::spectra::{
+    BlackbodySpectrum, DenselySampledSpectrum, LAMBDA_MAX, LAMBDA_MIN, PiecewiseLinearSpectrum,
+    SampledSpectrum, SampledWavelengths,
+};
+use shared::utils::math::square;
 
 pub struct DenselySampledSpectrumBuffer {
     pub device: DenselySampledSpectrum,
-    _storage: Vec<Float>,
+    pub _storage: Vec<Float>,
 }
 
 impl std::ops::Deref for DenselySampledSpectrumBuffer {

src/spectra/mod.rs

@@ -1,10 +1,13 @@
 use crate::spectra::colorspace::RGBColorSpaceData;
-use once_cell::unsync::Lazy;
+use shared::core::color::RGBToSpectrumTable;
 use shared::core::geometry::Point2f;
 use shared::core::spectrum::Spectrum;
+use shared::core::spectrum::StandardSpectra;
+use shared::spectra::RGBColorSpace;
+use shared::spectra::StandardColorSpaces;
 use shared::spectra::cie::{CIE_D65, CIE_X, CIE_Y, CIE_Z};
-use shared::spectra::{RGBColorSpace, RGBToSpectrumTable, StandardColorSpaces};
 use std::sync::Arc;
+use std::sync::LazyLock;
 
 pub mod colorspace;
 pub mod data;
@@ -14,14 +17,14 @@ pub mod piecewise;
 // pub use data;
 pub use dense::DenselySampledSpectrumBuffer;
 
-static CIE_X_DATA: Lazy<DenselySampledSpectrumBuffer> =
-    Lazy::new(|| data::create_cie_buffer(&CIE_X));
-static CIE_Y_DATA: Lazy<DenselySampledSpectrumBuffer> =
-    Lazy::new(|| data::create_cie_buffer(&CIE_Y));
-static CIE_Z_DATA: Lazy<DenselySampledSpectrumBuffer> =
-    Lazy::new(|| data::create_cie_buffer(&CIE_Z));
-static CIE_D65_DATA: Lazy<DenselySampledSpectrumBuffer> =
-    Lazy::new(|| data::create_cie_buffer(&CIE_D65));
+static CIE_X_DATA: LazyLock<DenselySampledSpectrumBuffer> =
+    LazyLock::new(|| data::create_cie_buffer(&CIE_X));
+static CIE_Y_DATA: LazyLock<DenselySampledSpectrumBuffer> =
+    LazyLock::new(|| data::create_cie_buffer(&CIE_Y));
+static CIE_Z_DATA: LazyLock<DenselySampledSpectrumBuffer> =
+    LazyLock::new(|| data::create_cie_buffer(&CIE_Z));
+static CIE_D65_DATA: LazyLock<DenselySampledSpectrumBuffer> =
+    LazyLock::new(|| data::create_cie_buffer(&CIE_D65));
 
 pub fn cie_x() -> Spectrum {
     Spectrum::DenselySampled(CIE_X_DATA.view)
@@ -45,15 +48,14 @@ pub fn get_spectra_context() -> StandardSpectra {
         y: CIE_Y_DATA.view,
         z: CIE_Z_DATA.view,
         d65: CIE_D65_DATA.view,
-        cie_y_integral: CIE_Y_INTEGRAL,
     }
 }
 
-pub static SRGB_DATA: Lazy<RGBColorSpaceData> = Lazy::new(|| {
+pub static SRGB: LazyLock<RGBColorSpaceData> = LazyLock::new(|| {
     let illum = DenselySampledSpectrumBuffer::new(
-        D65_BUFFER.view.lambda_min,
-        D65_BUFFER.view.lambda_max,
-        D65_BUFFER._storage.clone(),
+        CIE_D65_DATA.view.lambda_min,
+        CIE_D65_DATA.view.lambda_max,
+        CIE_D65_DATA._storage.clone(),
     );
 
     RGBColorSpaceData::new(
@@ -65,11 +67,11 @@ pub static SRGB_DATA: Lazy<RGBColorSpaceData> = Lazy::new(|| {
     )
 });
 
-pub static DCI_P3: Lazy<RGBColorSpaceData> = Lazy::new(|| {
+pub static DCI_P3: LazyLock<RGBColorSpaceData> = LazyLock::new(|| {
     let illum = DenselySampledSpectrumBuffer::new(
-        D65_BUFFER.view.lambda_min,
-        D65_BUFFER.view.lambda_max,
-        D65_BUFFER._storage.clone(),
+        CIE_D65_DATA.view.lambda_min,
+        CIE_D65_DATA.view.lambda_max,
+        CIE_D65_DATA._storage.clone(),
     );
 
     let r = Point2f::new(0.680, 0.320);
@@ -79,11 +81,11 @@ pub static DCI_P3: Lazy<RGBColorSpaceData> = Lazy::new(|| {
     RGBColorSpaceData::new(r, g, b, illum, RGBToSpectrumTable)
 });
 
-pub static REC2020: Lazy<Arc<RGBColorSpace>> = Lazy::new(|| {
+pub static REC2020: LazyLock<Arc<RGBColorSpace>> = LazyLock::new(|| {
     let illum = DenselySampledSpectrumBuffer::new(
-        D65_BUFFER.view.lambda_min,
-        D65_BUFFER.view.lambda_max,
-        D65_BUFFER._storage.clone(),
+        CIE_D65_DATA.view.lambda_min,
+        CIE_D65_DATA.view.lambda_max,
+        CIE_D65_DATA._storage.clone(),
     );
 
     let r = Point2f::new(0.708, 0.292);
@@ -95,7 +97,7 @@ pub static REC2020: Lazy<Arc<RGBColorSpace>> = Lazy::new(|| {
     RGBColorSpace::new(r, g, b, illum, table)
 });
 
-pub static ACES: Lazy<Arc<RGBColorSpace>> = Lazy::new(|| {
+pub static ACES: LazyLock<Arc<RGBColorSpace>> = LazyLock::new(|| {
     let r = Point2f::new(0.7347, 0.2653);
     let g = Point2f::new(0.0000, 1.0000);
     let b = Point2f::new(0.0001, -0.0770);
@@ -108,9 +110,9 @@ pub static ACES: Lazy<Arc<RGBColorSpace>> = Lazy::new(|| {
 
 pub fn get_colorspace_context() -> StandardColorSpaces {
     StandardColorSpaces {
-        srgb: &SRGB_DATA.view,
-        dci_p3: &DCI_P3_DATA.view,
-        rec2020: &REC2020.view,
-        aces2065_1: &ACES.view,
+        srgb: SRGB,
+        dci_p3: DCI_P3,
+        rec2020: REC2020,
+        aces2065_1: ACES,
     }
 }

src/spectra/piecewise.rs

@@ -1,6 +1,8 @@
 use crate::utils::read_float_file;
 use shared::Float;
 use shared::spectra::PiecewiseLinearSpectrum;
+use std::ops::Deref;
+use std::sync::atomic::Ordering;
 
 pub struct PiecewiseLinearSpectrumBuffer {
     pub view: PiecewiseLinearSpectrum,

src/textures/image.rs

@@ -1,10 +1,11 @@
+use crate::core::texture::{FloatTextureTrait, SpectrumTextureTrait};
 use crate::core::texture::{TexInfo, get_texture_cache};
 use crate::utils::mipmap::{MIPMap, MIPMapFilterOptions};
 use shared::Float;
 use shared::core::color::ColorEncoding;
 use shared::core::color::RGB;
 use shared::core::geometry::Vector2f;
-use shared::core::image::{WrapMode, WrapMode2D};
+use shared::core::image::WrapMode;
 use shared::core::texture::{SpectrumType, TextureEvalContext, TextureMapping2D};
 use shared::spectra::{
     RGBAlbedoSpectrum, RGBIlluminantSpectrum, RGBUnboundedSpectrum, SampledSpectrum,
@@ -174,4 +175,8 @@ impl FloatImageTexture {
     }
 }
 
-impl FloatTextureTrait for FloatImageTexture {}
+impl FloatTextureTrait for FloatImageTexture {
+    fn evaluate(&self, _ctx: &TextureEvalContext) -> Float {
+        todo!()
+    }
+}

src/textures/mix.rs

@@ -1,8 +1,10 @@
 use crate::core::texture::FloatTexture;
+use crate::core::texture::{FloatTextureTrait, SpectrumTextureTrait};
 use crate::utils::{Arena, FileLoc, TextureParameterDictionary};
 use shared::Float;
 use shared::core::geometry::Vector3f;
-use shared::core::texture::{FloatTextureTrait, SpectrumTextureTrait, TextureEvalContext};
+use shared::core::texture::TextureEvalContext;
+use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use shared::utils::Transform;
 use std::sync::Arc;
 
@@ -76,12 +78,24 @@ impl FloatDirectionMixTexture {
     }
 }
 
-impl FloatTextureTrait for FloatDirectionMixTexture {}
+impl FloatTextureTrait for FloatDirectionMixTexture {
+    fn evaluate(&self, _ctx: &TextureEvalContext) -> Float {
+        todo!()
+    }
+}
 
 #[derive(Debug, Clone)]
 pub struct SpectrumMixTexture;
-impl SpectrumTextureTrait for SpectrumMixTexture {}
+impl SpectrumTextureTrait for SpectrumMixTexture {
+    fn evaluate(&self, _ctx: &TextureEvalContext, _lambda: &SampledWavelengths) -> SampledSpectrum {
+        todo!()
+    }
+}
 
 #[derive(Debug, Clone)]
 pub struct SpectrumDirectionMixTexture;
-impl SpectrumTextureTrait for SpectrumDirectionMixTexture {}
+impl SpectrumTextureTrait for SpectrumDirectionMixTexture {
+    fn evaluate(&self, _ctx: &TextureEvalContext, _lambda: &SampledWavelengths) -> SampledSpectrum {
+        todo!()
+    }
+}

src/textures/ptex.rs

@@ -1,13 +1,11 @@
 // use crate::utils::{Arena, FileLoc, TextureParameterDictionary};
-use shared::core::color::{ColorEncoding, RGB};
-use shared::core::spectrum::SpectrumTextureTrait;
-use shared::core::texture::{SpectrumType, TextureEvalContext};
-use shared::Float;
-use shared::spectra::{SampledSpectrum, SampledWavelengths};
-
-use image_rs::imageops::FilterType;
+use crate::core::texture::{FloatTextureTrait, SpectrumTextureTrait};
 use ptex::Cache;
-use ptex_sys::ffi;
+use ptex_filter::{PtexFilterOptions, PtexFilterType, ptex_filter_create};
+use shared::Float;
+use shared::core::color::{ColorEncoding, RGB};
+use shared::core::texture::{SpectrumType, TextureEvalContext};
+use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use std::sync::OnceLock;
 
 struct PtexCache(Cache);
@@ -38,7 +36,7 @@ pub struct PtexTextureBase {
 
 impl PtexTextureBase {
     pub fn new(filename: String, encoding: ColorEncoding, scale: Float) -> Self {
-        let cache = get_ptex_cache();
+        let cache: &mut Cache = get_ptex_cache();
 
         // Attempt to get the texture to verify it exists and is valid
         let (valid, num_channels) = match cache.get(&filename) {
@@ -95,15 +93,15 @@ impl PtexTextureBase {
         let nc = texture.num_channels();
         let mut result = [0.0; 4];
         unsafe {
-            let opts = ffi::PtexFilter_Options {
-                filter: FilterType::f_bspline,
+            let opts = PtexFilterOptions {
+                filter: PtexFilterType::BSpline,
                 lerp: 1,
                 sharpness: 0.0,
                 noedgeblend: 0,
-                __structSize: std::mem::size_of::<ffi::PtexFilter_Options>() as i32,
+                // __structSize: std::mem::size_of::<PtexFilterOptions>() as i32,
             };
 
-            let filter_ptr = ffi::PtexFilter_getFilter(texture.as_ptr(), &opts);
+            let filter_ptr = ptex_filter_create(texture.as_ptr(), &opts);
             if filter_ptr.is_null() {
                 return None;
             }
@@ -146,12 +144,11 @@ pub struct FloatPtexTexture {
     pub base: PtexTextureBase,
 }
 
-// impl FloatPtexTexture {
-//     pub fn evaluate(&self, _ctx: &TextureEvalContext) -> Float {
-//         let
-//
-//     }
-// }
+impl FloatTextureTrait for FloatPtexTexture {
+    fn evaluate(&self, _ctx: &TextureEvalContext) -> Float {
+        todo!()
+    }
+}
 
 #[derive(Clone, Debug)]
 pub struct SpectrumPtexTexture {

src/textures/scaled.rs

@@ -1,9 +1,10 @@
 use crate::core::texture::{FloatTexture, SpectrumTexture};
+use crate::core::texture::{FloatTextureTrait, SpectrumTextureTrait};
 use crate::utils::{Arena, FileLoc, TextureParameterDictionary};
-use shared::core::texture::{FloatTextureTrait, SpectrumTextureTrait, TextureEvalContext};
+use shared::Float;
+use shared::core::texture::TextureEvalContext;
 use shared::spectra::{SampledSpectrum, SampledWavelengths};
 use shared::utils::Transform;
-use shared::Float;
 use std::sync::Arc;
 
 #[derive(Debug, Clone)]

src/utils/arena.rs

@@ -1,13 +1,16 @@
 use crate::core::image::Image;
 use crate::core::texture::{FloatTexture, SpectrumTexture};
+use crate::shapes::{BilinearPatchMesh, TriangleMesh};
 use crate::utils::sampling::PiecewiseConstant2D;
 use shared::core::color::RGBToSpectrumTable;
+use shared::core::image::DeviceImage;
 use shared::core::shape::Shape;
 use shared::core::spectrum::Spectrum;
 use shared::core::texture::{GPUFloatTexture, GPUSpectrumTexture};
 use shared::spectra::{RGBColorSpace, StandardColorSpaces};
 use shared::textures::*;
 use shared::utils::Ptr;
+use shared::utils::mesh::{DeviceBilinearPatchMesh, DeviceTriangleMesh};
 use shared::utils::sampling::{DevicePiecewiseConstant1D, DevicePiecewiseConstant2D};
 use std::alloc::Layout;
 use std::sync::Arc;
@@ -115,7 +118,7 @@ impl Upload for Shape {
 }
 
 impl Upload for Image {
-    type Target = Image;
+    type Target = DeviceImage;
     fn upload(&self, arena: &mut Arena) -> Ptr<Self::Target> {
         arena.alloc(self.clone())
     }
@@ -323,16 +326,104 @@ impl Upload for PiecewiseConstant2D {
         let conditionals_ptr = arena.alloc_slice(&conditionals_shared);
 
         let shared_2d = DevicePiecewiseConstant2D {
-            domain: self.domain,
-            p_marginal: marginal_shared,
-            n_conditionals: self.p_conditionals.len(),
-            p_conditional_v: conditionals_ptr,
+            conditionals: conditionals_ptr,
+            marginal: marginal_shared,
+            ..self.device
         };
 
         arena.alloc(shared_2d)
     }
 }
 
+impl Upload for TriangleMesh {
+    type Target = DeviceTriangleMesh;
+
+    fn upload(&self, arena: &mut Arena) -> Ptr<Self::Target> {
+        let storage = &self._storage;
+
+        // Upload all arrays to arena
+        let vertex_indices_ptr = arena.alloc_slice(&storage.vertex_indices);
+        let p_ptr = arena.alloc_slice(&storage.p);
+
+        let n_ptr = if storage.n.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.n).as_ptr()
+        };
+
+        let s_ptr = if storage.s.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.s).as_ptr()
+        };
+
+        let uv_ptr = if storage.uv.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.uv).as_ptr()
+        };
+
+        let face_indices_ptr = if storage.face_indices.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.face_indices).as_ptr()
+        };
+
+        let device = DeviceTriangleMesh {
+            vertex_indices: vertex_indices_ptr.as_ptr(),
+            p: p_ptr.as_ptr(),
+            n: n_ptr,
+            s: s_ptr,
+            uv: uv_ptr,
+            face_indices: face_indices_ptr,
+            ..self.device // Copy n_triangles, n_vertices, flags
+        };
+
+        arena.alloc(device)
+    }
+}
+
+impl Upload for BilinearPatchMesh {
+    type Target = DeviceBilinearPatchMesh;
+
+    fn upload(&self, arena: &mut Arena) -> Ptr<Self::Target> {
+        let storage = &self._storage;
+
+        let vertex_indices_ptr = arena.alloc_slice(&storage.vertex_indices);
+        let p_ptr = arena.alloc_slice(&storage.p);
+
+        let n_ptr = if storage.n.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.n).as_ptr()
+        };
+
+        let uv_ptr = if storage.uv.is_empty() {
+            std::ptr::null()
+        } else {
+            arena.alloc_slice(&storage.uv).as_ptr()
+        };
+
+        let image_dist_ptr = if let Some(ref dist) = storage.image_distribution {
+            let uploaded = dist.upload(arena);
+            uploaded.as_ptr()
+        } else {
+            std::ptr::null()
+        };
+
+        let device = DeviceBilinearPatchMesh {
+            vertex_indices: vertex_indices_ptr.as_ptr(),
+            p: p_ptr.as_ptr(),
+            n: n_ptr,
+            uv: uv_ptr,
+            image_distribution: image_dist_ptr,
+            ..self.device
+        };
+
+        arena.alloc(device)
+    }
+}
+
 impl<T: Upload> Upload for Option<T> {
     type Target = T::Target;
     fn upload(&self, arena: &mut Arena) -> Ptr<Self::Target> {
@@ -350,14 +441,3 @@ impl<T: Upload> Upload for Arc<T> {
         (**self).upload(arena)
     }
 }
-
-impl<T: Upload> Upload for Option<Arc<T>> {
-    type Target = T::Target;
-
-    fn upload(&self, arena: &mut Arena) -> Ptr<Self::Target> {
-        match self {
-            Some(v) => v.upload(arena),
-            None => Ptr::null(),
-        }
-    }
-}

src/utils/parameters.rs

@@ -214,7 +214,7 @@ impl PBRTParameter for String {
     const TYPE_NAME: &'static str = "string";
     const N_PER_ITEM: usize = 1;
     fn convert(v: &[Self::Raw]) -> Self {
-        v[0]
+        v[0].clone()
     }
     fn get_values(param: &ParsedParameter) -> &[Self::Raw] {
         param.strings
@@ -238,11 +238,12 @@ pub struct ParameterDictionary {
 
 impl ParameterDictionary {
     pub fn new(mut params: ParsedParameterVector, color_space: Option<Arc<RGBColorSpace>>) -> Self {
+        let n_owned_params = params.len();
         params.reverse();
         let dict = Self {
             params,
             color_space,
-            n_owned_params: params.len(),
+            n_owned_params,
         };
         dict.check_parameter_types();
         dict

src/utils/sampling.rs

@@ -149,7 +149,7 @@ impl PiecewiseConstant2D {
             .into_boxed_slice();
 
         let device = DevicePiecewiseConstant2D {
-            conditional: conditional_devices.as_ptr(),
+            conditionals: conditional_devices.as_ptr(),
             marginal: marginal.device,
             n_u: n_u as u32,
             n_v: n_v as u32,