From b6db483189e3b96e0b79ece647bbe0e47d370cb1 Mon Sep 17 00:00:00 2001 From: shadow cat Date: Wed, 11 Sep 2024 01:54:00 -0400 Subject: [PATCH] switched to stack and compute shaders --- src/client/handle_input.rs | 2 +- src/client/render/mod.rs | 16 +- src/client/render/util/storage.rs | 9 +- src/client/render/util/texture.rs | 86 +++++- src/client/render/util/uniform.rs | 2 +- .../ray_oct/{shader_old.wgsl => compute.wgsl} | 101 +++---- src/client/render/voxel/ray_oct/mod.rs | 173 +++++++++-- src/client/render/voxel/ray_oct/render.wgsl | 47 +++ src/client/render/voxel/ray_oct/shader.wgsl | 138 +++------ .../render/voxel/ray_oct/shader_broken.wgsl | 277 ------------------ src/util/oct_tree.rs | 33 +-- 11 files changed, 373 insertions(+), 511 deletions(-) rename src/client/render/voxel/ray_oct/{shader_old.wgsl => compute.wgsl} (80%) create mode 100644 src/client/render/voxel/ray_oct/render.wgsl delete mode 100644 src/client/render/voxel/ray_oct/shader_broken.wgsl diff --git a/src/client/handle_input.rs b/src/client/handle_input.rs index 315f954..132327f 100644 --- a/src/client/handle_input.rs +++ b/src/client/handle_input.rs @@ -84,7 +84,7 @@ impl Client<'_> { } // camera position - let move_dist = 2.0 * 16.0 * dt; + let move_dist = 64.0 * dt; if input.pressed(Key::KeyW) { state.camera.pos += *state.camera.forward() * move_dist; } diff --git a/src/client/render/mod.rs b/src/client/render/mod.rs index dd484af..4bdc6ec 100644 --- a/src/client/render/mod.rs +++ b/src/client/render/mod.rs @@ -6,7 +6,7 @@ pub use command::*; use super::camera::Camera; use crate::client::rsc::CLEAR_COLOR; use nalgebra::Vector2; -use util::Texture; +use util::DepthTexture; use voxel::VoxelPipeline; use winit::dpi::PhysicalSize; @@ -20,7 +20,7 @@ pub struct Renderer<'a> { staging_belt: wgpu::util::StagingBelt, voxel_pipeline: VoxelPipeline, camera: Camera, - depth_texture: Texture, + depth_texture: DepthTexture, } impl<'a> Renderer<'a> { @@ -78,7 +78,7 @@ 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 = Texture::create_depth_texture(&device, &config, "depth_texture"); + let depth_texture = DepthTexture::init(&device, &config, "depth_texture"); Self { camera: Camera::default(), @@ -106,6 +106,13 @@ impl<'a> Renderer<'a> { let view = output .texture .create_view(&wgpu::TextureViewDescriptor::default()); + let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: None, + timestamp_writes: None, + }); + self.voxel_pipeline.compute(&mut compute_pass, self.config.width, self.config.height); + drop(compute_pass); + let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("Render Pass"), color_attachments: &[Some(wgpu::RenderPassColorAttachment { @@ -141,9 +148,10 @@ impl<'a> Renderer<'a> { self.config.width = size.width; self.config.height = size.height; self.surface.configure(&self.device, &self.config); + self.voxel_pipeline.resize(&self.device, self.size); self.depth_texture = - Texture::create_depth_texture(&self.device, &self.config, "depth_texture"); + DepthTexture::init(&self.device, &self.config, "depth_texture"); self.voxel_pipeline.update_view( &self.device, &mut self.encoder, diff --git a/src/client/render/util/storage.rs b/src/client/render/util/storage.rs index 9b73cd4..a18537a 100644 --- a/src/client/render/util/storage.rs +++ b/src/client/render/util/storage.rs @@ -4,20 +4,22 @@ use wgpu::BufferUsages; pub struct Storage { binding: u32, buf: ArrBuf, + visibility: wgpu::ShaderStages, } impl Storage { - pub fn init(device: &wgpu::Device, label: &str, binding: u32) -> Self { + 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, label: &str, binding: u32, data: &[T]) -> Self { + pub fn init_with(device: &wgpu::Device, visibility: wgpu::ShaderStages, label: &str, binding: u32, data: &[T]) -> Self { Self { buf: ArrBuf::init_with( device, @@ -25,6 +27,7 @@ impl Storage { BufferUsages::STORAGE, data ), + visibility, binding, } } @@ -34,7 +37,7 @@ impl Storage { pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry { wgpu::BindGroupLayoutEntry { binding: self.binding, - visibility: wgpu::ShaderStages::FRAGMENT, + visibility: self.visibility, ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Storage { read_only: true }, has_dynamic_offset: false, diff --git a/src/client/render/util/texture.rs b/src/client/render/util/texture.rs index a72b490..8615f83 100644 --- a/src/client/render/util/texture.rs +++ b/src/client/render/util/texture.rs @@ -1,19 +1,18 @@ -pub struct Texture { +pub struct DepthTexture { pub texture: wgpu::Texture, pub view: wgpu::TextureView, pub sampler: wgpu::Sampler, } -impl Texture { - pub const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float; // 1. +impl DepthTexture { + pub const DEPTH_FORMAT: wgpu::TextureFormat = wgpu::TextureFormat::Depth32Float; - pub fn create_depth_texture( + pub fn init( device: &wgpu::Device, config: &wgpu::SurfaceConfiguration, label: &str, ) -> Self { let size = wgpu::Extent3d { - // 2. width: config.width, height: config.height, depth_or_array_layers: 1, @@ -25,22 +24,20 @@ impl Texture { sample_count: 1, dimension: wgpu::TextureDimension::D2, format: Self::DEPTH_FORMAT, - usage: wgpu::TextureUsages::RENDER_ATTACHMENT // 3. - | wgpu::TextureUsages::TEXTURE_BINDING, + usage: wgpu::TextureUsages::RENDER_ATTACHMENT | 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 { - // 4. 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::Nearest, - compare: Some(wgpu::CompareFunction::LessEqual), // 5. + compare: Some(wgpu::CompareFunction::LessEqual), lod_min_clamp: 0.0, lod_max_clamp: 100.0, ..Default::default() @@ -53,3 +50,74 @@ impl Texture { } } } + +pub struct StorageTexture { + binding: u32, + visibility: wgpu::ShaderStages, + pub texture: wgpu::Texture, + pub view: wgpu::TextureView, + pub sampler: wgpu::Sampler, +} + +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::Nearest, + compare: None, + lod_min_clamp: 0.0, + lod_max_clamp: 100.0, + ..Default::default() + }); + Self { + visibility, + binding, + 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), + }; + } +} diff --git a/src/client/render/util/uniform.rs b/src/client/render/util/uniform.rs index c78342d..48f59aa 100644 --- a/src/client/render/util/uniform.rs +++ b/src/client/render/util/uniform.rs @@ -37,7 +37,7 @@ impl Uniform { pub fn bind_group_layout_entry(&self) -> wgpu::BindGroupLayoutEntry { wgpu::BindGroupLayoutEntry { binding: self.binding, - visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT, + visibility: wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::COMPUTE, ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, has_dynamic_offset: false, diff --git a/src/client/render/voxel/ray_oct/shader_old.wgsl b/src/client/render/voxel/ray_oct/compute.wgsl similarity index 80% rename from src/client/render/voxel/ray_oct/shader_old.wgsl rename to src/client/render/voxel/ray_oct/compute.wgsl index 403c812..12e3c49 100644 --- a/src/client/render/voxel/ray_oct/shader_old.wgsl +++ b/src/client/render/voxel/ray_oct/compute.wgsl @@ -1,13 +1,18 @@ -// Vertex shader +@group(0) @binding(0) +var view: View; +@group(0) @binding(1) +var voxels: array; +@group(0) @binding(2) +var voxel_groups: array; +@group(0) @binding(3) +var global_lights: array; +@group(0) @binding(4) +var output: texture_storage_2d; struct GlobalLight { dir: vec3, }; -struct VertexOutput { - @builtin(position) clip_position: vec4, -}; - struct View { transform: mat4x4, width: u32, @@ -22,45 +27,16 @@ struct VoxelGroup { offset: u32, }; -@group(0) @binding(0) -var view: View; -@group(0) @binding(1) -var voxels: array; -@group(0) @binding(2) -var voxel_groups: array; -@group(0) @binding(3) -var global_lights: array; - -@vertex -fn vs_main( - @builtin(vertex_index) vi: u32, - @builtin(instance_index) ii: u32, -) -> VertexOutput { - var out: VertexOutput; - - var pos = vec2( - f32(vi % 2u) * 2.0 - 1.0, - f32(vi / 2u) * 2.0 - 1.0, - ) ; - out.clip_position = vec4(pos.x, pos.y, 0.0, 1.0); - return out; -} - -// Fragment shader - -@fragment -fn fs_main( - in: VertexOutput, -) -> @location(0) vec4 { +@compute +@workgroup_size(16, 16, 1) +fn main(@builtin(global_invocation_id) cell: vec3) { // get position of the pixel; eye at origin, pixel on plane z = 1 let win_dim = vec2(f32(view.width), f32(view.height)); let aspect = win_dim.y / win_dim.x; let pixel_pos = vec3( - (in.clip_position.xy / win_dim - vec2(0.5)) * vec2(2.0, -2.0 * aspect), + (vec2(cell.xy) / win_dim - vec2(0.5)) * vec2(2.0, 2.0 * aspect), view.zoom ); - - // move to position in world let pos = view.transform * vec4(pixel_pos, 1.0); let dir = view.transform * vec4(normalize(pixel_pos), 0.0); @@ -69,7 +45,7 @@ fn fs_main( let sky_color = light_mult * vec3(1.0, 1.0, 1.0); color += vec4(sky_color * (1.0 - color.a), 1.0 - color.a); color.a = 1.0; - return color; + textureStore(output, cell.xy, color); } const ZERO3F = vec3(0.0); @@ -145,20 +121,22 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { var data_start = 1u; var i = 0u; var axis = 0; - var hits = 0; + var parents = array(); + var scale = 0; for (var safety = 0; safety < 1000; safety += 1) { let node = voxels[group.offset + i]; if node >= LEAF_BIT { // leaf - hits += 1; - let vcolor = get_color(node & LEAF_MASK); - if vcolor.a > 0.0 { + let leaf = node & LEAF_MASK; + if leaf != 0 { + let vcolor = get_color(leaf); let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); let ambient = 0.2; let lighting = max(diffuse, ambient); let new_color = min(vcolor.xyz * lighting, vec3(1.0)); color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); if color.a > .999 { + // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); return color; } } @@ -177,33 +155,34 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { vox_pos[axis] += dir_i[axis]; } else if inside3i(vox_pos, vec3(0), vec3(side_len - 1)) { // node - let node_pos = data_start + node; + parents[scale] = (data_start << 3) + (data_start - i - 1); + scale += 1; + + let children_pos = data_start + node; side_len /= 2; let vcorner = vox_pos / side_len; - vox_pos -= vcorner * side_len; - let j = u32(vcorner.x * 4 + vcorner.y * 2 + vcorner.z); - i = node_pos + j; - data_start = node_pos + 9; + let child_pos = u32(vcorner.x * 4 + vcorner.y * 2 + vcorner.z); + i = children_pos + child_pos; + data_start = children_pos + 8; - low_corner += vec3(dir_to_vec(j)) * i32(side_len); + vox_pos -= vcorner * side_len; + low_corner += vec3(dir_to_vec(child_pos)) * i32(side_len); continue; } - // idrk what to put here tbh but this prolly works; don't zoom out if max - if side_len == 256 { + // exit if highest node + if scale == 0 { + // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); return color; } // get parent info and reset "pointers" to parent - let parent_info_i = data_start - 1; - let parent_info = voxels[group.offset + parent_info_i]; - let parent_root = parent_info_i - (parent_info >> 3); - let parent_loc = parent_info & 7; - let loc = 8 - (data_start - 1 - i); - // let test = (parent_root + 9 + voxels[group.offset + parent_root + parent_loc] + loc) == i; - i = parent_root + parent_loc; - data_start = parent_root + 9; + scale -= 1; + let parent_info = parents[scale]; + let loc = 8 - (data_start - i); + data_start = parent_info >> 3; + i = data_start - ((parent_info & 7) + 1); // adjust corner back to parent let low_corner_adj = vec3(dir_to_vec(loc)) * i32(side_len); @@ -213,8 +192,6 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { vox_pos += low_corner_adj; side_len *= 2; - // return vec4(vec3(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0); - // return vec4(vec3(f32(test)), 1.0); } return vec4(1.0, 0.0, 1.0, 1.0); } @@ -258,7 +235,7 @@ fn get_color(id: u32) -> vec4 { return vec4(0.5, 0.5, 0.5, 1.0); } case 2u: { - return vec4(0.5, 1.0, 0.5, 1.0); + return vec4(0.8, 0.2, 0.2, 1.0); } case 3u: { return vec4(0.5, 0.5, 1.0, 0.5); diff --git a/src/client/render/voxel/ray_oct/mod.rs b/src/client/render/voxel/ray_oct/mod.rs index 75aad27..7666bbd 100644 --- a/src/client/render/voxel/ray_oct/mod.rs +++ b/src/client/render/voxel/ray_oct/mod.rs @@ -8,13 +8,17 @@ pub use color::*; use wgpu::include_wgsl; use super::super::UpdateGridTransform; -use crate::{client::{ - camera::Camera, - render::{ - util::{ArrBufUpdate, Storage, Texture, Uniform}, - AddChunk, CreateVoxelGrid, +use crate::{ + client::{ + camera::Camera, + render::{ + util::{ArrBufUpdate, DepthTexture, Storage, StorageTexture, Uniform}, + AddChunk, CreateVoxelGrid, + }, }, -}, common::component::chunk, util::oct_tree::OctNode}; + common::component::chunk, + util::oct_tree::OctNode, +}; use bevy_ecs::entity::Entity; use light::GlobalLight; use nalgebra::{Projective3, Transform3, Translation3, Vector2, Vector3}; @@ -23,6 +27,11 @@ use std::{collections::HashMap, ops::Deref}; use {group::VoxelGroup, view::View}; pub struct VoxelPipeline { + compute_pipeline: wgpu::ComputePipeline, + texture: StorageTexture, + cbind_group_layout: wgpu::BindGroupLayout, + cbind_group: wgpu::BindGroup, + pipeline: wgpu::RenderPipeline, view: Uniform, bind_group_layout: wgpu::BindGroupLayout, @@ -36,27 +45,54 @@ pub struct VoxelPipeline { impl VoxelPipeline { pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self { // shaders - let shader = device.create_shader_module(include_wgsl!("shader.wgsl")); + let shader = device.create_shader_module(include_wgsl!("render.wgsl")); let view = Uniform::init(device, "view", 0); - let voxels = Storage::init(device, "voxels", 1); - let voxel_groups = Storage::init(device, "voxel groups", 2); + let voxels = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxels", 1); + let voxel_groups = Storage::init(device, wgpu::ShaderStages::COMPUTE, "voxel groups", 2); let global_lights = Storage::init_with( device, + wgpu::ShaderStages::COMPUTE, "global lights", 3, &[GlobalLight { direction: Vector3::new(-0.5, -4.0, 2.0).normalize(), }], ); + let texture = StorageTexture::init( + device, + wgpu::Extent3d { + width: 1920, + height: 1080, + depth_or_array_layers: 1, + }, + "idk man im tired", + wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, + 4, + ); // bind groups let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { entries: &[ view.bind_group_layout_entry(), - voxels.bind_group_layout_entry(), - voxel_groups.bind_group_layout_entry(), - global_lights.bind_group_layout_entry(), + wgpu::BindGroupLayoutEntry { + binding: 1, + visibility: wgpu::ShaderStages::FRAGMENT, + ty: wgpu::BindingType::Texture { + multisampled: false, + view_dimension: wgpu::TextureViewDimension::D2, + sample_type: wgpu::TextureSampleType::Float { filterable: true }, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + visibility: wgpu::ShaderStages::FRAGMENT, + // This should match the filterable field of the + // corresponding Texture entry above. + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, ], label: Some("tile_bind_group_layout"), }); @@ -65,9 +101,14 @@ impl VoxelPipeline { layout: &bind_group_layout, entries: &[ view.bind_group_entry(), - voxels.bind_group_entry(), - voxel_groups.bind_group_entry(), - global_lights.bind_group_entry(), + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::TextureView(&texture.view), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: wgpu::BindingResource::Sampler(&texture.sampler), + }, ], label: Some("tile_bind_group"), }); @@ -109,7 +150,7 @@ impl VoxelPipeline { conservative: false, }, depth_stencil: Some(wgpu::DepthStencilState { - format: Texture::DEPTH_FORMAT, + format: DepthTexture::DEPTH_FORMAT, depth_write_enabled: true, depth_compare: wgpu::CompareFunction::Less, stencil: wgpu::StencilState::default(), @@ -124,7 +165,51 @@ impl VoxelPipeline { cache: None, }); + let cbind_group_layout = + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + entries: &[ + view.bind_group_layout_entry(), + voxels.bind_group_layout_entry(), + voxel_groups.bind_group_layout_entry(), + global_lights.bind_group_layout_entry(), + texture.bind_group_layout_entry(), + ], + label: Some("voxel compute"), + }); + + let cbind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: &cbind_group_layout, + entries: &[ + view.bind_group_entry(), + voxels.bind_group_entry(), + voxel_groups.bind_group_entry(), + global_lights.bind_group_entry(), + texture.bind_group_entry(), + ], + label: Some("voxel compute"), + }); + + let cpipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("voxel compute"), + bind_group_layouts: &[&cbind_group_layout], + push_constant_ranges: &[], + }); + + let cs_module = device.create_shader_module(include_wgsl!("compute.wgsl")); + let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("voxel"), + layout: Some(&cpipeline_layout), + module: &cs_module, + entry_point: "main", + compilation_options: Default::default(), + cache: None, + }); + Self { + compute_pipeline, + texture, + cbind_group_layout, + cbind_group, pipeline: render_pipeline, view, bind_group, @@ -196,20 +281,12 @@ impl VoxelPipeline { device: &wgpu::Device, encoder: &mut wgpu::CommandEncoder, belt: &mut wgpu::util::StagingBelt, - AddChunk { - id, - pos, - tree, - .. - }: AddChunk, + AddChunk { id, pos, tree, .. }: AddChunk, ) { let offset = self.voxels.len(); let data = tree.raw(); - let updates = [ArrBufUpdate { - offset, - data, - }]; + let updates = [ArrBufUpdate { offset, data }]; let size = offset + data.len(); self.voxels.update(device, encoder, belt, size, &updates); @@ -232,14 +309,48 @@ impl VoxelPipeline { .update(device, encoder, belt, size, &updates); self.id_map.insert(id, (i, group)); + self.update_bind_group(device); + } - self.bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - layout: &self.bind_group_layout, + pub fn update_bind_group(&mut self, device: &wgpu::Device) { + self.cbind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: &self.cbind_group_layout, entries: &[ self.view.bind_group_entry(), self.voxels.bind_group_entry(), self.voxel_groups.bind_group_entry(), self.global_lights.bind_group_entry(), + self.texture.bind_group_entry(), + ], + label: Some("tile_bind_group"), + }); + } + + pub fn resize(&mut self, device: &wgpu::Device, size: Vector2) { + self.texture = StorageTexture::init( + 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.update_bind_group(device); + self.bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: &self.bind_group_layout, + entries: &[ + self.view.bind_group_entry(), + wgpu::BindGroupEntry { + binding: 1, + resource: wgpu::BindingResource::TextureView(&self.texture.view), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: wgpu::BindingResource::Sampler(&self.texture.sampler), + }, ], label: Some("tile_bind_group"), }); @@ -293,4 +404,10 @@ impl VoxelPipeline { render_pass.set_bind_group(0, &self.bind_group, &[]); render_pass.draw(0..4, 0..1); } + + pub fn compute(&self, pass: &mut wgpu::ComputePass, w: u32, h: u32) { + pass.set_pipeline(&self.compute_pipeline); + pass.set_bind_group(0, &self.cbind_group, &[]); + pass.dispatch_workgroups(w / 16, h / 16, 1); + } } diff --git a/src/client/render/voxel/ray_oct/render.wgsl b/src/client/render/voxel/ray_oct/render.wgsl new file mode 100644 index 0000000..d325d67 --- /dev/null +++ b/src/client/render/voxel/ray_oct/render.wgsl @@ -0,0 +1,47 @@ +// Vertex shader + +struct VertexOutput { + @builtin(position) clip_position: vec4, +}; + +struct View { + transform: mat4x4, + width: u32, + height: u32, + zoom: f32, +}; + +@group(0) @binding(0) +var view: View; +@group(0) @binding(1) +var t_diffuse: texture_2d; +@group(0) @binding(2) +var s_diffuse: sampler; + +@vertex +fn vs_main( + @builtin(vertex_index) vi: u32, + @builtin(instance_index) ii: u32, +) -> VertexOutput { + var out: VertexOutput; + + var pos = vec2( + f32(vi % 2u) * 2.0 - 1.0, + f32(vi / 2u) * 2.0 - 1.0, + ); + out.clip_position = vec4(pos.x, pos.y, 0.0, 1.0); + return out; +} + +// Fragment shader + +@fragment +fn fs_main( + in: VertexOutput, +) -> @location(0) vec4 { + let win_dim = vec2(f32(view.width), f32(view.height)); + var pos = in.clip_position.xy / win_dim; + pos.y = 1.0 - pos.y; + return textureSample(t_diffuse, s_diffuse, pos); +} + diff --git a/src/client/render/voxel/ray_oct/shader.wgsl b/src/client/render/voxel/ray_oct/shader.wgsl index ac878b3..bee2c87 100644 --- a/src/client/render/voxel/ray_oct/shader.wgsl +++ b/src/client/render/voxel/ray_oct/shader.wgsl @@ -91,6 +91,7 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { let dir = (group.transform_inv * dir_view).xyz; let dir_if = sign(dir); + let dir_uf = max(dir_if, vec3(0.0)); @@ -103,118 +104,65 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { var next_normal = vec3(0.0, 0.0, 0.0); // find where ray intersects with group - let plane_point = (vec3(1.0) - dir_if) / 2.0 * dim_f; + let pos_min = (vec3(1.0) - dir_uf) * dim_f; + let pos_max = dir_uf * dim_f; var pos = pos_start; + // time of intersection; x = td + p, solve for t + let t_min = (pos_min - pos) / dir; + let t_max = (pos_max - pos) / dir; var t = 0.0; if outside3f(pos, ZERO3F, dim_f) { - // time of intersection; x = td + p, solve for t - let t_i = (plane_point - pos) / dir; // points of intersection - let px = pos + t_i.x * dir; - let py = pos + t_i.y * dir; - let pz = pos + t_i.z * dir; + let px = pos + t_min.x * dir; + let py = pos + t_min.y * dir; + let pz = pos + t_min.z * dir; // check if point is in bounds let hit = vec3( inside2f(px.yz, ZERO2F, dim_f.yz), inside2f(py.xz, ZERO2F, dim_f.xz), inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_i > ZERO3F); + ) && (t_min > ZERO3F); if !any(hit) { return vec4(0.0); } pos = select(select(pz, py, hit.y), px, hit.x); - t = select(select(t_i.z, t_i.y, hit.y), t_i.x, hit.x); + t = select(select(t_min.z, t_min.y, hit.y), t_min.x, hit.x); next_normal = select(select(normals[2], normals[1], hit.y), normals[0], hit.x); } - // voxel position relative to low_corner - var vox_pos = clamp(vec3(pos), vec3(0), dim_i - vec3(1)); - - - - let dir_i = vec3(dir_if); - let dir_u = ((dir_i + vec3(1)) / 2); - let dir_bits = u32(dir_u.x * 4 + dir_u.y * 2 + dir_u.z); - // time to move 1 unit using dir let inc_t = abs(1.0 / dir); - var side_len = 256; - // "unsigned" minimum cube coords of current tree - var low_corner = vec3(0); - // time of next 1 unit plane hit in each direction - var color = vec4(0.0); - var data_start = 1u; + let dir_i = vec3(dir_if); + let dir_u = vec3((dir_i + vec3(1)) / 2); var i = 0u; - var axis = 0; - for (var safety = 0; safety < 1000; safety += 1) { + var data_start = 1u; + var t_center = (t_max + t_min) / 2.0; + var half_t_span = f32(256 / 2) * inc_t; + for (var safety = 0; safety < 9; safety += 1) { let node = voxels[group.offset + i]; if node >= LEAF_BIT { - // leaf let vcolor = get_color(node & LEAF_MASK); if vcolor.a > 0.0 { - let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); - let ambient = 0.2; - let lighting = max(diffuse, ambient); - let new_color = min(vcolor.xyz * lighting, vec3(1.0)); - color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); - color = vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); - if color.a > .999 { - return color; + // let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); + // let ambient = 0.2; + // let lighting = max(diffuse, ambient); + // let new_color = min(vcolor.xyz * lighting, vec3(1.0)); + // color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); + if vcolor.a > .999 { + return vcolor; } } + return vcolor; + } else { + half_t_span *= 0.5; + let dir_idx = vec3(vec3(t) < t_center); + t_center += half_t_span * (1.0 - vec3(dir_idx * 2)); - // move to next face of cube - let half_len = f32(side_len) / 2.0; - let corner = vec3(low_corner) + vec3(half_len) + dir_if * half_len; - let next_t = inc_t * abs(corner - pos_start); - axis = select(select(2, 1, next_t.y < next_t.z), 0, next_t.x < next_t.y && next_t.x < next_t.z); - t = next_t[axis]; - next_normal = normals[axis]; - pos = pos_start + t * dir; - let old = vox_pos[axis]; - vox_pos = vec3(pos) - low_corner; - vox_pos = clamp(vox_pos, vec3(0), side_len - vec3(1)); - vox_pos[axis] += dir_i[axis]; - } else if inside3i(vox_pos, vec3(0), vec3(side_len - 1)) { - // node + let child_i = vec_to_dir(dir_idx ^ dir_u); let node_pos = data_start + node; - side_len /= 2; - let vcorner = vox_pos / side_len; - vox_pos -= vcorner * side_len; - let j = u32(vcorner.x * 4 + vcorner.y * 2 + vcorner.z); - i = node_pos + j; - data_start = node_pos + 9; - - low_corner += vec3(dir_to_vec(j)) * i32(side_len); - + i = node_pos + child_i; + data_start = node_pos + 8; continue; } - - // idrk what to put here tbh but this prolly works; don't zoom out if max - if side_len == 256 { - let a = f32(safety) / 1000.0; - return vec4(0.0, 0.0, a, 1.0); - } - - // get parent info and reset "pointers" to parent - let parent_info_i = data_start - 1; - let parent_info = voxels[group.offset + parent_info_i]; - let parent_root = parent_info_i - (parent_info >> 3); - let parent_loc = parent_info & 7; - let loc = 8 - (data_start - 1 - i); - // let test = (parent_root + 9 + voxels[group.offset + parent_root + parent_loc] + loc) == i; - i = parent_root + parent_loc; - data_start = parent_root + 9; - - // adjust corner back to parent - let low_corner_adj = vec3(dir_to_vec(loc)) * i32(side_len); - low_corner -= low_corner_adj; - - // update vox pos to be relative to parent - vox_pos += low_corner_adj; - - side_len *= 2; - // return vec4(vec3(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0); - // return vec4(vec3(f32(test)), 1.0); } return vec4(1.0, 0.0, 1.0, 1.0); } @@ -227,26 +175,8 @@ fn dir_to_vec(bits: u32) -> vec3 { return vec3(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u)); } -fn get_voxel(offset: u32, pos_: vec3) -> u32 { - var data_start = 1u; - var i = 0u; - var pos = pos_; - var side_len: u32 = 256; - var safety = 0; - while voxels[offset + i] < LEAF_BIT { - let node_pos = data_start + voxels[offset + i]; - side_len /= 2u; - let corner = pos / side_len; - pos -= corner * side_len; - let j = corner.x * 4 + corner.y * 2 + corner.z; - i = node_pos + j; - data_start = node_pos + 8; - if safety == 10 { - return 10u; - } - safety += 1; - } - return voxels[offset + i] & LEAF_MASK; +fn vec_to_dir(vec: vec3) -> u32 { + return vec.x * 4 + vec.y * 2 + vec.z * 1; } fn get_color(id: u32) -> vec4 { diff --git a/src/client/render/voxel/ray_oct/shader_broken.wgsl b/src/client/render/voxel/ray_oct/shader_broken.wgsl deleted file mode 100644 index 6588f96..0000000 --- a/src/client/render/voxel/ray_oct/shader_broken.wgsl +++ /dev/null @@ -1,277 +0,0 @@ -// Vertex shader - -struct GlobalLight { - dir: vec3, -}; - -struct VertexOutput { - @builtin(position) clip_position: vec4, -}; - -struct View { - transform: mat4x4, - width: u32, - height: u32, - zoom: f32, -}; - -struct VoxelGroup { - transform: mat4x4, - transform_inv: mat4x4, - dimensions: vec3, - offset: u32, -}; - -@group(0) @binding(0) -var view: View; -@group(0) @binding(1) -var voxels: array; -@group(0) @binding(2) -var voxel_groups: array; -@group(0) @binding(3) -var global_lights: array; - -@vertex -fn vs_main( - @builtin(vertex_index) vi: u32, - @builtin(instance_index) ii: u32, -) -> VertexOutput { - var out: VertexOutput; - - var pos = vec2( - f32(vi % 2u) * 2.0 - 1.0, - f32(vi / 2u) * 2.0 - 1.0, - ) ; - out.clip_position = vec4(pos.x, pos.y, 0.0, 1.0); - return out; -} - -// Fragment shader - -@fragment -fn fs_main( - in: VertexOutput, -) -> @location(0) vec4 { - // get position of the pixel; eye at origin, pixel on plane z = 1 - let win_dim = vec2(f32(view.width), f32(view.height)); - let aspect = win_dim.y / win_dim.x; - let pixel_pos = vec3( - (in.clip_position.xy / win_dim - vec2(0.5)) * vec2(2.0, -2.0 * aspect), - 1.0 - ); - - // move to position in world - let pos = view.transform * vec4(pixel_pos, 1.0); - let dir = view.transform * vec4(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 sky_color = light_mult * vec3(1.0, 1.0, 1.0); - color += vec4(sky_color * (1.0 - color.a), 1.0 - color.a); - color.a = 1.0; - return color; -} - -const ZERO3F = vec3(0.0); -const ZERO2F = vec2(0.0); -const DEPTH = 16u; -const FULL_ALPHA = 0.9999; - -fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { - let gi = 0; - let group = voxel_groups[gi]; - if group.dimensions.x == 0 { - return vec4(0.0); - } - let dim_f = vec3(group.dimensions); - let dim_i = vec3(group.dimensions); - - // transform so that group is at 0,0 - let pos_start = (group.transform_inv * pos_view).xyz; - let dir = (group.transform_inv * dir_view).xyz; - - let dir_if = sign(dir); - - - - // calculate normals - var normals = mat3x3( - (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, - ); - var next_normal = vec3(0.0, 0.0, 0.0); - - // find where ray intersects with group - let plane_point = (vec3(1.0) - dir_if) / 2.0 * dim_f; - var pos = pos_start; - var t = 0.0; - if outside3f(pos, ZERO3F, dim_f) { - // time of intersection; x = td + p, solve for t - let t_i = (plane_point - pos) / dir; - // points of intersection - let px = pos + t_i.x * dir; - let py = pos + t_i.y * dir; - let pz = pos + t_i.z * dir; - - // check if point is in bounds - let hit = vec3( - inside2f(px.yz, ZERO2F, dim_f.yz), - inside2f(py.xz, ZERO2F, dim_f.xz), - inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_i > ZERO3F); - if !any(hit) { - return vec4(0.0); - } - pos = select(select(pz, py, hit.y), px, hit.x); - t = select(select(t_i.z, t_i.y, hit.y), t_i.x, hit.x); - next_normal = select(select(normals[2], normals[1], hit.y), normals[0], hit.x); - } - var vox_pos = clamp(vec3(pos), vec3(0), dim_i - vec3(1)); - - - - let dir_i = vec3(dir_if); - let dir_u = ((dir_i + vec3(1)) / 2); - let dir_bits = u32(dir_u.x * 4 + dir_u.y * 2 + dir_u.z); - // time to move 1 unit using dir - let inc_t = abs(1.0 / dir); - var side_len = 256; - // "unsigned" minimum cube coords of current tree - var low_corner = vec3(0); - // time of next 1 unit plane hit in each direction - var color = vec4(0.0); - var data_start = 1u; - var i = 0u; - var axis = 0; - for (var safety = 0; safety < 100; safety += 1) { - let node = voxels[group.offset + i]; - if node >= LEAF_BIT { - let vcolor = get_color(node & LEAF_MASK); - if vcolor.a > 0.0 { - let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); - let ambient = 0.2; - let lighting = max(diffuse, ambient); - let new_color = min(vcolor.xyz * lighting, vec3(1.0)); - color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); - if color.a > .999 { - return color; - } - } - - // move to next face of cube - let half_len = f32(side_len) / 2.0; - let corner = vec3(low_corner) + vec3(half_len) + dir_if * half_len; - let next_t = inc_t * abs(corner - pos_start); - axis = select(select(2, 1, next_t.y < next_t.z), 0, next_t.x < next_t.y && next_t.x < next_t.z); - t = next_t[axis]; - next_normal = normals[axis]; - pos = pos_start + t * dir; - let old = vox_pos[axis]; - vox_pos = vec3(pos) - low_corner; - vox_pos[axis] = old + side_len * dir_i[axis]; - // var axis_c = vec3(0.0); - // axis_c[axis] = 1.0; - // return vec4(axis_c, 1.0); - } else if inside3i(vox_pos, vec3(0), vec3(side_len)) { - let node_pos = data_start + node; - side_len /= 2; - let vcorner = vox_pos / side_len; - vox_pos -= vcorner * side_len; - let j = u32(vcorner.x * 4 + vcorner.y * 2 + vcorner.z); - i = node_pos + j; - data_start = node_pos + 9; - - low_corner += vec3(dir_to_vec(j)) * i32(side_len); - - continue; - } - - // get parent info and reset "pointers" to parent - let parent_info_i = data_start - 1; - if parent_info_i == 0 { - return color; - } - let parent_info = voxels[group.offset + parent_info_i]; - let parent_root = parent_info_i - (parent_info >> 3); - let parent_loc = parent_info & 7; - let loc = 8 - (data_start - 1 - i); - // let test = (parent_root + 9 + voxels[group.offset + parent_root + parent_loc] + loc) == i; - i = parent_root + parent_loc; - data_start = parent_root + 9; - - // adjust corner back to parent - let low_corner_adj = vec3(dir_to_vec(loc)) * i32(side_len); - low_corner -= low_corner_adj; - - // update vox pos to be relative to parent - vox_pos += low_corner_adj; - - side_len *= 2; - // return vec4(vec3(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0); - // return vec4(vec3(f32(test)), 1.0); - } - return color; -} - -const LEAF_BIT = 1u << 31u; -const LEAF_MASK = ~LEAF_BIT; - -// there's no way this is efficient, mod is faster for all I know -fn dir_to_vec(bits: u32) -> vec3 { - return vec3(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u)); -} - -fn get_voxel(offset: u32, pos_: vec3) -> u32 { - var data_start = 1u; - var i = 0u; - var pos = pos_; - var side_len: u32 = 256; - var safety = 0; - while voxels[offset + i] < LEAF_BIT { - let node_pos = data_start + voxels[offset + i]; - side_len /= 2u; - let corner = pos / side_len; - pos -= corner * side_len; - let j = corner.x * 4 + corner.y * 2 + corner.z; - i = node_pos + j; - data_start = node_pos + 8; - if safety == 10 { - return 10u; - } - safety += 1; - } - return voxels[offset + i] & LEAF_MASK; -} - -fn get_color(id: u32) -> vec4 { - switch id { - case 0u: { - return vec4(0.0); - } - case 1u: { - return vec4(0.5, 0.5, 0.5, 1.0); - } - case 2u: { - return vec4(0.5, 1.0, 0.5, 1.0); - } - case 3u: { - return vec4(0.5, 0.5, 1.0, 0.5); - } - default: { - return vec4(1.0, 0.0, 0.0, 1.0); - } - } -} - -fn outside3f(v: vec3, low: vec3, high: vec3) -> bool { - return any(v < low) || any(v > high); -} - -fn inside2f(v: vec2, low: vec2, high: vec2) -> bool { - return all(v >= low) && all(v <= high); -} - -fn inside3i(v: vec3, low: vec3, high: vec3) -> bool { - return all(v >= low) && all(v <= high); -} diff --git a/src/util/oct_tree.rs b/src/util/oct_tree.rs index 95bcae8..b8a3730 100644 --- a/src/util/oct_tree.rs +++ b/src/util/oct_tree.rs @@ -1,25 +1,22 @@ -use std::{collections::VecDeque, fmt::Debug}; +use std::fmt::Debug; use nalgebra::Vector3; use ndarray::ArrayView3; const LEAF_BIT: u32 = 1 << 31; -const DATA_OFFSET: usize = 9; +const DATA_OFFSET: usize = 8; #[repr(C)] #[derive(Debug, Clone, Copy, PartialEq, bytemuck::Pod, bytemuck::Zeroable)] pub struct OctNode(u32); impl OctNode { - pub fn new_node(addr: u32) -> Self { + pub const fn new_node(addr: u32) -> Self { Self(addr) } - pub fn new_leaf(data: u32) -> Self { + pub const fn new_leaf(data: u32) -> Self { Self(data | LEAF_BIT) } - pub fn new_parent(offset: u32, corner: u32) -> Self { - Self((offset << 3) + corner) - } - pub fn is_leaf(&self) -> bool { + pub const fn is_leaf(&self) -> bool { self.0 >= LEAF_BIT } pub fn is_node(&self) -> bool { @@ -69,10 +66,7 @@ impl OctTree { ) -> Self { let mut data = Vec::new(); data.push(OctNode::new_node(0)); - // #######N P SSSSSSSS P - // --------------------| 17 - // -------| 7 - Self::from_fn_offset_inner(f, &mut data, levels, offset, OctNode::new_parent(17, 7)); + Self::from_fn_offset_inner(f, &mut data, levels, offset); if data.len() == 2 { data.remove(0); } @@ -87,7 +81,6 @@ impl OctTree { accumulator: &mut Vec, level: u32, offset: Vector3, - parent: OctNode, ) { if level == 0 { accumulator.push(OctNode::new_leaf(f(offset))); @@ -95,27 +88,23 @@ impl OctTree { } else if level == 1 { let leaves: [OctNode; 8] = core::array::from_fn(|i| OctNode::new_leaf(f(offset + CORNERS[i]))); - if leaves.iter().all(|l| *l == leaves[0]) { + if leaves[1..].iter().all(|l| *l == leaves[0]) { accumulator.push(leaves[0]); } else { accumulator.extend_from_slice(&leaves); - accumulator.push(parent); } return; } let i = accumulator.len(); accumulator.resize(i + 8, OctNode::new_node(0)); - accumulator.push(parent); let mut data_start = 0; for (j, corner_offset) in CORNERS.iter().enumerate() { let sub_start = accumulator.len(); - let sub_parent_offset = 9 + data_start + 8; Self::from_fn_offset_inner( f, accumulator, level - 1, offset + corner_offset * 2usize.pow(level - 1), - OctNode::new_parent(sub_parent_offset as u32, j as u32), ); let len = accumulator.len() - sub_start; if len == 1 { @@ -128,12 +117,13 @@ impl OctTree { } if data_start == 0 { let first = accumulator[i]; - if accumulator[i..i + 8].iter().all(|l| *l == first) { + if accumulator[i + 1..i + 8].iter().all(|l| *l == first) { accumulator.truncate(i); - accumulator.push(first) + accumulator.push(first); } } } + pub fn from_arr(arr: ArrayView3, levels: u32) -> Self { Self::from_fn(&mut |p| arr[(p.x, p.y, p.z)], levels) } @@ -155,8 +145,7 @@ impl OctTree { pub fn raw(&self) -> &[OctNode] { &self.data } - pub fn mesh(&self) { - } + pub fn mesh(&self) {} } pub struct OctTreeIter<'a> {