cleaned up collision in shader, abstracted out raycast, and added gpu timing

src/client/mod.rs

@@ -19,7 +19,7 @@ use system::render::add_grid;
 
 use crate::{
     common::{ClientMessage, ServerHandle, ServerMessage},
-    server::Server,
+    server::Server, util::timer::Timer,
 };
 
 use self::{input::Input, render::Renderer, ClientState};
@@ -37,6 +37,7 @@ pub struct Client<'a> {
     window: Arc<Window>,
     state: ClientState,
     renderer: Renderer<'a>,
+    render_timer: Timer,
     render_commands: Vec<RenderCommand>,
     exit: bool,
     input: Input,
@@ -78,6 +79,7 @@ impl Client<'_> {
             window,
             exit: false,
             renderer,
+            render_timer: Timer::new(60),
             render_commands: Vec::new(),
             state,
             input: Input::new(),
@@ -132,17 +134,18 @@ impl Client<'_> {
             commands.extend(world_cmds);
             self.renderer.handle_commands(commands);
             self.renderer.draw();
+            self.render_timer.add(self.renderer.timer().duration(0));
         }
 
         if now >= self.second_target {
             self.second_target += Duration::from_secs(1);
-            // let timer = self.renderer.timer();
+            let timer = &self.render_timer;
-            // println!(
+            println!(
-            //     "avg: {:4?}; max: {:4?}; fps: {:4?}",
+                "avg: {:4?}; max: {:4?}; fps: {:4?}",
-            //     timer.avg(),
+                timer.avg(),
-            //     timer.max(),
+                timer.max(),
-            //     timer.per_sec(),
+                timer.per_sec(),
-            // );
+            );
         }
 
         if self.exit {

src/client/render/command.rs

@@ -7,8 +7,6 @@ use super::{voxel::VoxelColor, Renderer};
 use bevy_ecs::entity::Entity;
 use nalgebra::{Rotation3, Vector3};
 use ndarray::Array3;
-use std::sync::Arc;
-use winit::window::Window;
 
 #[derive(Debug, Clone)]
 pub enum RenderCommand {
@@ -76,7 +74,6 @@ impl<'a> Renderer<'a> {
                 &self.device,
                 &mut self.encoder,
                 &mut self.staging_belt,
-                self.size,
                 &self.camera,
             );
         }

src/client/render/mod.rs

@@ -8,7 +8,7 @@ pub use command::*;
 use super::camera::Camera;
 use crate::client::rsc::CLEAR_COLOR;
 use nalgebra::Vector2;
-use util::DepthTexture;
+use util::GPUTimer;
 use voxel::VoxelPipeline;
 use winit::{dpi::PhysicalSize, window::Window};
 
@@ -21,14 +21,12 @@ pub struct Renderer<'a> {
     config: wgpu::SurfaceConfiguration,
     staging_belt: wgpu::util::StagingBelt,
     voxel_pipeline: VoxelPipeline,
+    timer: GPUTimer,
     camera: Camera,
-    depth_texture: DepthTexture,
 }
 
 impl<'a> Renderer<'a> {
-    pub fn new(
+    pub fn new(window: Arc<Window>) -> Self {
-        window: Arc<Window>,
-    ) -> Self {
         let size = window.inner_size();
 
         let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
@@ -47,11 +45,13 @@ impl<'a> Renderer<'a> {
         }))
         .expect("Could not get adapter!");
 
-        let buf_size = (10u32.pow(9) * 15) / 10;
+        let buf_size = (10f32.powi(9) * 1.5) as u32;
         let (device, queue) = pollster::block_on(adapter.request_device(
             &wgpu::DeviceDescriptor {
                 label: None,
-                required_features: wgpu::Features::empty(),
+                required_features: wgpu::Features::TIMESTAMP_QUERY
+                    | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS
+                    | wgpu::Features::TIMESTAMP_QUERY_INSIDE_PASSES,
                 required_limits: wgpu::Limits {
                     max_storage_buffer_binding_size: buf_size,
                     max_buffer_size: buf_size as u64,
@@ -94,8 +94,6 @@ impl<'a> Renderer<'a> {
         // doesn't affect performance much and depends on "normal" zoom
         let staging_belt = wgpu::util::StagingBelt::new(4096 * 4);
 
-        let depth_texture = DepthTexture::init(&device, &config, "depth_texture");
-
         Self {
             camera: Camera::default(),
             size: Vector2::new(size.width, size.height),
@@ -103,10 +101,10 @@ impl<'a> Renderer<'a> {
             staging_belt,
             surface,
             encoder: Self::create_encoder(&device),
+            timer: GPUTimer::new(&device, 1),
             device,
             config,
             queue,
-            depth_texture,
         }
     }
 
@@ -135,7 +133,9 @@ impl<'a> Renderer<'a> {
             label: None,
             timestamp_writes: None,
         });
+        self.timer.start_compute(&mut compute_pass, 0);
         self.voxel_pipeline.compute(&mut compute_pass);
+        self.timer.stop_compute(&mut compute_pass, 0);
         drop(compute_pass);
 
         let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
@@ -148,14 +148,6 @@ impl<'a> Renderer<'a> {
                     store: wgpu::StoreOp::Store,
                 },
             })],
-            // depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment {
-            //     view: &self.depth_texture.view,
-            //     depth_ops: Some(wgpu::Operations {
-            //         load: wgpu::LoadOp::Clear(1.0),
-            //         store: wgpu::StoreOp::Store,
-            //     }),
-            //     stencil_ops: None,
-            // }),
             depth_stencil_attachment: None,
             timestamp_writes: None,
             occlusion_query_set: None,
@@ -163,10 +155,14 @@ impl<'a> Renderer<'a> {
         self.voxel_pipeline.draw(&mut render_pass);
         drop(render_pass);
 
+        self.timer.resolve(&mut encoder);
+
         self.staging_belt.finish();
         self.queue.submit(std::iter::once(encoder.finish()));
         output.present();
         self.staging_belt.recall();
+
+        self.timer.finish(&self.device);
     }
 
     pub fn resize(&mut self, size: PhysicalSize<u32>) {
@@ -175,14 +171,15 @@ impl<'a> Renderer<'a> {
         self.config.height = size.height;
         self.surface.configure(&self.device, &self.config);
         self.voxel_pipeline.resize(&self.device, self.size);
-
-        self.depth_texture = DepthTexture::init(&self.device, &self.config, "depth_texture");
         self.voxel_pipeline.update_view(
             &self.device,
             &mut self.encoder,
             &mut self.staging_belt,
-            self.size,
             &self.camera,
         );
     }
+
+    pub fn timer(&self) -> &GPUTimer {
+        &self.timer
+    }
 }

src/client/render/util/array.rs

@@ -1,7 +1,7 @@
 use std::marker::PhantomData;
 use wgpu::{util::DeviceExt, BufferAddress, BufferUsages};
 
-pub struct ArrBuf<T: bytemuck::Pod> {
+pub struct ArrayBuffer<T: bytemuck::Pod> {
     len: usize,
     buffer: wgpu::Buffer,
     label: String,
@@ -10,7 +10,7 @@ pub struct ArrBuf<T: bytemuck::Pod> {
     moves: Vec<BufMove>,
 }
 
-impl<T: bytemuck::Pod> ArrBuf<T> {
+impl<T: bytemuck::Pod> ArrayBuffer<T> {
     pub fn update(
         &mut self,
         device: &wgpu::Device,
@@ -62,6 +62,42 @@ impl<T: bytemuck::Pod> ArrBuf<T> {
         resized
     }
 
+    pub fn add(
+        &mut self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        belt: &mut wgpu::util::StagingBelt,
+        data: &[T],
+    ) {
+        self.update(
+            device,
+            encoder,
+            belt,
+            self.len + data.len(),
+            &[ArrBufUpdate {
+                offset: self.len,
+                data,
+            }],
+        );
+    }
+
+    pub fn set(
+        &mut self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        belt: &mut wgpu::util::StagingBelt,
+        offset: usize,
+        data: &[T],
+    ) {
+        self.update(
+            device,
+            encoder,
+            belt,
+            self.len,
+            &[ArrBufUpdate { offset, data }],
+        );
+    }
+
     pub fn init(device: &wgpu::Device, label: &str, usage: BufferUsages) -> Self {
         let label = &(label.to_owned() + " Buffer");
         Self {
@@ -116,10 +152,6 @@ impl<T: bytemuck::Pod> ArrBuf<T> {
         })
     }
 
-    pub fn buffer(&self) -> &wgpu::Buffer {
-        &self.buffer
-    }
-
     pub fn mov(&mut self, mov: BufMove) {
         self.moves.push(mov);
     }
@@ -127,6 +159,31 @@ impl<T: bytemuck::Pod> ArrBuf<T> {
     pub fn len(&self) -> usize {
         self.len
     }
+
+    pub fn bind_group_layout_entry(
+        &self,
+        binding: u32,
+        visibility: wgpu::ShaderStages,
+        ty: wgpu::BufferBindingType,
+    ) -> wgpu::BindGroupLayoutEntry {
+        wgpu::BindGroupLayoutEntry {
+            binding,
+            visibility,
+            ty: wgpu::BindingType::Buffer {
+                ty,
+                has_dynamic_offset: false,
+                min_binding_size: None,
+            },
+            count: None,
+        }
+    }
+
+    pub fn bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry {
+        return wgpu::BindGroupEntry {
+            binding,
+            resource: self.buffer.as_entire_binding(),
+        };
+    }
 }
 
 pub struct ArrBufUpdate<'a, T> {

src/client/render/util/instance.rs

@@ -1,78 +0,0 @@
-use wgpu::{BufferUsages, VertexAttribute};
-
-use super::buf::{ArrBuf, ArrBufUpdate, BufMove};
-
-pub struct Instances<T: bytemuck::Pod> {
-    buf: ArrBuf<T>,
-    location: u32,
-    attrs: Vec<VertexAttribute>,
-}
-
-impl<T: bytemuck::Pod> Instances<T> {
-    pub fn update(
-        &mut self,
-        device: &wgpu::Device,
-        encoder: &mut wgpu::CommandEncoder,
-        belt: &mut wgpu::util::StagingBelt,
-        size: usize,
-        updates: &[ArrBufUpdate<T>],
-    ) -> bool {
-        self.buf.update(device, encoder, belt, size, updates)
-    }
-
-    pub fn init(
-        device: &wgpu::Device,
-        label: &str,
-        location: u32,
-        attrs: &[wgpu::VertexAttribute],
-    ) -> Self {
-        Self {
-            buf: ArrBuf::init(
-                device,
-                &(label.to_owned() + " Instance"),
-                BufferUsages::VERTEX,
-            ),
-            location,
-            attrs: attrs.to_vec(),
-        }
-    }
-
-    pub fn init_with(
-        device: &wgpu::Device,
-        label: &str,
-        location: u32,
-        attrs: &[wgpu::VertexAttribute],
-        data: &[T],
-    ) -> Self {
-        Self {
-            buf: ArrBuf::init_with(
-                device,
-                &(label.to_owned() + " Instance"),
-                BufferUsages::VERTEX,
-                data,
-            ),
-            location,
-            attrs: attrs.to_vec(),
-        }
-    }
-
-    pub fn set_in<'a>(&'a self, render_pass: &mut wgpu::RenderPass<'a>) {
-        render_pass.set_vertex_buffer(self.location, self.buf.buffer().slice(..));
-    }
-
-    pub fn desc(&self) -> wgpu::VertexBufferLayout {
-        wgpu::VertexBufferLayout {
-            array_stride: std::mem::size_of::<T>() as wgpu::BufferAddress,
-            step_mode: wgpu::VertexStepMode::Instance,
-            attributes: &self.attrs,
-        }
-    }
-
-    pub fn mov(&mut self, mov: BufMove) {
-        self.buf.mov(mov);
-    }
-
-    pub fn len(&self) -> usize {
-        self.buf.len()
-    }
-}

src/client/render/util/mod.rs

@@ -1,13 +1,9 @@
-mod buf;
+mod array;
-mod instance;
-mod storage;
 mod uniform;
-mod vertex;
 mod texture;
+mod timer;
 
-pub use buf::*;
+pub use array::*;
-pub use instance::*;
-pub use storage::*;
 pub use uniform::*;
-pub use vertex::*;
 pub use texture::*;
+pub use timer::*;

src/client/render/util/storage.rs

@@ -1,73 +0,0 @@
-use super::buf::{ArrBuf, ArrBufUpdate, BufMove};
-use wgpu::BufferUsages;
-
-pub struct Storage<T: bytemuck::Pod> {
-    binding: u32,
-    buf: ArrBuf<T>,
-    visibility: wgpu::ShaderStages,
-}
-
-impl<T: bytemuck::Pod> Storage<T> {
-    pub fn init(device: &wgpu::Device, visibility: wgpu::ShaderStages, label: &str, binding: u32) -> Self {
-        Self {
-            buf: ArrBuf::init(
-                device,
-                &(label.to_owned() + " Storage"),
-                BufferUsages::STORAGE,
-            ),
-            visibility,
-            binding,
-        }
-    }
-    pub fn init_with(device: &wgpu::Device, visibility: wgpu::ShaderStages, label: &str, binding: u32, data: &[T]) -> Self {
-        Self {
-            buf: ArrBuf::init_with(
-                device,
-                &(label.to_owned() + " Storage"),
-                BufferUsages::STORAGE,
-                data
-            ),
-            visibility,
-            binding,
-        }
-    }
-}
-
-impl<T: bytemuck::Pod> Storage<T> {
-    pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry {
-        wgpu::BindGroupLayoutEntry {
-            binding: self.binding,
-            visibility: self.visibility,
-            ty: wgpu::BindingType::Buffer {
-                ty: wgpu::BufferBindingType::Storage { read_only: true },
-                has_dynamic_offset: false,
-                min_binding_size: None,
-            },
-            count: None,
-        }
-    }
-    pub fn bind_group_entry(&self) -> wgpu::BindGroupEntry {
-        return wgpu::BindGroupEntry {
-            binding: self.binding,
-            resource: self.buf.buffer().as_entire_binding(),
-        };
-    }
-    pub fn update(
-        &mut self,
-        device: &wgpu::Device,
-        encoder: &mut wgpu::CommandEncoder,
-        belt: &mut wgpu::util::StagingBelt,
-        size: usize,
-        updates: &[ArrBufUpdate<T>],
-    ) -> bool {
-        self.buf.update(device, encoder, belt, size, updates)
-    }
-
-    pub fn mov(&mut self, mov: BufMove) {
-        self.buf.mov(mov);
-    }
-
-    pub fn len(&mut self) -> usize {
-        self.buf.len()
-    }
-}

src/client/render/util/texture.rs

@@ -1,23 +1,26 @@
-pub struct DepthTexture {
+pub struct Texture {
+    texture_desc: wgpu::TextureDescriptor<'static>,
+    view_desc: wgpu::TextureViewDescriptor<'static>,
+    sampler_desc: wgpu::SamplerDescriptor<'static>,
     pub texture: wgpu::Texture,
     pub view: wgpu::TextureView,
     pub sampler: wgpu::Sampler,
 }
 
-impl DepthTexture {
+impl Texture {
     pub const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float;
 
-    pub fn init(
+    pub fn init_depth(
         device: &wgpu::Device,
         config: &wgpu::SurfaceConfiguration,
-        label: &str,
+        label: &'static str,
     ) -> Self {
         let size = wgpu::Extent3d {
             width: config.width + 1,
             height: config.height + 1,
             depth_or_array_layers: 1,
         };
-        let desc = wgpu::TextureDescriptor {
+        let texture_desc = wgpu::TextureDescriptor {
             label: Some(label),
             size,
             mip_level_count: 1,
@@ -27,97 +30,62 @@ impl DepthTexture {
             usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::TEXTURE_BINDING,
             view_formats: &[],
         };
-        let texture = device.create_texture(&desc);
+        Self::init(
-
+            device,
-        let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
+            texture_desc,
-        let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
+            wgpu::TextureViewDescriptor::default(),
-            address_mode_u: wgpu::AddressMode::ClampToEdge,
+            wgpu::SamplerDescriptor {
-            address_mode_v: wgpu::AddressMode::ClampToEdge,
+                address_mode_u: wgpu::AddressMode::ClampToEdge,
-            address_mode_w: wgpu::AddressMode::ClampToEdge,
+                address_mode_v: wgpu::AddressMode::ClampToEdge,
-            mag_filter: wgpu::FilterMode::Linear,
+                address_mode_w: wgpu::AddressMode::ClampToEdge,
-            min_filter: wgpu::FilterMode::Linear,
+                mag_filter: wgpu::FilterMode::Linear,
-            mipmap_filter: wgpu::FilterMode::Nearest,
+                min_filter: wgpu::FilterMode::Linear,
-            compare: Some(wgpu::CompareFunction::LessEqual),
+                mipmap_filter: wgpu::FilterMode::Nearest,
-            lod_min_clamp: 0.0,
+                compare: Some(wgpu::CompareFunction::LessEqual),
-            lod_max_clamp: 100.0,
+                lod_min_clamp: 0.0,
-            ..Default::default()
+                lod_max_clamp: 100.0,
-        });
+                ..Default::default()
+            },
+        )
+    }
 
+    pub fn init(
+        device: &wgpu::Device,
+        texture_desc: wgpu::TextureDescriptor<'static>,
+        view_desc: wgpu::TextureViewDescriptor<'static>,
+        sampler_desc: wgpu::SamplerDescriptor<'static>,
+    ) -> Self {
+        let texture = device.create_texture(&texture_desc);
+        let view = texture.create_view(&view_desc);
+        let sampler = device.create_sampler(&sampler_desc);
         Self {
+            texture_desc,
+            view_desc,
+            sampler_desc,
             texture,
             view,
             sampler,
         }
     }
-}
 
-pub struct StorageTexture {
+    pub fn resize(&mut self, device: &wgpu::Device, size: wgpu::Extent3d) {
-    binding: u32,
+        self.texture_desc.size = size;
-    visibility: wgpu::ShaderStages,
+        self.texture = device.create_texture(&self.texture_desc);
-    pub buf: wgpu::Texture,
+        self.view = self.texture.create_view(&self.view_desc);
-    pub view: wgpu::TextureView,
+    }
-    pub sampler: wgpu::Sampler,
+    pub fn view_bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry {
-}
+        wgpu::BindGroupEntry {
-
-impl StorageTexture {
-    pub const FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Rgba8Unorm;
-
-    pub fn init(
-        device: &wgpu::Device,
-        size: wgpu::Extent3d,
-        label: &str,
-        visibility: wgpu::ShaderStages,
-        binding: u32,
-    ) -> Self {
-        let desc = wgpu::TextureDescriptor {
-            label: Some(label),
-            size,
-            mip_level_count: 1,
-            sample_count: 1,
-            dimension: wgpu::TextureDimension::D2,
-            format: Self::FORMAT,
-            usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
-            view_formats: &[],
-        };
-        let texture = device.create_texture(&desc);
-
-        let view = texture.create_view(&wgpu::TextureViewDescriptor::default());
-        let sampler = device.create_sampler(&wgpu::SamplerDescriptor {
-            address_mode_u: wgpu::AddressMode::ClampToEdge,
-            address_mode_v: wgpu::AddressMode::ClampToEdge,
-            address_mode_w: wgpu::AddressMode::ClampToEdge,
-            mag_filter: wgpu::FilterMode::Linear,
-            min_filter: wgpu::FilterMode::Linear,
-            mipmap_filter: wgpu::FilterMode::Linear,
-            compare: None,
-            lod_min_clamp: 0.0,
-            lod_max_clamp: 100.0,
-            ..Default::default()
-        });
-        Self {
-            visibility,
             binding,
-            buf: texture,
-            view,
-            sampler,
-        }
-    }
-    pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry {
-        wgpu::BindGroupLayoutEntry {
-            binding: self.binding,
-            visibility: self.visibility,
-            ty: wgpu::BindingType::StorageTexture {
-                access: wgpu::StorageTextureAccess::WriteOnly,
-                format: Self::FORMAT,
-                view_dimension: wgpu::TextureViewDimension::D2,
-            },
-            count: None,
-        }
-    }
-    pub fn bind_group_entry(&self) -> wgpu::BindGroupEntry {
-        return wgpu::BindGroupEntry {
-            binding: self.binding,
             resource: wgpu::BindingResource::TextureView(&self.view),
-        };
+        }
+    }
+    pub fn sampler_bind_group_entry(&self, binding: u32) -> wgpu::BindGroupEntry {
+        wgpu::BindGroupEntry {
+            binding,
+            resource: wgpu::BindingResource::Sampler(&self.sampler),
+        }
+    }
+    pub fn format(&self) -> wgpu::TextureFormat {
+        self.texture_desc.format
     }
 }

src/client/render/util/timer.rs

@@ -0,0 +1,85 @@
+use std::time::Duration;
+
+pub struct GPUTimer {
+    resolve_buf: wgpu::Buffer,
+    map_buf: wgpu::Buffer,
+    query_set: wgpu::QuerySet,
+    timestamps: Vec<u64>,
+}
+
+impl GPUTimer {
+    pub fn new(device: &wgpu::Device, count: u32) -> Self {
+        let count = count * 2;
+        let timestamp_set = device.create_query_set(&wgpu::QuerySetDescriptor {
+            count,
+            label: Some("voxel timestamp"),
+            ty: wgpu::QueryType::Timestamp,
+        });
+        let timestamp_resolve_buf = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("voxel timestamp"),
+            mapped_at_creation: false,
+            size: 8 * count as u64,
+            usage: wgpu::BufferUsages::QUERY_RESOLVE | wgpu::BufferUsages::COPY_SRC,
+        });
+
+        let timestamp_mapped_buf = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("voxel timestamp"),
+            mapped_at_creation: false,
+            size: 8 * count as u64,
+            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+        });
+
+        Self {
+            query_set: timestamp_set,
+            resolve_buf: timestamp_resolve_buf,
+            map_buf: timestamp_mapped_buf,
+            timestamps: vec![0; count as usize],
+        }
+    }
+
+    pub fn resolve(&self, encoder: &mut wgpu::CommandEncoder) {
+        encoder.resolve_query_set(&self.query_set, 0..2, &self.resolve_buf, 0);
+        encoder.copy_buffer_to_buffer(&self.resolve_buf, 0, &self.map_buf, 0, self.map_buf.size());
+    }
+
+    pub fn duration(&self, i: u32) -> Duration {
+        let i = i as usize * 2;
+        Duration::from_nanos(self.timestamps[i + 1] - self.timestamps[i])
+    }
+
+    pub fn finish(&mut self, device: &wgpu::Device) {
+        let (s, r) = std::sync::mpsc::channel();
+        self.map_buf
+            .slice(..)
+            .map_async(wgpu::MapMode::Read, move |v| {
+                s.send(v).expect("what");
+            });
+        device.poll(wgpu::Maintain::wait()).panic_on_timeout();
+        if let Ok(Ok(())) = r.recv() {
+            let data = self.map_buf.slice(..).get_mapped_range();
+            self.timestamps.copy_from_slice(bytemuck::cast_slice(&data));
+            drop(data);
+            self.map_buf.unmap();
+        }
+    }
+
+    #[allow(dead_code)]
+    pub fn start(&self, encoder: &mut wgpu::CommandEncoder, i: u32) {
+        encoder.write_timestamp(&self.query_set, i * 2);
+    }
+
+    #[allow(dead_code)]
+    pub fn stop(&self, encoder: &mut wgpu::CommandEncoder, i: u32) {
+        encoder.write_timestamp(&self.query_set, i * 2 + 1);
+    }
+
+    #[allow(dead_code)]
+    pub fn start_compute(&self, pass: &mut wgpu::ComputePass, i: u32) {
+        pass.write_timestamp(&self.query_set, i * 2);
+    }
+
+    #[allow(dead_code)]
+    pub fn stop_compute(&self, pass: &mut wgpu::ComputePass, i: u32) {
+        pass.write_timestamp(&self.query_set, i * 2 + 1);
+    }
+}

src/client/render/util/vertex.rs

@@ -1,86 +0,0 @@
-use wgpu::{BufferUsages, VertexAttribute};
-
-use super::buf::{ArrBuf, ArrBufUpdate, BufMove};
-
-pub struct Vertices<T: bytemuck::Pod> {
-    buf: ArrBuf<T>,
-    location: u32,
-    attrs: [VertexAttribute; 1],
-}
-
-impl<T: bytemuck::Pod> Vertices<T> {
-    pub fn update(
-        &mut self,
-        device: &wgpu::Device,
-        encoder: &mut wgpu::CommandEncoder,
-        belt: &mut wgpu::util::StagingBelt,
-        size: usize,
-        updates: &[ArrBufUpdate<T>],
-    ) -> bool {
-        self.buf.update(device, encoder, belt, size, updates)
-    }
-
-    pub fn init(
-        device: &wgpu::Device,
-        label: &str,
-        location: u32,
-        format: wgpu::VertexFormat,
-    ) -> Self {
-        Self {
-            buf: ArrBuf::init(
-                device,
-                &(label.to_owned() + " Instance"),
-                BufferUsages::VERTEX,
-            ),
-            location,
-            attrs: [wgpu::VertexAttribute {
-                format,
-                offset: 0,
-                shader_location: location,
-            }],
-        }
-    }
-
-    pub fn init_with(
-        device: &wgpu::Device,
-        label: &str,
-        location: u32,
-        format: wgpu::VertexFormat,
-        data: &[T],
-    ) -> Self {
-        Self {
-            buf: ArrBuf::init_with(
-                device,
-                &(label.to_owned() + " Instance"),
-                BufferUsages::VERTEX,
-                data,
-            ),
-            location,
-            attrs: [wgpu::VertexAttribute {
-                format,
-                offset: 0,
-                shader_location: location,
-            }],
-        }
-    }
-
-    pub fn set_in<'a>(&'a self, render_pass: &mut wgpu::RenderPass<'a>) {
-        render_pass.set_vertex_buffer(self.location, self.buf.buffer().slice(..));
-    }
-
-    pub fn desc(&self) -> wgpu::VertexBufferLayout {
-        wgpu::VertexBufferLayout {
-            array_stride: std::mem::size_of::<T>() as wgpu::BufferAddress,
-            step_mode: wgpu::VertexStepMode::Vertex,
-            attributes: &self.attrs,
-        }
-    }
-
-    pub fn mov(&mut self, mov: BufMove) {
-        self.buf.mov(mov);
-    }
-
-    pub fn len(&self) -> usize {
-        self.buf.len()
-    }
-}

src/client/render/voxel/ray_oct/chunk.rs

@@ -0,0 +1,7 @@
+#[repr(C, align(16))]
+#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Zeroable)]
+pub struct Chunk {
+    pub offset: u32,
+}
+
+unsafe impl bytemuck::Pod for Chunk {}

src/client/render/voxel/ray_oct/group.rs

@@ -1,12 +0,0 @@
-use nalgebra::{Projective3, Vector3};
-
-#[repr(C, align(16))]
-#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Zeroable)]
-pub struct VoxelGroup {
-    pub transform: Projective3<f32>,
-    pub transform_inv: Projective3<f32>,
-    pub scale: u32,
-    pub offset: u32,
-}
-
-unsafe impl bytemuck::Pod for VoxelGroup {}

src/client/render/voxel/ray_oct/layout.rs

@@ -1,18 +1,18 @@
-use wgpu::TextureFormat;
+use wgpu::{BufferUsages, ShaderStages, TextureFormat};
 
-use super::{group::VoxelGroup, light::GlobalLight, view::View};
+use super::{chunk::Chunk, light::GlobalLight, view::View};
 use crate::{
-    client::render::util::{Storage, StorageTexture, Uniform},
+    client::render::util::{ArrayBuffer, Texture, Uniform},
     util::oct_tree::OctNode,
 };
 use nalgebra::Vector3;
 
 pub struct Layout {
-    pub texture: StorageTexture,
     pub view: Uniform<View>,
-    pub voxel_groups: Storage<VoxelGroup>,
+    pub chunks: ArrayBuffer<Chunk>,
-    pub voxels: Storage<OctNode>,
+    pub voxel_data: ArrayBuffer<OctNode>,
-    pub global_lights: Storage<GlobalLight>,
+    pub global_lights: ArrayBuffer<GlobalLight>,
+    pub texture: Texture,
     render_bind_layout: wgpu::BindGroupLayout,
     compute_bind_layout: wgpu::BindGroupLayout,
     render_pipeline_layout: wgpu::PipelineLayout,
@@ -23,27 +23,40 @@ pub struct Layout {
 impl Layout {
     pub fn init(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self {
         let view = Uniform::init(device, "view", 0);
-        let voxels = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxels", 1);
+        let chunks = ArrayBuffer::init(device, "chunks", BufferUsages::STORAGE);
-        let voxel_groups = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxel groups", 2);
+        let voxel_data = ArrayBuffer::init_with(
-        let global_lights = Storage::init_with(
+            device,
+            "voxel data",
+            BufferUsages::STORAGE,
+            &[OctNode::new_leaf(0)],
+        );
+        let global_lights = ArrayBuffer::init_with(
             device,
-            wgpu::ShaderStages::COMPUTE,
             "global lights",
-            3,
+            BufferUsages::STORAGE,
             &[GlobalLight {
                 direction: Vector3::new(-1.0, -2.3, 2.0).normalize(),
             }],
         );
-        let texture = StorageTexture::init(
+        let desc = wgpu::TextureDescriptor {
-            device,
+            label: Some("compute output"),
-            wgpu::Extent3d {
+            size: wgpu::Extent3d {
                 width: config.width,
                 height: config.height,
                 depth_or_array_layers: 1,
             },
-            "compute output",
+            mip_level_count: 1,
-            wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
+            sample_count: 1,
-            4,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Rgba8Unorm,
+            usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
+            view_formats: &[],
+        };
+        let texture = Texture::init(
+            device,
+            desc,
+            wgpu::TextureViewDescriptor::default(),
+            wgpu::SamplerDescriptor::default(),
         );
         let render_bind_layout =
             device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
@@ -73,10 +86,31 @@ impl Layout {
             device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
                 entries: &[
                     view.bind_group_layout_entry(),
-                    voxels.bind_group_layout_entry(),
+                    chunks.bind_group_layout_entry(
-                    voxel_groups.bind_group_layout_entry(),
+                        1,
-                    global_lights.bind_group_layout_entry(),
+                        ShaderStages::COMPUTE,
-                    texture.bind_group_layout_entry(),
+                        wgpu::BufferBindingType::Storage { read_only: true },
+                    ),
+                    voxel_data.bind_group_layout_entry(
+                        2,
+                        ShaderStages::COMPUTE,
+                        wgpu::BufferBindingType::Storage { read_only: true },
+                    ),
+                    global_lights.bind_group_layout_entry(
+                        3,
+                        ShaderStages::COMPUTE,
+                        wgpu::BufferBindingType::Storage { read_only: true },
+                    ),
+                    wgpu::BindGroupLayoutEntry {
+                        binding: 4,
+                        visibility: ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::StorageTexture {
+                            access: wgpu::StorageTextureAccess::WriteOnly,
+                            format: texture.format(),
+                            view_dimension: wgpu::TextureViewDimension::D2,
+                        },
+                        count: None,
+                    },
                 ],
                 label: Some("voxel compute"),
             });
@@ -94,8 +128,8 @@ impl Layout {
             });
         Self {
             view,
-            voxels,
+            voxel_data,
-            voxel_groups,
+            chunks,
             global_lights,
             texture,
             render_bind_layout,
@@ -110,16 +144,10 @@ impl Layout {
         device.create_bind_group(&wgpu::BindGroupDescriptor {
             layout: &self.render_bind_layout,
             entries: &[
-                wgpu::BindGroupEntry {
+                self.texture.view_bind_group_entry(0),
-                    binding: 0,
+                self.texture.sampler_bind_group_entry(1),
-                    resource: wgpu::BindingResource::TextureView(&self.texture.view),
-                },
-                wgpu::BindGroupEntry {
-                    binding: 1,
-                    resource: wgpu::BindingResource::Sampler(&self.texture.sampler),
-                },
             ],
-            label: Some("tile_bind_group"),
+            label: Some("voxel render"),
         })
     }
 
@@ -128,10 +156,10 @@ impl Layout {
             layout: &self.compute_bind_layout,
             entries: &[
                 self.view.bind_group_entry(),
-                self.voxels.bind_group_entry(),
+                self.chunks.bind_group_entry(1),
-                self.voxel_groups.bind_group_entry(),
+                self.voxel_data.bind_group_entry(2),
-                self.global_lights.bind_group_entry(),
+                self.global_lights.bind_group_entry(3),
-                self.texture.bind_group_entry(),
+                self.texture.view_bind_group_entry(4),
             ],
             label: Some("voxel compute"),
         })

src/client/render/voxel/ray_oct/mod.rs

@@ -1,28 +1,25 @@
+mod chunk;
 mod color;
 mod grid;
-mod group;
 mod layout;
 mod light;
 mod view;
 
 use super::super::UpdateGridTransform;
-use crate::{
+use crate::client::{
-    client::{
+    camera::Camera,
-        camera::Camera,
+    render::{AddChunk, CreateVoxelGrid},
-        render::{
-            util::{ArrBufUpdate, StorageTexture},
-            AddChunk, CreateVoxelGrid,
-        },
-    },
-    common::component::chunk,
 };
 use bevy_ecs::entity::Entity;
 pub use color::*;
 use layout::Layout;
-use nalgebra::{Projective3, Transform3, Translation3, Vector2, Vector3};
+use nalgebra::{Transform3, Translation3, Vector2};
-use std::{collections::HashMap, ops::Deref};
+use std::{
+    collections::HashMap,
+    sync::{Arc, Mutex},
+};
 use wgpu::include_wgsl;
-use {group::VoxelGroup, view::View};
+use {chunk::Chunk, view::View};
 
 pub struct VoxelPipeline {
     layout: Layout,
@@ -30,7 +27,7 @@ pub struct VoxelPipeline {
     compute_bind_group: wgpu::BindGroup,
     render_pipeline: wgpu::RenderPipeline,
     render_bind_group: wgpu::BindGroup,
-    id_map: HashMap<Entity, (usize, VoxelGroup)>,
+    id_map: HashMap<Entity, (usize, Chunk)>,
 }
 
 const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/render.wgsl");
@@ -82,7 +79,8 @@ impl VoxelPipeline {
 
     pub fn update_shader(&mut self, device: &wgpu::Device) {
         let Ok(shader) = std::fs::read_to_string(
-            env!("CARGO_MANIFEST_DIR").to_owned() + "/src/client/render/voxel/ray_oct/shader/compute.wgsl",
+            env!("CARGO_MANIFEST_DIR").to_owned()
+                + "/src/client/render/voxel/ray_oct/shader/compute.wgsl",
         ) else {
             println!("Failed to reload shader!");
             return;
@@ -110,49 +108,27 @@ impl VoxelPipeline {
         belt: &mut wgpu::util::StagingBelt,
         AddChunk { id, pos, tree, .. }: AddChunk,
     ) {
-        let offset = self.layout.voxels.len();
+        let offset = self.layout.voxel_data.len() as u32;
-
-        let data = tree.raw();
-        let updates = [ArrBufUpdate { offset, data }];
-        let size = offset + data.len();
         self.layout
-            .voxels
+            .voxel_data
-            .update(device, encoder, belt, size, &updates);
+            .add(device, encoder, belt, tree.raw());
 
-        let proj = Projective3::identity()
+        let chunk = Chunk { offset };
-            * Translation3::from((pos.deref() * chunk::SIDE_LENGTH as i32).cast())
+        let i = self.layout.chunks.len();
-            * Translation3::from(-chunk::DIMENSIONS.cast() / 2.0);
+        self.layout.chunks.add(device, encoder, belt, &[chunk]);
-        let group = VoxelGroup {
-            transform: proj,
-            transform_inv: proj.inverse(),
-            scale: chunk::SCALE,
-            offset: offset as u32,
-        };
-        let updates = [ArrBufUpdate {
-            offset: self.layout.voxel_groups.len(),
-            data: &[group],
-        }];
-        let i = self.layout.voxel_groups.len();
-        let size = i + 1;
-        self.layout
-            .voxel_groups
-            .update(device, encoder, belt, size, &updates);
 
-        self.id_map.insert(id, (i, group));
+        self.id_map.insert(id, (i, chunk));
         self.compute_bind_group = self.layout.compute_bind_group(device);
     }
 
     pub fn resize(&mut self, device: &wgpu::Device, size: Vector2<u32>) {
-        self.layout.texture = StorageTexture::init(
+        self.layout.texture.resize(
             device,
             wgpu::Extent3d {
                 width: size.x,
                 height: size.y,
                 depth_or_array_layers: 1,
             },
-            "idk man im tired",
-            wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
-            4,
         );
         self.compute_bind_group = self.layout.compute_bind_group(device);
         self.render_bind_group = self.layout.render_bind_group(device);
@@ -166,21 +142,7 @@ impl VoxelPipeline {
         update: UpdateGridTransform,
     ) {
         if let Some((i, group)) = self.id_map.get_mut(&update.id) {
-            let offset = Vector3::from_element(-(2u32.pow(group.scale) as f32) / 2.0);
+            self.layout.chunks.set(device, encoder, belt, *i, &[*group]);
-            let proj = Projective3::identity()
-                * Translation3::from(update.pos)
-                * update.orientation
-                * Translation3::from(offset);
-            group.transform = proj;
-            group.transform_inv = proj.inverse();
-            let updates = [ArrBufUpdate {
-                offset: *i,
-                data: &[*group],
-            }];
-            let size = self.layout.voxel_groups.len();
-            self.layout
-                .voxel_groups
-                .update(device, encoder, belt, size, &updates);
         }
     }
 
@@ -189,7 +151,6 @@ impl VoxelPipeline {
         device: &wgpu::Device,
         encoder: &mut wgpu::CommandEncoder,
         belt: &mut wgpu::util::StagingBelt,
-        size: Vector2<u32>,
         camera: &Camera,
     ) {
         let transform =
@@ -197,6 +158,7 @@ impl VoxelPipeline {
         let data = View {
             zoom: camera.scale,
             transform,
+            ..Default::default()
         };
         self.layout.view.update(device, encoder, belt, data);
     }
@@ -212,7 +174,7 @@ impl VoxelPipeline {
     pub fn compute(&self, pass: &mut wgpu::ComputePass) {
         pass.set_pipeline(&self.compute_pipeline);
         pass.set_bind_group(0, &self.compute_bind_group, &[]);
-        let buf = &self.layout.texture.buf;
+        let buf = &self.layout.texture.texture;
         let x = (buf.width() - 1) / Self::WORKGROUP_SIZE + 1;
         let y = (buf.height() - 1) / Self::WORKGROUP_SIZE + 1;
         pass.dispatch_workgroups(x, y, 1);

src/client/render/voxel/ray_oct/shader/compute.wgsl

@@ -1,9 +1,9 @@
 @group(0) @binding(0)
 var<uniform> view: View;
 @group(0) @binding(1)
-var<storage, read> voxels: array<u32>;
+var<storage, read> chunks: array<Chunk>;
 @group(0) @binding(2)
-var<storage, read> voxel_groups: array<VoxelGroup>;
+var<storage, read> voxel_data: array<u32>;
 @group(0) @binding(3)
 var<storage, read> global_lights: array<GlobalLight>;
 @group(0) @binding(4)
@@ -16,12 +16,11 @@ struct GlobalLight {
 struct View {
     transform: mat4x4<f32>,
     zoom: f32,
+    chunk_scale: u32,
+    chunk_dist: u32,
 };
 
-struct VoxelGroup {
+struct Chunk {
-    transform: mat4x4<f32>,
-    transform_inv: mat4x4<f32>,
-    scale: u32,
     offset: u32,
 };
 
@@ -38,7 +37,8 @@ fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
     let pixel_pos = vec2<f32>(
         (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect)
     );
-    let pos = view.transform * vec4<f32>(pixel_pos, 1.0, 1.0);
+    let offset = vec3<f32>(f32(1u << (view.chunk_scale - 1)));
+    let pos = view.transform * vec4<f32>(pixel_pos, 1.0, 1.0) + vec4<f32>(offset, 0.0);
     let dir = view.transform * vec4<f32>(normalize(vec3<f32>(pixel_pos, view.zoom)), 0.0);
 
     var color = trace_full(pos, dir);
@@ -53,6 +53,7 @@ fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
 
 const LEAF_BIT = 1u << 31u;
 const LEAF_MASK = ~LEAF_BIT;
+const MAX_HITS = 10;
 
 const ZERO3F = vec3<f32>(0.0);
 const ZERO2F = vec2<f32>(0.0);
@@ -62,126 +63,122 @@ const MAX_ITERS = 10000;
 // NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
 // numbers are stored and the bit manipulation used
 const MAX_SCALE: u32 = 13;
-const AMBIENT: f32 = 0.2;
-const SPECULAR: f32 = 0.5;
-
 fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    let gi = 0;
+    if arrayLength(&voxel_data) == 1 {
-    let group = voxel_groups[gi];
-    if group.scale == 0 {
         return vec4<f32>(0.0);
     }
-    let dimensions = vec3<u32>(1u << group.scale);
+    let gi = 0;
+    let chunk = chunks[gi];
+    let side_len = 1u << view.chunk_scale;
+    let dimensions = vec3<u32>(side_len);
     let dim_f = vec3<f32>(dimensions);
-    let dim_i = vec3<i32>(dimensions);
 
-    // transform so that group is at 0,0
+    let pos_start = pos_view.xyz;
-    let pos_start = (group.transform_inv * pos_view).xyz;
+    var dir = dir_view.xyz;
-    var dir = (group.transform_inv * dir_view).xyz;
+    if dir.x == 0 { dir.x = EPSILON; }
-    if dir.x == 0 {dir.x = EPSILON;}
+    if dir.y == 0 { dir.y = EPSILON; }
-    if dir.y == 0 {dir.y = EPSILON;}
+    if dir.z == 0 { dir.z = EPSILON; }
-    if dir.z == 0 {dir.z = EPSILON;}
 
     let dir_if = sign(dir);
     let dir_uf = max(dir_if, vec3<f32>(0.0));
 
+    // find where ray intersects with group
+    // closest (min) and furthest (max) corners of cube relative to direction
+    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
+    let pos_max = dir_uf * dim_f;
+    // time of intersection; x = td + p, solve for t
+    let t_min = (pos_min - pos_start) / dir;
+    let t_max = (pos_max - pos_start) / dir;
+    // time of entrance and exit of the cube
+    let t_start = max(max(t_min.x, t_min.y), t_min.z);
+    let t_end = min(min(t_max.x, t_max.y), t_max.z);
+    if t_end < t_start { return vec4<f32>(0.0); }
+    // axis of intersection
+    let axis = select(select(2u, 1u, t_start == t_min.y), 0u, t_start == t_min.x);
+    // time to move entire side length in each direction
+    let inc_t = abs(1.0 / dir) * f32(side_len);
+    let t = max(0.0, t_start);
 
+    let inv_dir_bits = 7 - vec_to_dir(vec3<u32>(dir_uf));
+    let corner_adj = t_min - inc_t;
 
     // calculate normals
     var normals = mat3x3<f32>(
-        (group.transform * vec4<f32>(dir_if.x, 0.0, 0.0, 0.0)).xyz,
+        vec3<f32>(dir_if.x, 0.0, 0.0),
-        (group.transform * vec4<f32>(0.0, dir_if.y, 0.0, 0.0)).xyz,
+        vec3<f32>(0.0, dir_if.y, 0.0),
-        (group.transform * vec4<f32>(0.0, 0.0, dir_if.z, 0.0)).xyz,
+        vec3<f32>(0.0, 0.0, dir_if.z),
     );
-    var axis = 0u;
 
-    // find where ray intersects with group
+    let result = cast_ray(chunk.offset, t, axis, inv_dir_bits, inc_t, corner_adj);
-    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
+    return shade_ray(result, pos_start, dir_view.xyz, t_end, normals);
-    // time of intersection; x = td + p, solve for t
+}
-    var t_min = (pos_min - pos_start) / dir;
-    if outside3f(pos_start, ZERO3F, dim_f) {
-        // points of intersection
-        let px = pos_start + t_min.x * dir;
-        let py = pos_start + t_min.y * dir;
-        let pz = pos_start + t_min.z * dir;
 
-        // check if point is in bounds
+fn shade_ray(result: RayResult, pos_start: vec3<f32>, dir: vec3<f32>, t_end: f32, normals: mat3x3<f32>) -> vec4<f32> {
-        let hit = vec3<bool>(
+    var hits = result.hits;
-            inside2f(px.yz, ZERO2F, dim_f.yz),
-            inside2f(py.xz, ZERO2F, dim_f.xz),
-            inside2f(pz.xy, ZERO2F, dim_f.xy),
-        ) && (t_min > ZERO3F);
-        if !any(hit) {
-            return vec4<f32>(0.0);
-        }
-        axis = select(select(2u, 1u, hit.y), 0u, hit.x);
-    }
-    let t_mult = f32(1u << (MAX_SCALE - group.scale));
-    t_min *= t_mult;
-    // time to move 1 unit in each direction
-    let full = f32(1u << MAX_SCALE);
-    let inc_t = abs(1.0 / dir) * full;
-    let t_offset = max(max(t_min.x, t_min.y), t_min.z);
-    var t = max(0.0, t_offset);
 
-    let dir_i = vec3<i32>(dir_if);
-    let dir_u = vec3<u32>(dir_uf);
-    let dir_bits = vec_to_dir(dir_u);
-    let inv_dir_bits = 7 - dir_bits;
-
-    var node_start = 1u;
-    var scale = MAX_SCALE - 1;
-    var scale_exp2 = 0.5;
     var color = vec4<f32>(0.0);
-    var parents = array<u32, MAX_SCALE>();
+    for (var i = 0u; i < result.len; i += 1u) {
-    var prev = LEAF_BIT;
+        let hit = hits[i];
-    var old_t = t / t_mult;
+        let id = hit.id;
+        let t = hit.t;
+        let axis = hit.axis;
 
-    var child = 0u;
+        let next_t = select(hits[i + 1].t, t_end, i == result.len - 1);
+
+        var pos = pos_start + dir * t;
+        pos[axis] = round(pos[axis]) - f32(dir[axis] < 0.0);
+        let normal = select(select(normals[0], normals[1], axis == 1), normals[2], axis == 2);
+        let vcolor = shade(id, pos, normal, dir, next_t - t);
+        color += vcolor * (1.0 - color.a);
+        if color.a > FULL_ALPHA { break; }
+    }
+    return color;
+}
+
+struct RayHit {
+    t: f32,
+    id: u32,
+    axis: u32,
+}
+
+struct RayResult {
+    hits: array<RayHit, MAX_HITS>,
+    len: u32,
+}
+
+fn cast_ray(
+    data_offset: u32, t_start: f32, axis_start: u32,
+    inv_dir_bits: u32, inc_t: vec3<f32>, corner_adj: vec3<f32>
+) -> RayResult {
+    var hits = array<RayHit, MAX_HITS>();
+    var depth = 0u;
+    var min_alpha = 0.0;
+
+    var t = t_start;
+    var axis = axis_start;
+    var node_start = 0u;
+    var scale = MAX_SCALE;
+    var scale_exp2 = 1.0;
+    var parents = array<u32, MAX_SCALE>();
+    var child = inv_dir_bits;
     var vox_pos = vec3<f32>(1.0);
-    let t_center = t_min + scale_exp2 * inc_t;
+    var prev = 0u;
-    if t > t_center.x { vox_pos.x = 1.5; child |= 4u; }
-    if t > t_center.y { vox_pos.y = 1.5; child |= 2u; }
-    if t > t_center.z { vox_pos.z = 1.5; child |= 1u; }
-    let min_adj = t_min - inc_t;
 
     var iters = 0;
     loop {
-        if iters == MAX_ITERS {
+        if iters == MAX_ITERS { break; }
-            return vec4<f32>(1.0, 0.0, 1.0, 1.0);
-        }
         iters += 1;
-        let t_corner = vox_pos * inc_t + min_adj;
+        let t_corner = vox_pos * inc_t + corner_adj;
-        let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
+        let node = voxel_data[data_offset + node_start + (child ^ inv_dir_bits)];
         if node >= LEAF_BIT {
+            // ignore consecutive identical leaves
             if node != prev {
-                if node != LEAF_BIT {
+                let id = node & LEAF_MASK;
-                    let real_t = t / t_mult;
+                hits[depth] = RayHit(t, id, axis);
-                    let dist = real_t - old_t;
+                min_alpha += min_alpha(id) * (1.0 - min_alpha);
-                    old_t = real_t;
+                depth += 1u;
-                    let filt = min(dist / 64.0, 1.0);
-                    if prev == LEAF_BIT + 3 {
-                        color.a += filt * (1.0 - color.a);
-                        if color.a > FULL_ALPHA { break; }
-                    }
-                    var pos = (pos_view + dir_view * real_t).xyz;
-                    pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]);
-                    let vcolor = get_color(node & LEAF_MASK, pos);
-                    var normal = normals[axis];
-                    let light_color = vec3<f32>(1.0);
-                    let light_dir = global_lights[0].dir;
-
-                    let diffuse = max(dot(light_dir, normal), 0.0) * light_color;
-                    let ambient = AMBIENT * light_color;
-                    let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR;
-                    let specular = spec_val * light_color;
-                    let new_color = (ambient + diffuse + specular) * vcolor.xyz;
-                    let new_a = min(vcolor.a + spec_val, 1.0);
-
-                    color += vec4<f32>(new_color.xyz * new_a, new_a) * (1.0 - color.a);
-                    if color.a > FULL_ALPHA { break; }
-                }
                 prev = node;
+                if depth == 10 || min_alpha >= FULL_ALPHA { break; }
             }
 
             // move to next time point and determine which axis to move along
@@ -226,10 +223,15 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
             node_start = node;
         }
     }
-    // let fog = min(t / t_mult / 1000.0, 1.0);
+    return RayResult(hits, depth);
-    // return vec4<f32>(color.xyz * (1.0 - fog) + vec3<f32>(fog), color.a * (1.0 - fog) + fog);
+}
-    // return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
+
-    return color;
+fn trace_chunk(
+    offset: u32,
+    inv_dir_bits: u32,
+    t: f32, t_mult: f32, inc_t: vec3<f32>,
+    min_adj: vec3<f32>
+) {
 }
 
 fn dir_to_vec(bits: u32) -> vec3<u32> {
@@ -240,29 +242,53 @@ fn vec_to_dir(vec: vec3<u32>) -> u32 {
     return vec.x * 4 + vec.y * 2 + vec.z * 1;
 }
 
-fn get_color(id: u32, pos: vec3<f32>) -> vec4<f32> {
+fn min_alpha(id: u32) -> f32 {
+    switch id {
+        case 0u: {return 0.0;}
+        case 3u: {return 0.5;}
+        default: {return 1.0;}
+    }
+}
+
+const AMBIENT: f32 = 0.2;
+const SPECULAR: f32 = 0.5;
+
+// returns premultiplied
+fn shade(id: u32, pos: vec3<f32>, normal: vec3<f32>, dir_view: vec3<f32>, dist: f32) -> vec4<f32> {
+    var color = vec4<f32>(0.0);
+    if id == 0 {
+        return color;
+    }
     let random = random(floor(pos));
     let random2 = random(floor(pos) + vec3<f32>(0.0001));
     switch id {
         case 0u: {
-            return vec4<f32>(0.0);
+            color = vec4<f32>(0.0);
         }
         case 1u: {
-            let color = vec3<f32>(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8);
+            color = vec4<f32>(vec3<f32>(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8), 1.0);
-            return vec4<f32>(color, 1.0);
         }
         case 2u: {
-            let color = vec3<f32>(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9);
+            color = vec4<f32>(vec3<f32>(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9), 1.0);
-            return vec4<f32>(color, 1.0);
         }
         case 3u: {
-            let color = vec3<f32>(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8);
+            let fog = min(dist / 64.0, 1.0);
-            return vec4<f32>(color, 0.5);
+            let a = 0.5;
-        }
+            let rgb = vec3<f32>(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8);
-        default: {
+            color = vec4<f32>(rgb * (1.0 - fog * a), a + fog * (1.0 - a));
-            return vec4<f32>(1.0, 0.0, 0.0, 1.0);
         }
+        default: {}
     }
+    let light_color = vec3<f32>(1.0);
+    let light_dir = global_lights[0].dir;
+
+    let diffuse = max(dot(light_dir, normal), 0.0) * light_color;
+    let ambient = AMBIENT * light_color;
+    let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR;
+    let specular = spec_val * light_color;
+    let new_color = (ambient + diffuse + specular) * color.xyz;
+    let new_a = min(color.a + spec_val, 1.0);
+    return vec4<f32>(new_color * new_a, new_a);
 }
 
 fn random(pos: vec3<f32>) -> f32 {

src/client/render/voxel/ray_oct/shader/compute_shadow.wgsl

@@ -1,419 +0,0 @@
-@group(0) @binding(0)
-var<uniform> view: View;
-@group(0) @binding(1)
-var<storage, read> voxels: array<u32>;
-@group(0) @binding(2)
-var<storage, read> voxel_groups: array<VoxelGroup>;
-@group(0) @binding(3)
-var<storage, read> global_lights: array<GlobalLight>;
-@group(0) @binding(4)
-var output: texture_storage_2d<rgba8unorm, write>;
-
-struct GlobalLight {
-    dir: vec3<f32>,
-};
-
-struct View {
-    transform: mat4x4<f32>,
-    zoom: f32,
-};
-
-struct VoxelGroup {
-    transform: mat4x4<f32>,
-    transform_inv: mat4x4<f32>,
-    scale: u32,
-    offset: u32,
-};
-
-@compute
-@workgroup_size(8, 8, 1)
-fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
-    let view_dim = textureDimensions(output);
-    // get position of the pixel; eye at origin, pixel on plane z = 1
-    if cell.x >= view_dim.x || cell.y >= view_dim.y {
-        return;
-    }
-    let view_dim_f = vec2<f32>(view_dim);
-    let aspect = view_dim_f.y / view_dim_f.x;
-    let pixel_pos = vec3<f32>(
-        (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
-        view.zoom
-    );
-    let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
-    let dir = view.transform * vec4<f32>(normalize(pixel_pos), 0.0);
-
-    var color = trace_full(pos, dir);
-    let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0);
-    let sun_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
-    let sky_bg = vec3<f32>(0.3, 0.6, 1.0);
-    let sky_color = sun_color + sky_bg * (1.0 - light_mult);
-    color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
-    color.a = 1.0;
-    textureStore(output, cell.xy, color);
-}
-
-const LEAF_BIT = 1u << 31u;
-const LEAF_MASK = ~LEAF_BIT;
-
-const ZERO3F = vec3<f32>(0.0);
-const ZERO2F = vec2<f32>(0.0);
-const FULL_ALPHA = 0.999;
-const EPSILON = 0.00000000001;
-const MAX_ITERS = 2000;
-// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
-// numbers are stored and the bit manipulation used
-const MAX_SCALE: u32 = 10;
-
-fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    let gi = 0;
-    let group = voxel_groups[gi];
-    if group.scale == 0 {
-        return vec4<f32>(0.0);
-    }
-    let dimensions = vec3<u32>(1u << group.scale);
-    let dim_f = vec3<f32>(dimensions);
-    let dim_i = vec3<i32>(dimensions);
-
-    // transform so that group is at 0,0
-    let pos_start = (group.transform_inv * pos_view).xyz;
-    var dir = (group.transform_inv * dir_view).xyz;
-    if dir.x == 0 {dir.x = EPSILON;}
-    if dir.y == 0 {dir.y = EPSILON;}
-    if dir.z == 0 {dir.z = EPSILON;}
-
-    let dir_if = sign(dir);
-    let dir_uf = max(dir_if, vec3<f32>(0.0));
-
-
-
-    // calculate normals
-    var normals = mat3x3<f32>(
-        (group.transform * vec4<f32>(dir_if.x, 0.0, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, dir_if.y, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, 0.0, dir_if.z, 0.0)).xyz,
-    );
-    var axis = 0u;
-
-    // find where ray intersects with group
-    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
-    // time of intersection; x = td + p, solve for t
-    var t_min = (pos_min - pos_start) / dir;
-    if outside3f(pos_start, ZERO3F, dim_f) {
-        // points of intersection
-        let px = pos_start + t_min.x * dir;
-        let py = pos_start + t_min.y * dir;
-        let pz = pos_start + t_min.z * dir;
-
-        // check if point is in bounds
-        let hit = vec3<bool>(
-            inside2f(px.yz, ZERO2F, dim_f.yz),
-            inside2f(py.xz, ZERO2F, dim_f.xz),
-            inside2f(pz.xy, ZERO2F, dim_f.xy),
-        ) && (t_min > ZERO3F);
-        if !any(hit) {
-            return vec4<f32>(0.0);
-        }
-        axis = select(select(2u, 1u, hit.y), 0u, hit.x);
-    }
-    let t_mult = f32(1u << (MAX_SCALE - group.scale));
-    t_min *= t_mult;
-    // time to move 1 unit in each direction
-    let full = f32(1u << MAX_SCALE);
-    let inc_t = abs(1.0 / dir) * full;
-    let t_offset = max(max(t_min.x, t_min.y), t_min.z);
-    var t = max(0.0, t_offset);
-
-    let dir_i = vec3<i32>(dir_if);
-    let dir_u = vec3<u32>((dir_i + vec3<i32>(1)) / 2);
-    let dir_bits = vec_to_dir(dir_u);
-    let inv_dir_bits = 7 - dir_bits;
-
-    var node_start = 1u;
-    var scale = MAX_SCALE - 1;
-    var scale_exp2 = 0.5;
-    var skip = LEAF_BIT;
-    var color = vec4<f32>(0.0);
-    var parents = array<u32, MAX_SCALE>();
-
-    var child = 0u;
-    var vox_pos = vec3<f32>(1.0);
-    let t_center = t_min + scale_exp2 * inc_t;
-    if t > t_center.x { vox_pos.x = 1.5; child |= 4u; }
-    if t > t_center.y { vox_pos.y = 1.5; child |= 2u; }
-    if t > t_center.z { vox_pos.z = 1.5; child |= 1u; }
-    let min_adj = t_min - inc_t;
-
-    var iters = 0;
-    loop {
-        if iters == MAX_ITERS {
-            return vec4<f32>(1.0, 0.0, 1.0, 1.0);
-        }
-        iters += 1;
-        let t_corner = vox_pos * inc_t + min_adj;
-        let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
-        if node >= LEAF_BIT {
-            if node != skip && node != LEAF_BIT {
-                skip = node;
-                let normal = normals[axis];
-                let sun_dir = global_lights[0].dir;
-                let new_pos = pos_view + dir_view * t / t_mult - vec4<f32>(normals[axis] * 0.001, 0.0);
-
-                let light = trace_light(new_pos, vec4<f32>(-sun_dir, 0.0));
-                let diffuse = max(dot(sun_dir, normal) + 0.1, 0.0);
-                let ambient = 0.2;
-                let specular = (exp(max(
-                    -(dot(reflect(dir_view.xyz, normal), sun_dir) + 0.90) * 4.0, 0.0
-                )) - 1.0) * light;
-                let lighting = max(diffuse * light.a, ambient);
-
-                let vcolor = get_color(node & LEAF_MASK);
-                let new_rgb = min(vcolor.xyz * lighting + specular.xyz + light.xyz * vcolor.xyz, vec3<f32>(1.0));
-                let new_a = min(vcolor.a + specular.a, 1.0);
-                let new_color = vec4<f32>(new_rgb, new_a);
-                color += vec4<f32>(new_color.xyz * new_color.a, new_color.a) * (1.0 - color.a);
-                if color.a > FULL_ALPHA { break; }
-            }
-
-            // move to next time point and determine which axis to move along
-            let t_next = t_corner + scale_exp2 * inc_t;
-            t = min(min(t_next.x, t_next.y), t_next.z);
-            axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z);
-            let move_dir = 4u >> axis;
-
-            // check if need to pop stack
-            if (child & move_dir) > 0 {
-                // calculate new scale; first differing bit after adding
-                let axis_pos = vox_pos[axis];
-                // AWARE
-                let differing = bitcast<u32>(axis_pos) ^ bitcast<u32>(axis_pos + scale_exp2);
-                scale = (bitcast<u32>(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE);
-                scale_exp2 = bitcast<f32>((scale + 127 - MAX_SCALE) << 23);
-                if scale >= MAX_SCALE { break; }
-
-                // restore & recalculate parent
-                let parent_info = parents[scale];
-                node_start = parent_info >> 3;
-                child = parent_info & 7;
-                let scale_vec = vec3<u32>(scale + 23 - MAX_SCALE);
-                // remove bits lower than current scale
-                vox_pos = bitcast<vec3<f32>>((bitcast<vec3<u32>>(vox_pos) >> scale_vec) << scale_vec);
-            }
-            // move to next child and voxel position
-            child += move_dir;
-            vox_pos[axis] += scale_exp2;
-        } else {
-            // push current node to stack
-            parents[scale] = (node_start << 3) + child;
-            scale -= 1u;
-
-            // calculate child node vars
-            scale_exp2 *= 0.5;
-            child = 0u;
-            let t_center = t_corner + scale_exp2 * inc_t;
-            if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; }
-            if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; }
-            if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; }
-            node_start += 8 + node;
-        }
-    }
-    // return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
-    return color;
-}
-
-fn trace_light(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    let gi = 0;
-    let group = voxel_groups[gi];
-    if group.scale == 0 {
-        return vec4<f32>(0.0);
-    }
-    let dimensions = vec3<u32>(1u << group.scale);
-    let dim_f = vec3<f32>(dimensions);
-    let dim_i = vec3<i32>(dimensions);
-
-    // transform so that group is at 0,0
-    let pos_start = (group.transform_inv * pos_view).xyz;
-    var dir = (group.transform_inv * dir_view).xyz;
-    if dir.x == 0 {dir.x = EPSILON;}
-    if dir.y == 0 {dir.y = EPSILON;}
-    if dir.z == 0 {dir.z = EPSILON;}
-
-    let dir_if = sign(dir);
-    let dir_uf = max(dir_if, vec3<f32>(0.0));
-
-
-
-    // calculate normals
-    var normals = mat3x3<f32>(
-        (group.transform * vec4<f32>(dir_if.x, 0.0, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, dir_if.y, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, 0.0, dir_if.z, 0.0)).xyz,
-    );
-    var axis = 0u;
-
-    // find where ray intersects with group
-    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
-    // time of intersection; x = td + p, solve for t
-    var t_min = (pos_min - pos_start) / dir;
-    if outside3f(pos_start, ZERO3F, dim_f) {
-        // points of intersection
-        let px = pos_start + t_min.x * dir;
-        let py = pos_start + t_min.y * dir;
-        let pz = pos_start + t_min.z * dir;
-
-        // check if point is in bounds
-        let hit = vec3<bool>(
-            inside2f(px.yz, ZERO2F, dim_f.yz),
-            inside2f(py.xz, ZERO2F, dim_f.xz),
-            inside2f(pz.xy, ZERO2F, dim_f.xy),
-        ) && (t_min > ZERO3F);
-        if !any(hit) {
-            return vec4<f32>(0.0);
-        }
-        axis = select(select(2u, 1u, hit.y), 0u, hit.x);
-    }
-    let t_mult = f32(1u << (MAX_SCALE - group.scale));
-    t_min *= t_mult;
-    // time to move 1 unit in each direction
-    let full = f32(1u << MAX_SCALE);
-    let inc_t = abs(1.0 / dir) * full;
-    let t_offset = max(max(t_min.x, t_min.y), t_min.z);
-    var t = max(0.0, t_offset);
-    var old_t = t;
-
-    let dir_i = vec3<i32>(dir_if);
-    let dir_u = vec3<u32>((dir_i + vec3<i32>(1)) / 2);
-    let dir_bits = vec_to_dir(dir_u);
-    let inv_dir_bits = 7 - dir_bits;
-
-    var node_start = 1u;
-    var scale = MAX_SCALE - 1;
-    var scale_exp2 = 0.5;
-    var mask = vec4<f32>(0.0);
-    var skip = LEAF_BIT;
-    var parents = array<u32, MAX_SCALE>();
-
-    var child = 0u;
-    var vox_pos = vec3<f32>(1.0);
-    let t_center = t_min + scale_exp2 * inc_t;
-    if t > t_center.x { vox_pos.x = 1.5; child |= 4u; }
-    if t > t_center.y { vox_pos.y = 1.5; child |= 2u; }
-    if t > t_center.z { vox_pos.z = 1.5; child |= 1u; }
-    let min_adj = t_min - inc_t;
-    var data = 0u;
-
-    var iters = 0;
-    loop {
-        if iters == MAX_ITERS {
-            return vec4<f32>(1.0, 0.0, 1.0, 1.0);
-        }
-        iters += 1;
-        let t_corner = vox_pos * inc_t + min_adj;
-        let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
-        if node >= LEAF_BIT {
-            if node != skip && node != LEAF_BIT {
-                skip = node;
-                if data == 3 {
-                    let dist = (t - old_t) / t_mult;
-                    let vcolor = vec4<f32>(vec3<f32>(0.0), min(dist / 12.0, 1.0));
-                    mask += vec4<f32>(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a);
-                }
-                data = node & LEAF_MASK;
-                if data != 3 && data != 0 {
-                    let vcolor = get_color(data);
-                    mask += vec4<f32>(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a);
-                }
-                old_t = t;
-                if mask.a > FULL_ALPHA { break; }
-            }
-
-            // move to next time point and determine which axis to move along
-            let t_next = t_corner + scale_exp2 * inc_t;
-            t = min(min(t_next.x, t_next.y), t_next.z);
-            axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z);
-            let move_dir = 4u >> axis;
-
-            // check if need to pop stack
-            if (child & move_dir) > 0 {
-                // calculate new scale; first differing bit after adding
-                let axis_pos = vox_pos[axis];
-                // AWARE
-                let differing = bitcast<u32>(axis_pos) ^ bitcast<u32>(axis_pos + scale_exp2);
-                scale = (bitcast<u32>(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE);
-                scale_exp2 = bitcast<f32>((scale + 127 - MAX_SCALE) << 23);
-                if scale >= MAX_SCALE { break; }
-
-                // restore & recalculate parent
-                let parent_info = parents[scale];
-                node_start = parent_info >> 3;
-                child = parent_info & 7;
-                let scale_vec = vec3<u32>(scale + 23 - MAX_SCALE);
-                // remove bits lower than current scale
-                vox_pos = bitcast<vec3<f32>>((bitcast<vec3<u32>>(vox_pos) >> scale_vec) << scale_vec);
-            }
-            // move to next child and voxel position
-            child += move_dir;
-            vox_pos[axis] += scale_exp2;
-        } else {
-            // push current node to stack
-            parents[scale] = (node_start << 3) + child;
-            scale -= 1u;
-
-            // calculate child node vars
-            scale_exp2 *= 0.5;
-            child = 0u;
-            let t_center = t_corner + scale_exp2 * inc_t;
-            if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; }
-            if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; }
-            if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; }
-            node_start += 8 + node;
-        }
-    }
-    if data == 3 {
-        let dist = (t - old_t) / t_mult;
-        let vcolor = vec4<f32>(vec3<f32>(0.0), min(dist / 12.0, 1.0));
-        mask += vec4<f32>(vcolor.xyz * vcolor.a, vcolor.a) * (1.0 - mask.a);
-    }
-    mask.a = 1.0 - mask.a;
-    mask = vec4<f32>(mask.a * mask.xyz, mask.a);
-    return mask;
-}
-
-fn dir_to_vec(bits: u32) -> vec3<u32> {
-    return vec3<u32>(bits >> 2, (bits & 2) >> 1, bits & 1);
-}
-
-fn vec_to_dir(vec: vec3<u32>) -> u32 {
-    return vec.x * 4 + vec.y * 2 + vec.z * 1;
-}
-
-fn get_color(id: u32) -> vec4<f32> {
-    switch id {
-        case 0u: {
-            return vec4<f32>(0.0);
-        }
-        case 1u: {
-            return vec4<f32>(0.5, 0.5, 0.5, 1.0);
-        }
-        case 2u: {
-            return vec4<f32>(0.5, 1.0, 0.5, 1.0);
-        }
-        case 3u: {
-            return vec4<f32>(0.5, 0.5, 1.0, 0.5);
-        }
-        default: {
-            return vec4<f32>(1.0, 0.0, 0.0, 1.0);
-        }
-    }
-}
-
-fn outside3f(v: vec3<f32>, low: vec3<f32>, high: vec3<f32>) -> bool {
-    return any(v < low) || any(v > high);
-}
-
-fn inside2f(v: vec2<f32>, low: vec2<f32>, high: vec2<f32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}
-
-fn inside3i(v: vec3<i32>, low: vec3<i32>, high: vec3<i32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}

src/client/render/voxel/ray_oct/shader/compute_single.wgsl

@@ -35,67 +35,13 @@ fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
     }
     let view_dim_f = vec2<f32>(view_dim);
     let aspect = view_dim_f.y / view_dim_f.x;
-    let pixel_pos = vec3<f32>(
+    let pixel_pos = vec2<f32>(
-        (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
+        (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect)
-        view.zoom
     );
-    let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
+    let pos = view.transform * vec4<f32>(pixel_pos, 1.0, 1.0);
-    let dir = view.transform * vec4<f32>(normalize(pixel_pos), 0.0);
+    let dir = view.transform * vec4<f32>(normalize(vec3<f32>(pixel_pos, view.zoom)), 0.0);
 
-    let start = start_ray(pos, dir);
+    var color = trace_full(pos, dir);
-    var color = vec4<f32>(0.0);
-    let ambient = 0.2;
-
-    if start.hit {
-        var res = ray_next(start.ray, LEAF_BIT);
-        var normals = start.normals;
-        let specular = (exp(max(
-            -(dot(reflect(dir.xyz, normals[res.ray.axis]), global_lights[0].dir) + 0.90) * 4.0, 0.0
-        )) - 1.0);
-        while res.data != 0 {
-            let data = res.data & LEAF_MASK;
-            if data != 0 {
-                let vcolor = get_color(data);
-                let diffuse = max(dot(global_lights[0].dir, normals[res.ray.axis]) + 0.1, 0.0);
-                let light = max(diffuse, ambient);
-                let new_color = min(vcolor.xyz * light, vec3<f32>(1.0));
-                color += vec4<f32>(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a);
-                if color.a > FULL_ALPHA { break; }
-            }
-            let old_t = res.ray.t;
-            res = ray_next(res.ray, res.data);
-            if data == 3 {
-                let dist = (res.ray.t - old_t) / start.t_mult;
-                let a = min(dist / 12.0, 1.0);
-                color += vec4<f32>(vec3<f32>(0.0) * a, a) * (1.0 - color.a);
-            }
-        }
-        if color.a != 0 {
-            let pos = pos + dir * res.ray.t / start.t_mult - vec4<f32>(normals[res.ray.axis] * 0.001, 0.0);
-            let dir = vec4<f32>(-global_lights[0].dir, 0.0);
-            let start = start_ray(pos, dir);
-            res = ray_next(start.ray, LEAF_BIT);
-            var light = 1.0;
-            while res.data != 0 {
-                let data = res.data & LEAF_MASK;
-                if data != 0 {
-                    let vcolor = get_color(data);
-                    if data != 3 { light -= vcolor.a * light; }
-                    if light <= 0 { break; }
-                }
-                let old_t = res.ray.t;
-                res = ray_next(res.ray, res.data);
-                if data == 3 {
-                    let dist = (res.ray.t - old_t) / start.t_mult;
-                    let a = min(dist / 12.0, 1.0);
-                    light -= a;
-                }
-            }
-            color = vec4<f32>(color.xyz * max(light, ambient), color.a) + vec4<f32>(vec3<f32>(specular * light), 0.0);
-        }
-        // color = vec4<f32>(pos.xyz / 128.0, 1.0);
-    }
-    // var color = trace_full(pos, dir);
     let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0);
     let sun_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
     let sky_bg = vec3<f32>(0.3, 0.6, 1.0);
@@ -112,49 +58,25 @@ const ZERO3F = vec3<f32>(0.0);
 const ZERO2F = vec2<f32>(0.0);
 const FULL_ALPHA = 0.999;
 const EPSILON = 0.00000000001;
-const MAX_ITERS = 2000;
+const MAX_ITERS = 10000;
 // NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
 // numbers are stored and the bit manipulation used
-const MAX_SCALE: u32 = 10;
+const MAX_SCALE: u32 = 13;
+const AMBIENT: f32 = 0.2;
+const SPECULAR: f32 = 0.5;
 
-struct Ray {
+fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    t: f32,
-    vox_pos: vec3<f32>,
-    t_inc: vec3<f32>,
-    scale: u32,
-    min_adj: vec3<f32>,
-    child: u32,
-    axis: u32,
-    node_start: u32,
-    group_offset: u32,
-    inv_dir_bits: u32,
-    parents: array<u32, MAX_SCALE>,
-};
-
-struct RayResult {
-    ray: Ray,
-    data: u32,
-}
-
-struct RayStart {
-    hit: bool,
-    ray: Ray,
-    normals: mat3x3<f32>,
-    t_mult: f32,
-}
-
-fn start_ray(pos_view: vec4<f32>, dir_view: vec4<f32>) -> RayStart {
     let gi = 0;
     let group = voxel_groups[gi];
     if group.scale == 0 {
-        return RayStart();
+        return vec4<f32>(0.0);
     }
     let dimensions = vec3<u32>(1u << group.scale);
     let dim_f = vec3<f32>(dimensions);
     let dim_i = vec3<i32>(dimensions);
 
     // transform so that group is at 0,0
-    let pos = (group.transform_inv * pos_view).xyz;
+    let pos_start = (group.transform_inv * pos_view).xyz;
     var dir = (group.transform_inv * dir_view).xyz;
     if dir.x == 0 {dir.x = EPSILON;}
     if dir.y == 0 {dir.y = EPSILON;}
@@ -176,12 +98,12 @@ fn start_ray(pos_view: vec4<f32>, dir_view: vec4<f32>) -> RayStart {
     // find where ray intersects with group
     let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
     // time of intersection; x = td + p, solve for t
-    var t_min = (pos_min - pos) / dir;
+    var t_min = (pos_min - pos_start) / dir;
-    if outside3f(pos, ZERO3F, dim_f) {
+    if outside3f(pos_start, ZERO3F, dim_f) {
         // points of intersection
-        let px = pos + t_min.x * dir;
+        let px = pos_start + t_min.x * dir;
-        let py = pos + t_min.y * dir;
+        let py = pos_start + t_min.y * dir;
-        let pz = pos + t_min.z * dir;
+        let pz = pos_start + t_min.z * dir;
 
         // check if point is in bounds
         let hit = vec3<bool>(
@@ -190,81 +112,80 @@ fn start_ray(pos_view: vec4<f32>, dir_view: vec4<f32>) -> RayStart {
             inside2f(pz.xy, ZERO2F, dim_f.xy),
         ) && (t_min > ZERO3F);
         if !any(hit) {
-            return RayStart();
+            return vec4<f32>(0.0);
         }
         axis = select(select(2u, 1u, hit.y), 0u, hit.x);
     }
     let t_mult = f32(1u << (MAX_SCALE - group.scale));
     t_min *= t_mult;
-    // time to move 1 unit in each direction
     let full = f32(1u << MAX_SCALE);
-    let t_inc = abs(1.0 / dir) * full;
+    // time to move entire side length in each direction
+    let inc_t = abs(1.0 / dir) * full;
     let t_offset = max(max(t_min.x, t_min.y), t_min.z);
-    let t = max(0.0, t_offset);
+    var t = max(0.0, t_offset);
 
     let dir_i = vec3<i32>(dir_if);
-    let dir_u = vec3<u32>((dir_i + vec3<i32>(1)) / 2);
+    let dir_u = vec3<u32>(dir_uf);
     let dir_bits = vec_to_dir(dir_u);
     let inv_dir_bits = 7 - dir_bits;
 
-    let node_start = 1u;
+    var node_start = 1u;
-    let scale = MAX_SCALE - 1;
+    var scale = MAX_SCALE - 1;
-    let scale_exp2 = 0.5;
+    var scale_exp2 = 0.5;
-    let parents = array<u32, MAX_SCALE>();
+    var color = vec4<f32>(0.0);
+    var parents = array<u32, MAX_SCALE>();
+    var prev = LEAF_BIT;
+    var old_t = t / t_mult;
 
     var child = 0u;
     var vox_pos = vec3<f32>(1.0);
-    let t_center = t_min + scale_exp2 * t_inc;
+    let t_center = t_min + scale_exp2 * inc_t;
     if t > t_center.x { vox_pos.x = 1.5; child |= 4u; }
     if t > t_center.y { vox_pos.y = 1.5; child |= 2u; }
     if t > t_center.z { vox_pos.z = 1.5; child |= 1u; }
-    let min_adj = t_min - t_inc;
+    let min_adj = t_min - inc_t;
 
-    return RayStart(
+    var iters = 0;
-        true,
-        Ray(
-            t,
-            vox_pos,
-            t_inc,
-            scale,
-            min_adj,
-            child,
-            axis,
-            node_start,
-            group.offset,
-            inv_dir_bits,
-            parents,
-        ),
-        normals,
-        t_mult,
-    );
-}
-
-fn ray_next(ray: Ray, skip: u32) -> RayResult {
-    let group_offset = ray.group_offset;
-    let t_inc = ray.t_inc;
-    let min_adj = ray.min_adj;
-    let inv_dir_bits = ray.inv_dir_bits;
-    var scale = ray.scale;
-    var scale_exp2 = bitcast<f32>((scale + 127 - MAX_SCALE) << 23);
-    var vox_pos = ray.vox_pos;
-    var t = ray.t;
-    var node_start = ray.node_start;
-    var child = ray.child;
-    var parents = ray.parents;
-    var axis: u32;
-    var data = 0u;
     loop {
-        let t_corner = vox_pos * t_inc + min_adj;
+        if iters == MAX_ITERS {
-        let node = voxels[group_offset + node_start + (child ^ inv_dir_bits)];
+            return vec4<f32>(1.0, 0.0, 1.0, 1.0);
+        }
+        iters += 1;
+        let t_corner = vox_pos * inc_t + min_adj;
+        let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
         if node >= LEAF_BIT {
-            if node != skip {
+            if node != prev {
-                data = node;
+                if node != LEAF_BIT {
-                break;
+                    let real_t = t / t_mult;
+                    let dist = real_t - old_t;
+                    old_t = real_t;
+                    let filt = min(dist / 64.0, 1.0);
+                    if prev == LEAF_BIT + 3 {
+                        color.a += filt * (1.0 - color.a);
+                        if color.a > FULL_ALPHA { break; }
+                    }
+                    var pos = (pos_view + dir_view * real_t).xyz;
+                    pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]);
+                    let vcolor = get_color(node & LEAF_MASK, pos);
+                    var normal = normals[axis];
+                    let light_color = vec3<f32>(1.0);
+                    let light_dir = global_lights[0].dir;
+
+                    let diffuse = max(dot(light_dir, normal), 0.0) * light_color;
+                    let ambient = AMBIENT * light_color;
+                    let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR;
+                    let specular = spec_val * light_color;
+                    let new_color = (ambient + diffuse + specular) * vcolor.xyz;
+                    let new_a = min(vcolor.a + spec_val, 1.0);
+
+                    color += vec4<f32>(new_color.xyz * new_a, new_a) * (1.0 - color.a);
+                    if color.a > FULL_ALPHA { break; }
+                }
+                prev = node;
             }
 
             // move to next time point and determine which axis to move along
-            let t_next = t_corner + scale_exp2 * t_inc;
+            let t_next = t_corner + scale_exp2 * inc_t;
             t = min(min(t_next.x, t_next.y), t_next.z);
             axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z);
             let move_dir = 4u >> axis;
@@ -298,29 +219,17 @@ fn ray_next(ray: Ray, skip: u32) -> RayResult {
             // calculate child node vars
             scale_exp2 *= 0.5;
             child = 0u;
-            let t_center = t_corner + scale_exp2 * t_inc;
+            let t_center = t_corner + scale_exp2 * inc_t;
             if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; }
             if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; }
             if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; }
-            node_start += 8 + node;
+            node_start = node;
         }
     }
-    return RayResult(
+    // let fog = min(t / t_mult / 1000.0, 1.0);
-        Ray(
+    // return vec4<f32>(color.xyz * (1.0 - fog) + vec3<f32>(fog), color.a * (1.0 - fog) + fog);
-            t,
+    // return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
-            vox_pos,
+    return color;
-            t_inc,
-            scale,
-            min_adj,
-            child,
-            axis,
-            node_start,
-            group_offset,
-            inv_dir_bits,
-            parents,
-        ),
-        data
-    );
 }
 
 fn dir_to_vec(bits: u32) -> vec3<u32> {
@@ -331,19 +240,24 @@ fn vec_to_dir(vec: vec3<u32>) -> u32 {
     return vec.x * 4 + vec.y * 2 + vec.z * 1;
 }
 
-fn get_color(id: u32) -> vec4<f32> {
+fn get_color(id: u32, pos: vec3<f32>) -> vec4<f32> {
+    let random = random(floor(pos));
+    let random2 = random(floor(pos) + vec3<f32>(0.0001));
     switch id {
         case 0u: {
             return vec4<f32>(0.0);
         }
         case 1u: {
-            return vec4<f32>(0.5, 0.5, 0.5, 1.0);
+            let color = vec3<f32>(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8);
+            return vec4<f32>(color, 1.0);
         }
         case 2u: {
-            return vec4<f32>(0.5, 1.0, 0.5, 1.0);
+            let color = vec3<f32>(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9);
+            return vec4<f32>(color, 1.0);
         }
         case 3u: {
-            return vec4<f32>(0.5, 0.5, 1.0, 0.5);
+            let color = vec3<f32>(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8);
+            return vec4<f32>(color, 0.5);
         }
         default: {
             return vec4<f32>(1.0, 0.0, 0.0, 1.0);
@@ -351,6 +265,10 @@ fn get_color(id: u32) -> vec4<f32> {
     }
 }
 
+fn random(pos: vec3<f32>) -> f32 {
+    return fract(sin(dot(pos, vec3<f32>(12.9898, 78.233, 25.1279))) * 43758.5453123);
+}
+
 fn outside3f(v: vec3<f32>, low: vec3<f32>, high: vec3<f32>) -> bool {
     return any(v < low) || any(v > high);
 }

src/client/render/voxel/ray_oct/view.rs

@@ -1,10 +1,14 @@
 use nalgebra::Transform3;
 
+use crate::common::component::chunk::SCALE;
+
 #[repr(C, align(16))]
 #[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
 pub struct View {
     pub transform: Transform3<f32>,
     pub zoom: f32,
+    pub chunk_scale: u32,
+    pub chunk_radius: u32,
 }
 
 unsafe impl bytemuck::Pod for View {}
@@ -14,6 +18,8 @@ impl Default for View {
         Self {
             zoom: 1.0,
             transform: Transform3::identity(),
+            chunk_scale: SCALE,
+            chunk_radius: 2,
         }
     }
 }

src/common/component/chunk/mod.rs

@@ -8,7 +8,7 @@ use bevy_derive::{Deref, DerefMut};
 use bevy_ecs::{bundle::Bundle, component::Component, entity::Entity, system::Resource};
 use nalgebra::Vector3;
 
-pub const SCALE: u32 = 13;
+pub const SCALE: u32 = 9;
 pub const SIDE_LENGTH: usize = 2usize.pow(SCALE);
 pub const SHAPE: (usize, usize, usize) = (SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);
 pub const DIMENSIONS: Vector3<usize> = Vector3::new(SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);

src/util/oct_tree.rs

@@ -86,7 +86,7 @@ impl OctTree {
         assert!(levels > 0);
         let mut data = Vec::new();
         let mut map = OctNodeMap::default();
-        data.push(OctNode::new_node(0));
+        data.push(OctNode::new_node(1));
         Self::from_fn_offset_inner(f_leaf, f_node, &mut data, levels, offset, &mut map);
         if data.len() == 2 {
             data.remove(0);

src/util/timer.rs

@@ -25,6 +25,11 @@ impl Timer {
         self.times[self.pos] = Some(self.start);
         self.pos = (self.pos + 1) % self.times.len();
     }
+    pub fn add(&mut self, duration: Duration) {
+        self.durs[self.pos] = Some(duration);
+        self.times[self.pos] = Some(Instant::now());
+        self.pos = (self.pos + 1) % self.times.len();
+    }
     pub fn avg(&self) -> Duration {
         let filtered: Vec<_> = self.durs.iter().filter_map(|d| *d).collect();
         let len = filtered.len();
@@ -45,8 +50,9 @@ impl Timer {
     pub fn per_sec(&self) -> usize {
         let now = Instant::now();
         let mut count = 0;
-        while count < self.times.len() {
+        let len = self.times.len();
-            let i = (self.pos + count + 1) % self.times.len();
+        while count < len {
+            let i = (self.pos + len - count - 1) % len;
             let Some(t) = self.times[i] else { break };
             if now - t <= Duration::from_secs(1) {
                 count += 1;