shadowcat/voxel-game
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

switched to stack and compute shaders

Browse Source
This commit is contained in:
shadow cat
2024-09-11 01:54:00 -04:00
parent 3417250544
commit b6db483189
11 changed files with 373 additions and 511 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/client/handle_input.rs
View File

@@ -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;
}

16
src/client/render/mod.rs
View File

@@ -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,

9
src/client/render/util/storage.rs
View File

@@ -4,20 +4,22 @@ 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, 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<T: bytemuck::Pod> Storage<T> {
BufferUsages::STORAGE,
data
),
visibility,
binding,
}
}
@@ -34,7 +37,7 @@ impl<T: bytemuck::Pod> Storage<T> {
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,

86
src/client/render/util/texture.rs
View File

@@ -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),
};
}
}

2
src/client/render/util/uniform.rs
View File

@@ -37,7 +37,7 @@ impl<T: PartialEq + bytemuck::Pod> Uniform<T> {
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,

101
src/client/render/voxel/ray_oct/shader_old.wgsl → src/client/render/voxel/ray_oct/compute.wgsl
View File

@@ -1,13 +1,18 @@
// Vertex shader
@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 VertexOutput {
@builtin(position) clip_position: vec4<f32>,
};
struct View {
transform: mat4x4<f32>,
width: u32,
@@ -22,45 +27,16 @@ struct VoxelGroup {
offset: u32,
};
@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>;
@vertex
fn vs_main(
@builtin(vertex_index) vi: u32,
@builtin(instance_index) ii: u32,
) -> VertexOutput {
var out: VertexOutput;
var pos = vec2<f32>(
f32(vi % 2u) * 2.0 - 1.0,
f32(vi / 2u) * 2.0 - 1.0,
) ;
out.clip_position = vec4<f32>(pos.x, pos.y, 0.0, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(
in: VertexOutput,
) -> @location(0) vec4<f32> {
@compute
@workgroup_size(16, 16, 1)
fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
// get position of the pixel; eye at origin, pixel on plane z = 1
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
let aspect = win_dim.y / win_dim.x;
let pixel_pos = vec3<f32>(
(in.clip_position.xy / win_dim - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
(vec2<f32>(cell.xy) / win_dim - vec2<f32>(0.5)) * vec2<f32>(2.0, 2.0 * aspect),
view.zoom
);
// move to position in world
let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
let dir = view.transform * vec4<f32>(normalize(pixel_pos), 0.0);
@@ -69,7 +45,7 @@ fn fs_main(
let sky_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
color.a = 1.0;
return color;
textureStore(output, cell.xy, color);
}
const ZERO3F = vec3<f32>(0.0);
@@ -145,20 +121,22 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
var data_start = 1u;
var i = 0u;
var axis = 0;
var hits = 0;
var parents = array<u32, 8>();
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<f32>(1.0));
color += vec4<f32>(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a);
if color.a > .999 {
// return vec4<f32>(f32(safety) / 1000.0, 0.0, 0.0, 1.0);
return color;
}
}
@@ -177,33 +155,34 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
vox_pos[axis] += dir_i[axis];
} else if inside3i(vox_pos, vec3<i32>(0), vec3<i32>(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<i32>(dir_to_vec(j)) * i32(side_len);
vox_pos -= vcorner * side_len;
low_corner += vec3<i32>(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>(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<i32>(dir_to_vec(loc)) * i32(side_len);
@@ -213,8 +192,6 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
vox_pos += low_corner_adj;
side_len *= 2;
// return vec4<f32>(vec3<f32>(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0);
// return vec4<f32>(vec3<f32>(f32(test)), 1.0);
}
return vec4<f32>(1.0, 0.0, 1.0, 1.0);
}
@@ -258,7 +235,7 @@ fn get_color(id: u32) -> vec4<f32> {
return vec4<f32>(0.5, 0.5, 0.5, 1.0);
}
case 2u: {
return vec4<f32>(0.5, 1.0, 0.5, 1.0);
return vec4<f32>(0.8, 0.2, 0.2, 1.0);
}
case 3u: {
return vec4<f32>(0.5, 0.5, 1.0, 0.5);

173
src/client/render/voxel/ray_oct/mod.rs
View File

@@ -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<View>,
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<u32>) {
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);
}
}

47
src/client/render/voxel/ray_oct/render.wgsl Normal file
View File

@@ -0,0 +1,47 @@
// Vertex shader
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
};
struct View {
transform: mat4x4<f32>,
width: u32,
height: u32,
zoom: f32,
};
@group(0) @binding(0)
var<uniform> view: View;
@group(0) @binding(1)
var t_diffuse: texture_2d<f32>;
@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>(
f32(vi % 2u) * 2.0 - 1.0,
f32(vi / 2u) * 2.0 - 1.0,
);
out.clip_position = vec4<f32>(pos.x, pos.y, 0.0, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(
in: VertexOutput,
) -> @location(0) vec4<f32> {
let win_dim = vec2<f32>(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);
}

138
src/client/render/voxel/ray_oct/shader.wgsl
View File

@@ -91,6 +91,7 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
let dir = (group.transform_inv * dir_view).xyz;
let dir_if = sign(dir);
let dir_uf = max(dir_if, vec3<f32>(0.0));
@@ -103,118 +104,65 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
var next_normal = vec3<f32>(0.0, 0.0, 0.0);
// find where ray intersects with group
let plane_point = (vec3<f32>(1.0) - dir_if) / 2.0 * dim_f;
let pos_min = (vec3<f32>(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<bool>(
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<f32>(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<i32>(pos), vec3<i32>(0), dim_i - vec3<i32>(1));
let dir_i = vec3<i32>(dir_if);
let dir_u = ((dir_i + vec3<i32>(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<i32>(0);
// time of next 1 unit plane hit in each direction
var color = vec4<f32>(0.0);
var data_start = 1u;
let dir_i = vec3<i32>(dir_if);
let dir_u = vec3<u32>((dir_i + vec3<i32>(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<f32>(1.0));
color += vec4<f32>(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a);
color = vec4<f32>(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<f32>(1.0));
// color += vec4<f32>(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<u32>(vec3<f32>(t) < t_center);
t_center += half_t_span * (1.0 - vec3<f32>(dir_idx * 2));
// move to next face of cube
let half_len = f32(side_len) / 2.0;
let corner = vec3<f32>(low_corner) + vec3<f32>(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<i32>(pos) - low_corner;
vox_pos = clamp(vox_pos, vec3<i32>(0), side_len - vec3<i32>(1));
vox_pos[axis] += dir_i[axis];
} else if inside3i(vox_pos, vec3<i32>(0), vec3<i32>(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<i32>(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<f32>(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<i32>(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<f32>(vec3<f32>(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0);
// return vec4<f32>(vec3<f32>(f32(test)), 1.0);
}
return vec4<f32>(1.0, 0.0, 1.0, 1.0);
}
@@ -227,26 +175,8 @@ fn dir_to_vec(bits: u32) -> vec3<u32> {
return vec3<u32>(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u));
}
fn get_voxel(offset: u32, pos_: vec3<u32>) -> 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>) -> u32 {
return vec.x * 4 + vec.y * 2 + vec.z * 1;
}
fn get_color(id: u32) -> vec4<f32> {

277
src/client/render/voxel/ray_oct/shader_broken.wgsl
View File

@@ -1,277 +0,0 @@
// Vertex shader
struct GlobalLight {
dir: vec3<f32>,
};
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
};
struct View {
transform: mat4x4<f32>,
width: u32,
height: u32,
zoom: f32,
};
struct VoxelGroup {
transform: mat4x4<f32>,
transform_inv: mat4x4<f32>,
dimensions: vec3<u32>,
offset: u32,
};
@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>;
@vertex
fn vs_main(
@builtin(vertex_index) vi: u32,
@builtin(instance_index) ii: u32,
) -> VertexOutput {
var out: VertexOutput;
var pos = vec2<f32>(
f32(vi % 2u) * 2.0 - 1.0,
f32(vi / 2u) * 2.0 - 1.0,
) ;
out.clip_position = vec4<f32>(pos.x, pos.y, 0.0, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(
in: VertexOutput,
) -> @location(0) vec4<f32> {
// get position of the pixel; eye at origin, pixel on plane z = 1
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
let aspect = win_dim.y / win_dim.x;
let pixel_pos = vec3<f32>(
(in.clip_position.xy / win_dim - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
1.0
);
// move to position in world
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 sky_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
color.a = 1.0;
return color;
}
const ZERO3F = vec3<f32>(0.0);
const ZERO2F = vec2<f32>(0.0);
const DEPTH = 16u;
const FULL_ALPHA = 0.9999;
fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
let gi = 0;
let group = voxel_groups[gi];
if group.dimensions.x == 0 {
return vec4<f32>(0.0);
}
let dim_f = vec3<f32>(group.dimensions);
let dim_i = vec3<i32>(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<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 next_normal = vec3<f32>(0.0, 0.0, 0.0);
// find where ray intersects with group
let plane_point = (vec3<f32>(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<bool>(
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<f32>(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<i32>(pos), vec3<i32>(0), dim_i - vec3<i32>(1));
let dir_i = vec3<i32>(dir_if);
let dir_u = ((dir_i + vec3<i32>(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<i32>(0);
// time of next 1 unit plane hit in each direction
var color = vec4<f32>(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<f32>(1.0));
color += vec4<f32>(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<f32>(low_corner) + vec3<f32>(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<i32>(pos) - low_corner;
vox_pos[axis] = old + side_len * dir_i[axis];
// var axis_c = vec3<f32>(0.0);
// axis_c[axis] = 1.0;
// return vec4<f32>(axis_c, 1.0);
} else if inside3i(vox_pos, vec3<i32>(0), vec3<i32>(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<i32>(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<i32>(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<f32>(vec3<f32>(dir_to_vec(parent_loc)) * f32(loc) / 8.0, 1.0);
// return vec4<f32>(vec3<f32>(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<u32> {
return vec3<u32>(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u));
}
fn get_voxel(offset: u32, pos_: vec3<u32>) -> 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<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);
}

33
src/util/oct_tree.rs
View File

@@ -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<OctNode>,
level: u32,
offset: Vector3<usize>,
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<u32>, 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> {