Cleanup

Cargo.toml

@@ -8,7 +8,7 @@ default = []
 use_f64 = []
 use_gpu = []
 use_nvtx = []
-cuda = ["dep:cudarc"]
+cuda = ["dep:cudarc", "dep:cust", "dep:cust_raw"]
 
 [dependencies]
 anyhow = "1.0.100"
@@ -36,6 +36,7 @@ ptex-filter = { path = "crates/ptex-filter" }
 
 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 }
+cust_raw = { git = "https://github.com/Rust-GPU/Rust-CUDA", branch = "main", default-features = false, optional = true }
 ptex = "0.3.0"
 # ptex-sys = "0.3.0"
 slice = "0.0.4"

shared/src/core/film.rs

@@ -366,6 +366,10 @@ impl SpectralFilm {
     pub fn add_splat(&mut self, _p: Point2f, _v: SampledSpectrum, _lambda: &SampledWavelengths) {
         todo!()
     }
+
+    pub fn get_pixel_rgb(&self, _p: Point2i, _splat_scale: Option<Float>) -> RGB {
+        todo!()
+    }
 }
 
 #[repr(C)]

shared/src/textures/scaled.rs

@@ -12,11 +12,11 @@ pub struct GPUFloatScaledTexture {
 
 impl GPUFloatScaledTexture {
     pub fn evaluate(&self, ctx: &TextureEvalContext) -> Float {
-        let sc = self.scale.get().map(|t| t.evaluate(&ctx)).unwrap();
+        let sc = self.scale.get().map(|t| t.evaluate(ctx)).unwrap();
         if sc == 0. {
             return 0.;
         }
-        self.tex.get().map(|t| t.evaluate(&ctx)).unwrap_or(0.0) * sc
+        self.tex.get().map(|t| t.evaluate(ctx)).unwrap_or(0.0) * sc
     }
 }
 
@@ -33,11 +33,11 @@ impl GPUSpectrumScaledTexture {
         ctx: &TextureEvalContext,
         lambda: &SampledWavelengths,
     ) -> SampledSpectrum {
-        let sc = self.scale.get().map(|t| t.evaluate(&ctx)).unwrap_or(0.);
+        let sc = self.scale.get().map(|t| t.evaluate(ctx)).unwrap_or(0.);
 
         self.tex
             .get()
-            .map(|t| t.evaluate(&ctx, &lambda))
+            .map(|t| t.evaluate(ctx, lambda))
             .unwrap_or(SampledSpectrum::new(0.))
             * sc
     }

src/core/camera.rs

@@ -1,7 +1,7 @@
 use crate::core::image::ImageMetadata;
 use crate::core::image::{Image, ImageIO};
-use crate::utils::read_float_file;
 use crate::utils::{Arena, FileLoc, ParameterDictionary};
+use crate::utils::{Upload, read_float_file};
 use anyhow::{Result, anyhow};
 use shared::Ptr;
 use shared::cameras::*;
@@ -380,7 +380,7 @@ impl CameraFactory for Camera {
                     &lens_params,
                     focal_distance,
                     aperture_diameter,
-                    aperture_image,
+                    aperture_image.upload(arena),
                 );
 
                 arena.alloc(camera);

src/core/film.rs

@@ -2,6 +2,7 @@ use crate::Arena;
 use crate::core::image::{Image, ImageChannelDesc, ImageChannelValues, ImageIO, ImageMetadata};
 use crate::films::*;
 use crate::spectra::data::get_named_spectrum;
+use crate::spectra::piecewise::PiecewiseLinearSpectrumBuffer;
 use anyhow::{Result, anyhow};
 use rayon::iter::ParallelIterator;
 use rayon::prelude::IntoParallelIterator;
@@ -12,7 +13,7 @@ use shared::core::filter::{Filter, FilterTrait};
 use shared::core::geometry::{Bounds2f, Bounds2i, Point2f, Point2i};
 use shared::core::image::PixelFormat;
 use shared::core::spectrum::Spectrum;
-use shared::spectra::{PiecewiseLinearSpectrum, RGBColorSpace, cie::SWATCHES_RAW};
+use shared::spectra::{RGBColorSpace, cie::SWATCHES_RAW};
 use shared::utils::math::{SquareMatrix, linear_least_squares};
 use shared::{Float, Ptr};
 use std::sync::atomic::{AtomicUsize, Ordering};
@@ -27,17 +28,26 @@ const N_SWATCH_REFLECTANCES: usize = 24;
 const SWATCH_REFLECTANCES: LazyLock<[Spectrum; N_SWATCH_REFLECTANCES]> = LazyLock::new(|| {
     std::array::from_fn(|i| {
         let raw_data = SWATCHES_RAW[i];
-        let pls = PiecewiseLinearSpectrum::from_interleaved(raw_data, false);
+        let pls = PiecewiseLinearSpectrumBuffer::from_interleaved(raw_data, false);
-        Spectrum::Piecewise(pls)
+        Spectrum::Piecewise(pls.device)
     })
 });
 
-pub trait PixelSensorTrait: Sized {
+#[derive(Debug, Clone)]
-    fn get_swatches() -> Arc<[Spectrum; N_SWATCH_REFLECTANCES]> {
+struct SensorStorage {
-        Arc::new(*SWATCH_REFLECTANCES)
+    r_bar: DenselySampledSpectrumBuffer,
-    }
+    g_bar: DenselySampledSpectrumBuffer,
+    b_bar: DenselySampledSpectrumBuffer,
+}
 
-    fn create(
+#[derive(Debug, Clone)]
+pub struct PixelSensor {
+    device: DevicePixelSensor,
+    data: SensorStorage,
+}
+
+impl PixelSensor {
+    pub fn create(
         params: &ParameterDictionary,
         output_colorspace: Arc<RGBColorSpace>,
         exposure_time: Float,
@@ -103,23 +113,6 @@ pub trait PixelSensorTrait: Sized {
         }
     }
 
-    fn new(
-        r: &Spectrum,
-        g: &Spectrum,
-        b: &Spectrum,
-        output_colorspace: Arc<RGBColorSpace>,
-        sensor_illum: Option<&Spectrum>,
-        imaging_ratio: Float,
-    ) -> Self;
-
-    fn new_with_white_balance(
-        output_colorspace: &RGBColorSpace,
-        sensor_illum: Option<&Spectrum>,
-        imaging_ratio: Float,
-    ) -> Self;
-}
-
-impl PixelSensorTrait for DevicePixelSensor {
     fn new(
         r: &Spectrum,
         g: &Spectrum,
@@ -136,9 +129,9 @@ impl PixelSensorTrait for DevicePixelSensor {
             None => &Spectrum::Dense(output_colorspace.as_ref().illuminant),
         };
 
-        let r_bar = DenselySampledSpectrumBuffer::from_spectrum(r).device();
+        let r_bar = DenselySampledSpectrumBuffer::from_spectrum(r);
-        let g_bar = DenselySampledSpectrumBuffer::from_spectrum(g).device();
+        let g_bar = DenselySampledSpectrumBuffer::from_spectrum(g);
-        let b_bar = DenselySampledSpectrumBuffer::from_spectrum(b).device();
+        let b_bar = DenselySampledSpectrumBuffer::from_spectrum(b);
         let mut rgb_camera = [[0.; 3]; N_SWATCH_REFLECTANCES];
 
         let swatches = Self::get_swatches();
@@ -147,9 +140,9 @@ impl PixelSensorTrait for DevicePixelSensor {
             let rgb = DevicePixelSensor::project_reflectance::<RGB>(
                 &swatches[i],
                 illum,
-                &Spectrum::Dense(r_bar),
+                &Spectrum::Dense(r_bar.device()),
-                &Spectrum::Dense(g_bar),
+                &Spectrum::Dense(g_bar.device()),
-                &Spectrum::Dense(b_bar),
+                &Spectrum::Dense(b_bar.device()),
             );
             for c in 0..3 {
                 rgb_camera[i][c] = rgb[c];
@@ -158,7 +151,7 @@ impl PixelSensorTrait for DevicePixelSensor {
 
         let mut xyz_output = [[0.; 3]; N_SWATCH_REFLECTANCES];
         let spectra = get_spectra_context();
-        let sensor_white_g = illum.inner_product(&Spectrum::Dense(g_bar.clone()));
+        let sensor_white_g = illum.inner_product(&Spectrum::Dense(g_bar.device()));
         let sensor_white_y = illum.inner_product(&Spectrum::Dense(spectra.y));
         for i in 0..N_SWATCH_REFLECTANCES {
             let s = swatches[i].clone();
@@ -177,13 +170,21 @@ impl PixelSensorTrait for DevicePixelSensor {
         let xyz_from_sensor_rgb = linear_least_squares(rgb_camera, xyz_output)
             .expect("Could not convert sensor illuminance to XYZ space");
 
-        Self {
+        let data = SensorStorage {
-            xyz_from_sensor_rgb,
+            r_bar: r_bar.clone(),
-            r_bar,
+            g_bar: g_bar.clone(),
-            g_bar,
+            b_bar: b_bar.clone(),
-            b_bar,
+        };
+
+        let device = DevicePixelSensor {
+            r_bar: r_bar.device(),
+            g_bar: g_bar.device(),
+            b_bar: b_bar.device(),
             imaging_ratio,
-        }
+            xyz_from_sensor_rgb,
+        };
+
+        Self { device, data }
     }
 
     fn new_with_white_balance(
@@ -192,9 +193,9 @@ impl PixelSensorTrait for DevicePixelSensor {
         imaging_ratio: Float,
     ) -> Self {
         let spectra = get_spectra_context();
-        let r_bar = CIE_X_DATA.device();
+        let r_bar = CIE_X_DATA.clone();
-        let g_bar = CIE_Y_DATA.device();
+        let g_bar = CIE_Y_DATA.clone();
-        let b_bar = CIE_Z_DATA.device();
+        let b_bar = CIE_Z_DATA.clone();
         let xyz_from_sensor_rgb: SquareMatrix<Float, 3>;
 
         if let Some(illum) = sensor_illum {
@@ -205,13 +206,29 @@ impl PixelSensorTrait for DevicePixelSensor {
             xyz_from_sensor_rgb = SquareMatrix::<Float, 3>::default();
         }
 
-        DevicePixelSensor {
+        let data = SensorStorage {
+            r_bar: r_bar.clone(),
+            g_bar: g_bar.clone(),
+            b_bar: b_bar.clone(),
+        };
+
+        let device = DevicePixelSensor {
+            r_bar: r_bar.device(),
+            g_bar: g_bar.device(),
+            b_bar: b_bar.device(),
             xyz_from_sensor_rgb,
-            r_bar,
-            g_bar,
-            b_bar,
             imaging_ratio,
+        };
+
+        Self { data, device }
     }
+
+    pub fn device(&self) -> DevicePixelSensor {
+        self.device
+    }
+
+    fn get_swatches() -> Arc<[Spectrum; N_SWATCH_REFLECTANCES]> {
+        Arc::new(*SWATCH_REFLECTANCES)
     }
 }
 
@@ -347,7 +364,7 @@ pub trait FilmTrait: Sync {
             for (ix, rgb_chunk) in row_data.chunks_exact(3).enumerate() {
                 let p_offset = Point2i::new(ix as i32, iy as i32);
                 let values = ImageChannelValues::from(rgb_chunk);
-                image.set_channel(p_offset, &rgb_desc, &values);
+                image.set_channels(p_offset, &values);
             }
         }
 

src/core/image/mod.rs

@@ -254,6 +254,12 @@ impl Image {
     }
 
     // Read
+    pub fn as_f32_slice(&self) -> Option<&[f32]> {
+        match &self.pixels {
+            PixelStorage::F32(data) => Some(data.as_slice()),
+            _ => None,
+        }
+    }
     pub fn get_channel(&self, p: Point2i, c: i32) -> Float {
         self.get_channel_with_wrap(p, c, WrapMode::Clamp.into())
     }
@@ -336,6 +342,12 @@ impl Image {
         }
     }
 
+    pub fn set_channels(&mut self, p: Point2i, values: &ImageChannelValues) {
+        for i in 0..values.len() {
+            self.set_channel(p, i.try_into().unwrap(), values[i])
+        }
+    }
+
     // Descriptions
     pub fn get_channels_with_desc(
         &self,

src/core/scene/builder.rs

@@ -426,7 +426,7 @@ impl ParserTarget for BasicSceneBuilder {
         })
     }
 
-    fn world_begin(&mut self, loc: FileLoc, arena: &Arena) {
+    fn world_begin(&mut self, loc: FileLoc, arena: Arc<Arena>) {
         self.verify_options("WorldBegin", &loc);
         self.current_block = BlockState::WorldBlock;
         for i in 0..MAX_TRANSFORMS {
@@ -569,7 +569,7 @@ impl ParserTarget for BasicSceneBuilder {
         };
         let entity = TextureSceneEntity {
             base,
-            render_from_object: self.graphics_state.ctm[0].clone(),
+            render_from_object: AnimatedTransform::from_transform(&self.graphics_state.ctm[0]),
         };
 
         if type_name == "float" {

src/core/scene/scene.rs

@@ -109,7 +109,7 @@ impl BasicScene {
         sampler: SceneEntity,
         integ: SceneEntity,
         accel: SceneEntity,
-        arena: &Arena,
+        arena: Arc<Arena>,
     ) {
         *self.integrator.lock() = Some(integ);
         *self.accelerator.lock() = Some(accel);
@@ -131,7 +131,7 @@ impl BasicScene {
                 filt.expect("Must have a filter"),
                 Some(camera.camera_transform.clone()),
                 &film.loc,
-                arena,
+                &arena,
             )
             .expect("Must have a film"),
         );
@@ -141,14 +141,22 @@ impl BasicScene {
             job: None,
         };
 
+        let arena_sampler = Arc::clone(&arena);
         let sampler_film = Arc::clone(&film_instance);
         let sampler_job = run_async(move || {
             let res = sampler_film.as_ref().base().full_resolution;
-            Sampler::create(&sampler.name, &sampler.parameters, res, &sampler.loc, arena)
+            Sampler::create(
+                &sampler.name,
+                &sampler.parameters,
+                res,
+                &sampler.loc,
+                &arena_sampler,
+            )
             .expect("Sampler was not correctly created")
         });
         self.sampler_state.lock().job = Some(sampler_job);
 
+        let arena_camera = Arc::clone(&arena);
         let camera_film = Arc::clone(&film_instance);
         let scene_ptr = Arc::clone(self);
         let camera_job = run_async(move || {
@@ -160,7 +168,7 @@ impl BasicScene {
                 *medium.unwrap(),
                 camera_film,
                 &camera.base.loc,
-                arena,
+                &arena_camera,
             )
             .expect("Failed to create camera")
         });
@@ -736,7 +744,7 @@ impl BasicScene {
         if let Some(job) = state.jobs.remove(name) {
             let job: AsyncJob<Medium> = job;
             let result: Medium = job.wait();
-            let medium: Arc<Medium> = Arc::new(job.wait());
+            let medium: Arc<Medium> = Arc::new(result);
             state.map.insert(name.to_string(), medium.clone());
             return Some(medium);
         }

src/films/gbuffer.rs

@@ -1,8 +1,8 @@
 use super::*;
-use crate::core::film::{CreateFilmBase, PixelSensorTrait};
+use crate::core::film::{CreateFilmBase, PixelSensor};
 use crate::utils::containers::Array2D;
 use anyhow::{Result, anyhow};
-use shared::core::film::{DevicePixelSensor, FilmBase, GBufferFilm};
+use shared::core::film::{FilmBase, GBufferFilm};
 use shared::core::filter::FilterTrait;
 use shared::spectra::RGBColorSpace;
 use shared::utils::AnimatedTransform;
@@ -59,8 +59,8 @@ impl CreateFilm for GBufferFilm {
         let colorspace = params.color_space.as_ref().unwrap();
         let max_component_value = params.get_one_float("maxcomponentvalue", Float::INFINITY);
         let write_fp16 = params.get_one_bool("savefp16", true);
-        let sensor = DevicePixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
+        let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
-        let film_base = FilmBase::create(params, filter, Some(&sensor), loc);
+        let film_base = FilmBase::create(params, filter, Some(&sensor.device()), loc);
 
         let filename = params.get_one_string("filename", "pbrt.exr");
         if Path::new(&filename).extension() != Some("exr".as_ref()) {

src/films/rgb.rs

@@ -1,10 +1,10 @@
 use super::*;
 use crate::Arena;
-use crate::core::film::{CreateFilmBase, PixelSensorTrait};
+use crate::core::film::{CreateFilmBase, PixelSensor};
 use crate::utils::containers::Array2D;
 use anyhow::Result;
 use shared::core::camera::CameraTransform;
-use shared::core::film::{DevicePixelSensor, Film, FilmBase, RGBFilm, RGBPixel};
+use shared::core::film::{Film, FilmBase, RGBFilm, RGBPixel};
 use shared::core::filter::FilterTrait;
 use shared::spectra::RGBColorSpace;
 
@@ -72,8 +72,8 @@ impl CreateFilm for RGBFilm {
         let colorspace = params.color_space.as_ref().unwrap();
         let max_component_value = params.get_one_float("maxcomponentvalue", Float::INFINITY);
         let write_fp16 = params.get_one_bool("savefp16", true);
-        let sensor = DevicePixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
+        let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
-        let film_base = FilmBase::create(params, filter, Some(&sensor), loc);
+        let film_base = FilmBase::create(params, filter, Some(&sensor.device()), loc);
         let film = RGBFilmHost::new(film_base, &colorspace, max_component_value, write_fp16);
         Ok(Film::RGB(film.device))
     }

src/films/spectral.rs

@@ -1,11 +1,11 @@
 use super::*;
-use crate::core::film::{CreateFilmBase, PixelSensorTrait};
+use crate::core::film::{CreateFilmBase, PixelSensor};
 use crate::utils::containers::Array2D;
 use crate::{Arena, FileLoc, ParameterDictionary};
 use anyhow::{Result, anyhow};
 use shared::Float;
 use shared::core::camera::CameraTransform;
-use shared::core::film::{DevicePixelSensor, FilmBase, SpectralFilm, SpectralPixel};
+use shared::core::film::{FilmBase, SpectralFilm, SpectralPixel};
 use shared::core::filter::FilterTrait;
 use shared::spectra::{LAMBDA_MAX, LAMBDA_MIN, RGBColorSpace};
 use shared::utils::AtomicFloat;
@@ -99,8 +99,8 @@ impl CreateFilm for SpectralFilm {
         let colorspace = params.color_space.as_ref().unwrap();
         let max_component_value = params.get_one_float("maxcomponentvalue", Float::INFINITY);
         let write_fp16 = params.get_one_bool("savefp16", true);
-        let sensor = DevicePixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
+        let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, loc)?;
-        let film_base = FilmBase::create(params, filter, Some(&sensor), loc);
+        let film_base = FilmBase::create(params, filter, Some(&sensor.device()), loc);
 
         let filename = params.get_one_string("filename", "pbrt.exr");
         if Path::new(&filename).extension() != Some("exr".as_ref()) {

src/gpu/mod.rs

@@ -1,6 +1,8 @@
-mod context;
+pub mod context;
 
-pub use context::{GpuState, gpu_init, gpu_state, gpu_state_or_panic, gpu_thread_init};
+pub use context::{
+    GPU_STATE, GpuContext, GpuState, gpu_init, gpu_state, gpu_state_or_panic, gpu_thread_init,
+};
 
 pub mod wavefront;
 
@@ -15,5 +17,5 @@ pub enum GpuError {
 }
 
 pub fn gpu_unwrap() -> &'static GpuContext {
-    context::GPU_CONTEXT.get().expect("GPU not initialized")
+    context::GPU_STATE.get().expect("GPU not initialized")
 }

src/lights/projection.rs

@@ -18,13 +18,14 @@ use shared::spectra::RGBColorSpace;
 use shared::utils::math::{radians, square};
 use shared::utils::{Ptr, Transform};
 use std::path::Path;
+use std::sync::Arc;
 
 pub trait CreateProjectionLight {
     fn new(
         render_from_light: Transform,
         medium_interface: MediumInterface,
         scale: Float,
-        image: Ptr<Image>,
+        image: Arc<Image>,
         image_color_space: Ptr<RGBColorSpace>,
         fov: Float,
     ) -> Self;
@@ -35,7 +36,7 @@ impl CreateProjectionLight for ProjectionLight {
         render_from_light: Transform,
         medium_interface: MediumInterface,
         scale: Float,
-        image: Ptr<Image>,
+        image: Arc<Image>,
         image_color_space: Ptr<RGBColorSpace>,
         fov: Float,
     ) -> Self {
@@ -158,7 +159,7 @@ impl CreateLight for ProjectionLight {
             render_from_light_flip,
             medium.into(),
             scale,
-            image.upload(arena),
+            Arc::new(image),
             colorspace.upload(arena),
             fov,
         );

src/shapes/mesh.rs

@@ -22,7 +22,7 @@ pub struct TriQuadMesh {
 }
 
 #[derive(Debug)]
-struct TriangleMeshStorage {
+pub(crate) struct TriangleMeshStorage {
     pub p: Vec<Point3f>,
     pub n: Vec<Normal3f>,
     pub s: Vec<Vector3f>,
@@ -32,7 +32,7 @@ struct TriangleMeshStorage {
 }
 
 #[derive(Debug)]
-struct BilinearMeshStorage {
+pub(crate) struct BilinearMeshStorage {
     pub vertex_indices: Vec<i32>,
     pub p: Vec<Point3f>,
     pub n: Vec<Normal3f>,

src/textures/mod.rs

@@ -6,7 +6,6 @@ mod fbm;
 mod image;
 mod marble;
 mod mix;
-mod ptex;
 mod scaled;
 mod windy;
 mod wrinkled;

src/textures/ptex.rs

@@ -1,74 +0,0 @@
-// use crate::utils::{Arena, FileLoc, TextureParameterDictionary};
-use crate::core::texture::{FloatTextureTrait, SpectrumTextureTrait};
-use crate::spectra::get_colorspace_context;
-use shared::core::spectrum::SpectrumTrait;
-// use ptex::Cache;
-// use ptex_filter::{PtexFilter, PtexFilterOptions, PtexFilterType};
-
-use shared::Float;
-use shared::core::color::{ColorEncoding, RGB};
-use shared::core::texture::{SpectrumType, TextureEvalContext};
-use shared::spectra::{RGBIlluminantSpectrum, SampledSpectrum, SampledWavelengths};
-
-#[derive(Debug, Clone)]
-pub struct PtexTextureBase {
-    pub valid: bool,
-    pub filename: String,
-    pub encoding: ColorEncoding,
-    pub scale: Float,
-}
-
-impl PtexTextureBase {
-    pub fn new(filename: String, encoding: ColorEncoding, scale: Float) -> Self {
-        log::warn!(
-            "Ptex support is currently disabled. Texture '{}' will render as Magenta.",
-            filename
-        );
-
-        Self {
-            filename,
-            encoding,
-            scale,
-            valid: false,
-        }
-    }
-
-    pub fn sample_texture(&self, _ctx: &TextureEvalContext) -> Option<RGB> {
-        Some(RGB::new(1.0, 0.0, 1.0) * self.scale)
-    }
-}
-
-#[derive(Debug, Clone)]
-pub struct FloatPtexTexture {
-    pub base: PtexTextureBase,
-}
-
-impl FloatTextureTrait for FloatPtexTexture {
-    fn evaluate(&self, ctx: &TextureEvalContext) -> Float {
-        if let Some(rgb) = self.base.sample_texture(ctx) {
-            rgb.g
-        } else {
-            0.0
-        }
-    }
-}
-
-#[derive(Clone, Debug)]
-pub struct SpectrumPtexTexture {
-    pub base: PtexTextureBase,
-    pub spectrum_type: SpectrumType,
-}
-
-impl SpectrumTextureTrait for SpectrumPtexTexture {
-    fn evaluate(&self, ctx: &TextureEvalContext, lambda: &SampledWavelengths) -> SampledSpectrum {
-        if let Some(rgb) = self.base.sample_texture(ctx) {
-            let stdcs = get_colorspace_context();
-            let srgb = stdcs.srgb.as_ref();
-            RGBIlluminantSpectrum::new(srgb, rgb).sample(lambda)
-        } else {
-            SampledSpectrum::new(0.0)
-        }
-    }
-}
-
-pub trait CreateGPUPtexTexture {}

src/utils/arena.rs

@@ -78,23 +78,27 @@ impl Arena {
 
     #[cfg(feature = "cuda")]
     unsafe fn alloc_unified(&self, layout: Layout) -> *mut u8 {
-        use cuda_runtime_sys::*;
+        use cust::memory::{UnifiedPointer, cuda_free_unified, cuda_malloc_unified};
 
-        let mut ptr: *mut std::ffi::c_void = std::ptr::null_mut();
         let size = layout.size().max(layout.align());
-
+        if size == 0 {
-        let result = unsafe { cudaMallocManaged(&mut ptr, size, cudaMemAttachGlobal) };
+            return layout.align() as *mut u8;
-
-        if result != cudaError::cudaSuccess {
-            panic!("cudaMallocManaged failed: {:?}", result);
         }
 
-        self.buffer.push((ptr as *mut u8, layout));
+        let mut unified_ptr =
-        ptr as *mut u8
+            unsafe { cuda_malloc_unified::<u8>(size).expect("cuda_malloc_unified failed") };
+        let raw = unified_ptr.as_raw_mut();
+
+        let mut inner = self.inner.lock();
+        inner.buffer.push((raw, layout));
+        raw
     }
 
     #[cfg(not(feature = "cuda"))]
     unsafe fn alloc_unified(&self, layout: Layout) -> *mut u8 {
+        if layout.size() == 0 {
+            return layout.align() as *mut u8;
+        }
         let ptr = unsafe { std::alloc::alloc(layout) };
         let mut inner = self.inner.lock();
         inner.buffer.push((ptr, layout));
@@ -129,20 +133,23 @@ impl Arena {
         // 5. Return handle
         0
     }
-
-    // pub fn raw_data(&self) -> &[u8] {
-    //     &self.buffer
-    // }
 }
 
 impl Drop for Arena {
     fn drop(&mut self) {
         let inner = self.inner.get_mut();
         for (ptr, layout) in inner.buffer.drain(..) {
+            if layout.size() == 0 {
+                continue;
+            }
+
             unsafe {
                 #[cfg(feature = "cuda")]
                 {
-                    cuda_runtime_sys::cudaFree(ptr as *mut _);
+                    use cust::memory::{UnifiedPointer, cuda_free_unified};
+                    if layout.size() > 0 {
+                        let _ = cuda_free_unified(UnifiedPointer::wrap(ptr as *mut u8));
+                    }
                 }
                 #[cfg(not(feature = "cuda"))]
                 {

src/utils/mipmap.rs

@@ -5,10 +5,12 @@ use shared::core::geometry::{Point2f, Point2i, Vector2f, VectorLike};
 use shared::core::image::{WrapMode, WrapMode2D};
 use shared::spectra::RGBColorSpace;
 use shared::utils::math::{lerp, safe_sqrt, square};
-use std::path::Path;
-
 use std::hash::{Hash, Hasher};
 use std::ops::{Add, Mul, Sub};
+use std::path::Path;
+
+#[cfg(feature = "cuda")]
+use std::sync::OnceLock;
 
 #[repr(C)]
 #[derive(Debug, Clone, Copy, PartialEq, Eq, Hash)]
@@ -59,7 +61,6 @@ impl Eq for MIPMapFilterOptions {}
 impl Hash for MIPMapFilterOptions {
     fn hash<H: Hasher>(&self, state: &mut H) {
         self.filter.hash(state);
-        // Hash the bits, not the float value
         self.max_anisotropy.to_bits().hash(state);
     }
 }
@@ -76,11 +77,9 @@ impl MIPMapSample for Float {
     fn zero() -> Self {
         0.
     }
-
     fn sample_bilerp(image: &Image, st: Point2f, wrap: WrapMode2D) -> Self {
         image.bilerp_channel_with_wrap(st, 0, wrap)
     }
-
     fn sample_texel(image: &Image, st: Point2i, wrap: WrapMode2D) -> Self {
         image.get_channel_with_wrap(st, 0, wrap)
     }
@@ -90,7 +89,6 @@ impl MIPMapSample for RGB {
     fn zero() -> Self {
         RGB::new(0., 0., 0.)
     }
-
     fn sample_bilerp(image: &Image, st: Point2f, wrap: WrapMode2D) -> Self {
         let nc = image.n_channels();
         if nc >= 3 {
@@ -103,7 +101,6 @@ impl MIPMapSample for RGB {
             RGB::new(v, v, v)
         }
     }
-
     fn sample_texel(image: &Image, st: Point2i, wrap: WrapMode2D) -> Self {
         let nc = image.n_channels();
         if nc >= 3 {
@@ -124,6 +121,8 @@ pub struct MIPMap {
     pub color_space: Option<RGBColorSpace>,
     pub wrap_mode: WrapMode,
     pub options: MIPMapFilterOptions,
+    #[cfg(feature = "cuda")]
+    tex_obj: OnceLock<u64>,
 }
 
 impl MIPMap {
@@ -139,6 +138,8 @@ impl MIPMap {
             color_space,
             wrap_mode,
             options,
+            #[cfg(feature = "cuda")]
+            tex_obj: OnceLock::new(),
         }
     }
 
@@ -174,7 +175,6 @@ impl MIPMap {
         mut dst1: Vector2f,
     ) -> T {
         if self.options.filter != FilterFunction::Ewa {
-            // Compute largest change in texture coordinates
             let width = 2.0
                 * [
                     dst0.x().abs(),
@@ -186,8 +186,6 @@ impl MIPMap {
                 .reduce(Float::max)
                 .unwrap_or(0.0);
 
-            // Compute MIP Map level
-            // n_levels - 1 + log2(width) maps width=1.0 to the top level (1x1)
             let n_levels = self.levels() as Float;
             let level = n_levels - 1.0 + width.max(1e-8).log2();
 
@@ -196,7 +194,6 @@ impl MIPMap {
             }
 
             let i_level = level.floor() as usize;
-
             return match self.options.filter {
                 FilterFunction::Point => {
                     let resolution = self.level_resolution(i_level);
@@ -208,7 +205,6 @@ impl MIPMap {
                 }
                 FilterFunction::Bilinear => self.bilerp(i_level, st),
                 FilterFunction::Trilinear => {
-                    // Interpolate between current level and next level
                     let v0 = self.bilerp(i_level, st);
                     let v1 = self.bilerp(i_level + 1, st);
                     let t = level - i_level as Float;
@@ -221,12 +217,9 @@ impl MIPMap {
         if dst0.norm_squared() < dst1.norm_squared() {
             std::mem::swap(&mut dst0, &mut dst1);
         }
-
         let longer_len = dst0.norm();
         let mut shorter_len = dst1.norm();
 
-        // If ellipse is too thin, fatten the minor axis to limit the number
-        // of texels
         if shorter_len * self.options.max_anisotropy < longer_len && shorter_len > 0.0 {
             let scale = longer_len / (shorter_len * self.options.max_anisotropy);
             dst1 *= scale;
@@ -242,7 +235,6 @@ impl MIPMap {
 
         let v0 = self.ewa(ilod, st, dst0, dst1);
         let v1 = self.ewa(ilod + 1, st, dst0, dst1);
-
         lerp(lod - ilod as Float, v0, v1)
     }
 
@@ -250,12 +242,10 @@ impl MIPMap {
         if level >= self.levels() {
             panic!("MIPMap level out of bounds");
         }
-
         let image = &self.pyramid[level];
         let wrap_2d = WrapMode2D {
             uv: [self.wrap_mode; 2],
         };
-
         T::sample_texel(image, st, wrap_2d)
     }
 
@@ -277,6 +267,7 @@ impl MIPMap {
         if level > self.levels() {
             return self.texel(self.levels() - 1, Point2i::new(0, 0));
         }
+
         let level_res = self.level_resolution(level);
         st[0] = st[0] * level_res[0] as Float - 0.5;
         st[1] = st[1] * level_res[1] as Float - 0.5;
@@ -284,6 +275,7 @@ impl MIPMap {
         dst0[1] *= level_res[1] as Float;
         dst1[0] *= level_res[0] as Float;
         dst1[1] *= level_res[1] as Float;
+
         let mut a = square(dst0[1]) + square(dst1[1]) + 1.;
         let mut b = -2. * (dst0[0] + dst0[1] + dst1[1]);
         let mut c = square(dst0[0]) + square(dst1[0]) + 1.;
@@ -296,35 +288,30 @@ impl MIPMap {
         let inv_det = 1. / det;
         let u_sqrt = safe_sqrt(det * c);
         let v_sqrt = safe_sqrt(det * a);
+
         let s0: i32 = (st[0] - 2. * inv_det * u_sqrt).ceil() as i32;
         let s1: i32 = (st[0] + 2. * inv_det * u_sqrt).floor() as i32;
         let t0: i32 = (st[1] - 2. * inv_det * v_sqrt).ceil() as i32;
         let t1: i32 = (st[1] + 2. * inv_det * v_sqrt).floor() as i32;
+
         let mut sum = T::zero();
         let mut sum_wts = 0.;
+
         for it in t0..=t1 {
             let tt = it as Float - st[1];
             for is in s0..=s1 {
                 let ss = is as Float - st[0];
-
-                // Compute squared radius and filter texel if inside ellipse
                 let r2 = a * square(ss) + b * ss * tt + c * square(tt);
-
                 if r2 < 1.0 {
-                    // Map r2 to LUT index
                     let index = (r2 * MIP_FILTER_LUT_SIZE as Float)
                         .min((MIP_FILTER_LUT_SIZE - 1) as Float)
                         as usize;
-
                     let weight = MIP_FILTER_LUT[index];
-
-                    // Accumulate
                     sum = sum + self.texel::<T>(level, Point2i::new(is, it)) * weight;
                     sum_wts += weight;
                 }
             }
         }
-
         sum * (1. / sum_wts)
     }
 
@@ -336,9 +323,6 @@ impl MIPMap {
     ) -> Result<MIPMap, ()> {
         let image_and_metadata = Image::read(filename, Some(encoding)).unwrap();
         let image = image_and_metadata.image;
-        // if image.n_channels() != 1 {
-        //     let rgba_dsc = image.all_channels_desc();
-        // }
         Ok(MIPMap::new(
             image,
             image_and_metadata.metadata.colorspace,
@@ -369,7 +353,7 @@ fn create_cuda_texture(pyramid: &[Image], wrap_mode: WrapMode) -> u64 {
         base.resolution().x() as usize,
         base.resolution().y() as usize,
     );
-    let channels = base.n_channels();
+    let channels = base.n_channels() as usize;
 
     unsafe {
         let channel_desc = cudaCreateChannelDesc(
@@ -377,57 +361,60 @@ fn create_cuda_texture(pyramid: &[Image], wrap_mode: WrapMode) -> u64 {
             if channels > 1 { 32 } else { 0 },
             if channels > 2 { 32 } else { 0 },
             if channels > 3 { 32 } else { 0 },
-            cudaChannelFormatKindFloat,
+            cudaChannelFormatKind::cudaChannelFormatKindFloat,
         );
 
         let mut array: cudaArray_t = std::ptr::null_mut();
         cudaMallocArray(&mut array, &channel_desc, width, height, 0);
 
-        let pixels = base.as_slice();
+        // Get raw pixel data from the base image
+        let pixel_data = base
+            .as_f32_slice()
+            .expect("GPU upload requires Float encoded image");
+        let row_bytes = width * channels * std::mem::size_of::<f32>();
+
         cudaMemcpy2DToArray(
             array,
             0,
             0,
-            pixels.as_ptr() as *const _,
+            pixel_data.as_ptr() as *const _,
-            width * channels * std::mem::size_of::<f32>(),
+            row_bytes,
-            width * channels * std::mem::size_of::<f32>(),
+            row_bytes,
             height,
-            cudaMemcpyHostToDevice,
+            cudaMemcpyKind::cudaMemcpyHostToDevice,
         );
 
-        // 4. Create texture object
         let res_desc = cudaResourceDesc {
-            resType: cudaResourceTypeArray,
+            resType: cudaResourceType::cudaResourceTypeArray,
             res: cudaResourceDesc__bindgen_ty_1 {
                 array: cudaResourceDesc__bindgen_ty_1__bindgen_ty_1 { array },
             },
         };
 
         let address_mode = match wrap_mode {
-            WrapMode::Repeat => cudaAddressModeWrap,
+            WrapMode::Repeat => cudaTextureAddressMode::cudaAddressModeWrap,
-            WrapMode::Clamp => cudaAddressModeClamp,
+            WrapMode::Clamp => cudaTextureAddressMode::cudaAddressModeClamp,
-            WrapMode::Black => cudaAddressModeBorder,
+            WrapMode::Black => cudaTextureAddressMode::cudaAddressModeBorder,
-            WrapMode::OctahedralSphere => cudaAddressModeBorder,
+            WrapMode::OctahedralSphere => cudaTextureAddressMode::cudaAddressModeBorder,
         };
 
         let tex_desc = cudaTextureDesc {
             addressMode: [address_mode; 3],
-            filterMode: cudaFilterModeLinear,
+            filterMode: cudaTextureFilterMode::cudaFilterModeLinear,
-            readMode: cudaReadModeElementType,
+            readMode: cudaTextureReadMode::cudaReadModeElementType,
             normalizedCoords: 1,
             ..std::mem::zeroed()
         };
 
         let mut tex_obj: cudaTextureObject_t = 0;
         cudaCreateTextureObject(&mut tex_obj, &res_desc, &tex_desc, std::ptr::null());
-
         tex_obj
     }
 }
 
 static MIP_FILTER_LUT_SIZE: usize = 128;
+
 static MIP_FILTER_LUT: [Float; MIP_FILTER_LUT_SIZE] = [
-    // MIPMap EWA Lookup Table Values
     0.864664733,
     0.849040031,
     0.83365953,

src/utils/parser.rs

@@ -50,7 +50,7 @@ pub trait ParserTarget {
     fn medium_interface(&mut self, inside_name: &str, outside_name: &str, loc: FileLoc);
     fn sampler(&mut self, name: &str, params: &ParsedParameterVector, loc: FileLoc);
 
-    fn world_begin(&mut self, loc: FileLoc, arena: &Arena);
+    fn world_begin(&mut self, loc: FileLoc, arena: Arc<Arena>);
     fn attribute_begin(&mut self, loc: FileLoc);
     fn attribute_end(&mut self, loc: FileLoc);
     fn attribute(&mut self, target: &str, params: ParsedParameterVector, loc: FileLoc);
@@ -459,7 +459,7 @@ impl ParserTarget for FormattingParserTarget {
         println!("{}CoordSysTransform \"{}\"", self.indent(0), name);
     }
 
-    fn world_begin(&mut self, _loc: FileLoc, _arena: &Arena) {
+    fn world_begin(&mut self, _loc: FileLoc, _arena: Arc<Arena>) {
         println!("{}WorldBegin", self.indent(0));
         self.cat_indent_count += 4;
     }
@@ -1030,7 +1030,7 @@ impl<'a> SceneParser<'a> {
                 },
 
                 'W' => match token.text.as_str() {
-                    "WorldBegin" => self.target.world_begin(token.loc, &arena),
+                    "WorldBegin" => self.target.world_begin(token.loc, arena.clone()),
                     "WorldEnd" => {}
                     _ => {
                         return Err(ParserError::Generic(