From 0b04d54346ed4a2ae0682569cd821cc1a2983f2e Mon Sep 17 00:00:00 2001 From: Wito Wiala Date: Wed, 18 Feb 2026 19:53:13 +0000 Subject: [PATCH] Cleanup --- Cargo.toml | 3 +- shared/src/core/film.rs | 4 ++ shared/src/textures/scaled.rs | 8 +-- src/core/camera.rs | 4 +- src/core/film.rs | 111 ++++++++++++++++++++-------------- src/core/image/mod.rs | 12 ++++ src/core/scene/builder.rs | 4 +- src/core/scene/scene.rs | 20 ++++-- src/films/gbuffer.rs | 8 +-- src/films/rgb.rs | 8 +-- src/films/spectral.rs | 8 +-- src/gpu/mod.rs | 8 ++- src/lights/projection.rs | 7 ++- src/shapes/mesh.rs | 4 +- src/textures/mod.rs | 1 - src/textures/ptex.rs | 74 ----------------------- src/utils/arena.rs | 35 ++++++----- src/utils/mipmap.rs | 79 ++++++++++-------------- src/utils/parser.rs | 6 +- 19 files changed, 184 insertions(+), 220 deletions(-) delete mode 100644 src/textures/ptex.rs diff --git a/Cargo.toml b/Cargo.toml index c4a8733..d4fd833 100644 --- a/Cargo.toml +++ b/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" diff --git a/shared/src/core/film.rs b/shared/src/core/film.rs index a7bc018..5d9c3c2 100644 --- a/shared/src/core/film.rs +++ b/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) -> RGB { + todo!() + } } #[repr(C)] diff --git a/shared/src/textures/scaled.rs b/shared/src/textures/scaled.rs index 0d79b6f..de2c81a 100644 --- a/shared/src/textures/scaled.rs +++ b/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 } diff --git a/src/core/camera.rs b/src/core/camera.rs index 60f8054..b77e1ee 100644 --- a/src/core/camera.rs +++ b/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); diff --git a/src/core/film.rs b/src/core/film.rs index be83d9f..2017816 100644 --- a/src/core/film.rs +++ b/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); - Spectrum::Piecewise(pls) + let pls = PiecewiseLinearSpectrumBuffer::from_interleaved(raw_data, false); + Spectrum::Piecewise(pls.device) }) }); -pub trait PixelSensorTrait: Sized { - fn get_swatches() -> Arc<[Spectrum; N_SWATCH_REFLECTANCES]> { - Arc::new(*SWATCH_REFLECTANCES) - } +#[derive(Debug, Clone)] +struct SensorStorage { + 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, exposure_time: Float, @@ -103,23 +113,6 @@ pub trait PixelSensorTrait: Sized { } } - fn new( - r: &Spectrum, - g: &Spectrum, - b: &Spectrum, - output_colorspace: Arc, - 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 g_bar = DenselySampledSpectrumBuffer::from_spectrum(g).device(); - let b_bar = DenselySampledSpectrumBuffer::from_spectrum(b).device(); + let r_bar = DenselySampledSpectrumBuffer::from_spectrum(r); + let g_bar = DenselySampledSpectrumBuffer::from_spectrum(g); + 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::( &swatches[i], illum, - &Spectrum::Dense(r_bar), - &Spectrum::Dense(g_bar), - &Spectrum::Dense(b_bar), + &Spectrum::Dense(r_bar.device()), + &Spectrum::Dense(g_bar.device()), + &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 { - xyz_from_sensor_rgb, - r_bar, - g_bar, - b_bar, + 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(), 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 g_bar = CIE_Y_DATA.device(); - let b_bar = CIE_Z_DATA.device(); + let r_bar = CIE_X_DATA.clone(); + let g_bar = CIE_Y_DATA.clone(); + let b_bar = CIE_Z_DATA.clone(); let xyz_from_sensor_rgb: SquareMatrix; if let Some(illum) = sensor_illum { @@ -205,13 +206,29 @@ impl PixelSensorTrait for DevicePixelSensor { xyz_from_sensor_rgb = SquareMatrix::::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); } } diff --git a/src/core/image/mod.rs b/src/core/image/mod.rs index 68ee82e..416f676 100644 --- a/src/core/image/mod.rs +++ b/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, diff --git a/src/core/scene/builder.rs b/src/core/scene/builder.rs index f9c1ea2..8fede2a 100644 --- a/src/core/scene/builder.rs +++ b/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) { 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" { diff --git a/src/core/scene/scene.rs b/src/core/scene/scene.rs index bd35197..2aaa559 100644 --- a/src/core/scene/scene.rs +++ b/src/core/scene/scene.rs @@ -109,7 +109,7 @@ impl BasicScene { sampler: SceneEntity, integ: SceneEntity, accel: SceneEntity, - arena: &Arena, + arena: Arc, ) { *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) - .expect("Sampler was not correctly created") + 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 = job; let result: Medium = job.wait(); - let medium: Arc = Arc::new(job.wait()); + let medium: Arc = Arc::new(result); state.map.insert(name.to_string(), medium.clone()); return Some(medium); } diff --git a/src/films/gbuffer.rs b/src/films/gbuffer.rs index fa6f831..00ff578 100644 --- a/src/films/gbuffer.rs +++ b/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 film_base = FilmBase::create(params, filter, Some(&sensor), loc); + let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, 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()) { diff --git a/src/films/rgb.rs b/src/films/rgb.rs index ed321d2..2a55da6 100644 --- a/src/films/rgb.rs +++ b/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 film_base = FilmBase::create(params, filter, Some(&sensor), loc); + let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, 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)) } diff --git a/src/films/spectral.rs b/src/films/spectral.rs index 633616c..749ec47 100644 --- a/src/films/spectral.rs +++ b/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 film_base = FilmBase::create(params, filter, Some(&sensor), loc); + let sensor = PixelSensor::create(params, colorspace.clone(), exposure_time, 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()) { diff --git a/src/gpu/mod.rs b/src/gpu/mod.rs index 81d7230..88c4bf2 100644 --- a/src/gpu/mod.rs +++ b/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") } diff --git a/src/lights/projection.rs b/src/lights/projection.rs index b8c2b86..f195a4a 100644 --- a/src/lights/projection.rs +++ b/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: Arc, image_color_space: Ptr, fov: Float, ) -> Self; @@ -35,7 +36,7 @@ impl CreateProjectionLight for ProjectionLight { render_from_light: Transform, medium_interface: MediumInterface, scale: Float, - image: Ptr, + image: Arc, image_color_space: Ptr, 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, ); diff --git a/src/shapes/mesh.rs b/src/shapes/mesh.rs index b085770..2ec5dfe 100644 --- a/src/shapes/mesh.rs +++ b/src/shapes/mesh.rs @@ -22,7 +22,7 @@ pub struct TriQuadMesh { } #[derive(Debug)] -struct TriangleMeshStorage { +pub(crate) struct TriangleMeshStorage { pub p: Vec, pub n: Vec, pub s: Vec, @@ -32,7 +32,7 @@ struct TriangleMeshStorage { } #[derive(Debug)] -struct BilinearMeshStorage { +pub(crate) struct BilinearMeshStorage { pub vertex_indices: Vec, pub p: Vec, pub n: Vec, diff --git a/src/textures/mod.rs b/src/textures/mod.rs index d355a5b..6dad3ab 100644 --- a/src/textures/mod.rs +++ b/src/textures/mod.rs @@ -6,7 +6,6 @@ mod fbm; mod image; mod marble; mod mix; -mod ptex; mod scaled; mod windy; mod wrinkled; diff --git a/src/textures/ptex.rs b/src/textures/ptex.rs deleted file mode 100644 index a89bf6f..0000000 --- a/src/textures/ptex.rs +++ /dev/null @@ -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 { - 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 {} diff --git a/src/utils/arena.rs b/src/utils/arena.rs index 4da0fb5..42c2e84 100644 --- a/src/utils/arena.rs +++ b/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()); - - let result = unsafe { cudaMallocManaged(&mut ptr, size, cudaMemAttachGlobal) }; - - if result != cudaError::cudaSuccess { - panic!("cudaMallocManaged failed: {:?}", result); + if size == 0 { + return layout.align() as *mut u8; } - self.buffer.push((ptr as *mut u8, layout)); - ptr as *mut u8 + let mut unified_ptr = + unsafe { cuda_malloc_unified::(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"))] { diff --git a/src/utils/mipmap.rs b/src/utils/mipmap.rs index 51a08bd..db56acb 100644 --- a/src/utils/mipmap.rs +++ b/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(&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, pub wrap_mode: WrapMode, pub options: MIPMapFilterOptions, + #[cfg(feature = "cuda")] + tex_obj: OnceLock, } 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::(level, Point2i::new(is, it)) * weight; sum_wts += weight; } } } - sum * (1. / sum_wts) } @@ -336,9 +323,6 @@ impl MIPMap { ) -> Result { 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::(); + cudaMemcpy2DToArray( array, 0, 0, - pixels.as_ptr() as *const _, - width * channels * std::mem::size_of::(), - width * channels * std::mem::size_of::(), + pixel_data.as_ptr() as *const _, + row_bytes, + 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::Clamp => cudaAddressModeClamp, - WrapMode::Black => cudaAddressModeBorder, - WrapMode::OctahedralSphere => cudaAddressModeBorder, + WrapMode::Repeat => cudaTextureAddressMode::cudaAddressModeWrap, + WrapMode::Clamp => cudaTextureAddressMode::cudaAddressModeClamp, + WrapMode::Black => cudaTextureAddressMode::cudaAddressModeBorder, + WrapMode::OctahedralSphere => cudaTextureAddressMode::cudaAddressModeBorder, }; let tex_desc = cudaTextureDesc { addressMode: [address_mode; 3], - filterMode: cudaFilterModeLinear, - readMode: cudaReadModeElementType, + filterMode: cudaTextureFilterMode::cudaFilterModeLinear, + 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, diff --git a/src/utils/parser.rs b/src/utils/parser.rs index 90d937b..ec6b352 100644 --- a/src/utils/parser.rs +++ b/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); 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) { 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(