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

switch to noise lib that works on arm neon (apple)

Browse Source
This commit is contained in:
shadow cat
2024-09-20 22:07:40 -04:00
parent cb422ff5dd
commit e66c75ca44
24 changed files with 2313 additions and 114 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
src/client/mod.rs
View File

@@ -3,6 +3,7 @@ mod camera;
mod component;
mod handle_input;
mod input;
// mod render_vulkan;
pub mod render;
mod rsc;
mod state;
@@ -66,7 +67,7 @@ impl Client<'_> {
.expect("Failed to create window"),
);
let renderer = Renderer::spawn(window.clone());
let renderer = Renderer::new(window.clone());
world.insert_resource(RenderCommands(Vec::new()));
let state = ClientState::new();

14
src/client/render/command.rs
View File

@@ -43,20 +43,6 @@ pub struct UpdateGridTransform {
}
impl<'a> Renderer<'a> {
pub fn spawn(window: Arc<Window>) -> Renderer<'a> {
let size = window.inner_size();
let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
let surface = instance
.create_surface(window)
.expect("Could not create window surface!");
Self::new(instance, surface, size)
}
pub fn handle_commands(&mut self, commands: Vec<RenderCommand>) {
let mut new_camera = false;
for cmd in commands {

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

@@ -1,14 +1,16 @@
mod command;
mod util;
pub mod voxel;
use std::sync::Arc;
pub use command::*;
use super::camera::Camera;
use crate::{client::rsc::CLEAR_COLOR, util::timer::Timer};
use crate::client::rsc::CLEAR_COLOR;
use nalgebra::Vector2;
use util::DepthTexture;
use voxel::VoxelPipeline;
use winit::dpi::PhysicalSize;
use winit::{dpi::PhysicalSize, window::Window};
pub struct Renderer<'a> {
size: Vector2<u32>,
@@ -25,10 +27,19 @@ pub struct Renderer<'a> {
impl<'a> Renderer<'a> {
pub fn new(
instance: wgpu::Instance,
surface: wgpu::Surface<'a>,
size: PhysicalSize<u32>,
window: Arc<Window>,
) -> Self {
let size = window.inner_size();
let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
backends: wgpu::Backends::PRIMARY,
..Default::default()
});
let surface = instance
.create_surface(window)
.expect("Could not create window surface!");
let adapter = pollster::block_on(instance.request_adapter(&wgpu::RequestAdapterOptions {
power_preference: wgpu::PowerPreference::default(),
compatible_surface: Some(&surface),
@@ -36,11 +47,16 @@ impl<'a> Renderer<'a> {
}))
.expect("Could not get adapter!");
let buf_size = (10u32.pow(9) * 15) / 10;
let (device, queue) = pollster::block_on(adapter.request_device(
&wgpu::DeviceDescriptor {
label: None,
required_features: wgpu::Features::empty(),
required_limits: wgpu::Limits::default(),
required_limits: wgpu::Limits {
max_storage_buffer_binding_size: buf_size,
max_buffer_size: buf_size as u64,
..Default::default()
},
memory_hints: wgpu::MemoryHints::default(),
},
None, // Trace path
@@ -160,8 +176,7 @@ impl<'a> Renderer<'a> {
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.depth_texture = DepthTexture::init(&self.device, &self.config, "depth_texture");
self.voxel_pipeline.update_view(
&self.device,
&mut self.encoder,

53
src/client/render/voxel/ray/voxel/color.rs Normal file
View File

@@ -0,0 +1,53 @@
use rand::distributions::{Distribution, Standard};
#[repr(C)]
#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Zeroable)]
pub struct VoxelColor {
pub r: u8,
pub g: u8,
pub b: u8,
pub a: u8,
}
unsafe impl bytemuck::Pod for VoxelColor {}
impl VoxelColor {
pub fn none() -> Self {
Self {
r: 0,
g: 0,
b: 0,
a: 0,
}
}
pub fn black() -> Self {
Self {
r: 0,
g: 0,
b: 0,
a: 255,
}
}
pub fn white() -> Self {
Self {
r: 255,
g: 255,
b: 255,
a: 255,
}
}
pub fn random() -> Self {
rand::random()
}
}
impl Distribution<VoxelColor> for Standard {
fn sample<R: rand::prelude::Rng + ?Sized>(&self, rng: &mut R) -> VoxelColor {
VoxelColor {
r: rng.gen(),
g: rng.gen(),
b: rng.gen(),
a: rng.gen(),
}
}
}

24
src/client/render/voxel/ray/voxel/grid.rs Normal file
View File

@@ -0,0 +1,24 @@
use nalgebra::Matrix4x3;
// this has cost me more than a couple of hours trying to figure out alignment :skull:
// putting transform at the beginning so I don't have to deal with its alignment
// I should probably look into encase (crate)
#[repr(C, align(16))]
#[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
pub struct GridInfo {
pub transform: Matrix4x3<f32>,
pub width: u32,
pub height: u32,
}
unsafe impl bytemuck::Pod for GridInfo {}
impl Default for GridInfo {
fn default() -> Self {
Self {
transform: Matrix4x3::identity(),
width: 0,
height: 0,
}
}
}

12
src/client/render/voxel/ray/voxel/group.rs Normal file
View File

@@ -0,0 +1,12 @@
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 {}

198
src/client/render/voxel/ray/voxel/layout.rs Normal file
View File

@@ -0,0 +1,198 @@
use wgpu::TextureFormat;
use super::{group::VoxelGroup, light::GlobalLight, view::View};
use crate::{
client::render::util::{Storage, StorageTexture, Uniform},
util::oct_tree::OctNode,
};
use nalgebra::Vector3;
pub struct Layout {
pub texture: StorageTexture,
pub view: Uniform<View>,
pub voxel_groups: Storage<VoxelGroup>,
pub voxels: Storage<OctNode>,
pub global_lights: Storage<GlobalLight>,
render_bind_layout: wgpu::BindGroupLayout,
compute_bind_layout: wgpu::BindGroupLayout,
render_pipeline_layout: wgpu::PipelineLayout,
compute_pipeline_layout: wgpu::PipelineLayout,
format: TextureFormat,
}
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 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(-1.0, -2.3, 2.0).normalize(),
}],
);
let texture = StorageTexture::init(
device,
wgpu::Extent3d {
width: config.width,
height: config.height,
depth_or_array_layers: 1,
},
"compute output",
wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT,
4,
);
let render_bind_layout =
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
entries: &[
wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::VERTEX,
ty: wgpu::BindingType::Texture {
multisampled: false,
view_dimension: wgpu::TextureViewDimension::D2,
sample_type: wgpu::TextureSampleType::Float { filterable: true },
},
count: None,
},
wgpu::BindGroupLayoutEntry {
binding: 1,
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("voxel render"),
});
let compute_bind_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 render_pipeline_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("Tile Pipeline Layout"),
bind_group_layouts: &[&render_bind_layout],
push_constant_ranges: &[],
});
let compute_pipeline_layout =
device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("voxel compute"),
bind_group_layouts: &[&compute_bind_layout],
push_constant_ranges: &[],
});
Self {
view,
voxels,
voxel_groups,
global_lights,
texture,
render_bind_layout,
compute_bind_layout,
render_pipeline_layout,
compute_pipeline_layout,
format: config.format,
}
}
pub fn render_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
device.create_bind_group(&wgpu::BindGroupDescriptor {
layout: &self.render_bind_layout,
entries: &[
wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureView(&self.texture.view),
},
wgpu::BindGroupEntry {
binding: 1,
resource: wgpu::BindingResource::Sampler(&self.texture.sampler),
},
],
label: Some("tile_bind_group"),
})
}
pub fn compute_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
device.create_bind_group(&wgpu::BindGroupDescriptor {
layout: &self.compute_bind_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("voxel compute"),
})
}
pub fn render_pipeline(
&self,
device: &wgpu::Device,
shader: wgpu::ShaderModule,
) -> wgpu::RenderPipeline {
device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
label: Some("Voxel Pipeline"),
layout: Some(&self.render_pipeline_layout),
vertex: wgpu::VertexState {
module: &shader,
entry_point: "vs_main",
buffers: &[],
compilation_options: wgpu::PipelineCompilationOptions::default(),
},
fragment: Some(wgpu::FragmentState {
module: &shader,
entry_point: "fs_main",
targets: &[Some(wgpu::ColorTargetState {
format: self.format,
blend: Some(wgpu::BlendState::REPLACE),
write_mask: wgpu::ColorWrites::ALL,
})],
compilation_options: wgpu::PipelineCompilationOptions::default(),
}),
primitive: wgpu::PrimitiveState {
topology: wgpu::PrimitiveTopology::TriangleStrip,
strip_index_format: None,
front_face: wgpu::FrontFace::Ccw,
cull_mode: None,
polygon_mode: wgpu::PolygonMode::Fill,
unclipped_depth: false,
conservative: false,
},
depth_stencil: None,
multisample: wgpu::MultisampleState {
count: 1,
mask: !0,
alpha_to_coverage_enabled: true,
},
multiview: None,
cache: None,
})
}
pub fn compute_pipeline(
&self,
device: &wgpu::Device,
shader: &wgpu::ShaderModule,
) -> wgpu::ComputePipeline {
device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: Some("voxel"),
layout: Some(&self.compute_pipeline_layout),
module: shader,
entry_point: "main",
compilation_options: Default::default(),
cache: None,
})
}
}

9
src/client/render/voxel/ray/voxel/light.rs Normal file
View File

@@ -0,0 +1,9 @@
use nalgebra::Vector3;
#[repr(C, align(16))]
#[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
pub struct GlobalLight {
pub direction: Vector3<f32>,
}
unsafe impl bytemuck::Pod for GlobalLight {}

220
src/client/render/voxel/ray/voxel/mod.rs Normal file
View File

@@ -0,0 +1,220 @@
mod color;
mod grid;
mod group;
mod layout;
mod light;
mod view;
use super::super::UpdateGridTransform;
use crate::{
client::{
camera::Camera,
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 std::{collections::HashMap, ops::Deref};
use wgpu::include_wgsl;
use {group::VoxelGroup, view::View};
pub struct VoxelPipeline {
layout: Layout,
compute_pipeline: wgpu::ComputePipeline,
compute_bind_group: wgpu::BindGroup,
render_pipeline: wgpu::RenderPipeline,
render_bind_group: wgpu::BindGroup,
id_map: HashMap<Entity, (usize, VoxelGroup)>,
}
const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/render.wgsl");
const COMPUTE_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/compute.wgsl");
impl VoxelPipeline {
pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self {
// shaders
let layout = Layout::init(device, config);
let render_bind_group = layout.render_bind_group(device);
let shader = device.create_shader_module(RENDER_SHADER);
let render_pipeline = layout.render_pipeline(device, shader);
let compute_bind_group = layout.compute_bind_group(device);
let shader = device.create_shader_module(COMPUTE_SHADER);
let compute_pipeline = layout.compute_pipeline(device, &shader);
Self {
layout,
compute_pipeline,
compute_bind_group,
render_pipeline,
render_bind_group,
id_map: HashMap::new(),
}
}
pub fn reset_shader(&mut self, device: &wgpu::Device) {
let shader = device.create_shader_module(COMPUTE_SHADER);
self.compute_pipeline = self.layout.compute_pipeline(device, &shader);
}
pub fn add_group(
&mut self,
device: &wgpu::Device,
encoder: &mut wgpu::CommandEncoder,
belt: &mut wgpu::util::StagingBelt,
CreateVoxelGrid {
id,
pos,
orientation,
dimensions,
grid,
}: CreateVoxelGrid,
) {
}
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",
) else {
println!("Failed to reload shader!");
return;
};
device.push_error_scope(wgpu::ErrorFilter::Validation);
let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Shader"),
source: wgpu::ShaderSource::Wgsl(shader.into()),
});
if pollster::block_on(device.pop_error_scope()).is_some() {
let comp_info = pollster::block_on(shader.get_compilation_info());
println!("Failed to compile shaders:");
for msg in comp_info.messages {
println!("{}", msg.message);
}
} else {
self.compute_pipeline = self.layout.compute_pipeline(device, &shader);
}
}
pub fn add_chunk(
&mut self,
device: &wgpu::Device,
encoder: &mut wgpu::CommandEncoder,
belt: &mut wgpu::util::StagingBelt,
AddChunk { id, pos, tree, .. }: AddChunk,
) {
let offset = self.layout.voxels.len();
let data = tree.raw();
let updates = [ArrBufUpdate { offset, data }];
let size = offset + data.len();
self.layout
.voxels
.update(device, encoder, belt, size, &updates);
let proj = Projective3::identity()
* Translation3::from((pos.deref() * chunk::SIDE_LENGTH as i32).cast())
* Translation3::from(-chunk::DIMENSIONS.cast() / 2.0);
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.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(
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);
}
pub fn update_transform(
&mut self,
device: &wgpu::Device,
encoder: &mut wgpu::CommandEncoder,
belt: &mut wgpu::util::StagingBelt,
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);
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);
}
}
pub fn update_view(
&mut self,
device: &wgpu::Device,
encoder: &mut wgpu::CommandEncoder,
belt: &mut wgpu::util::StagingBelt,
size: Vector2<u32>,
camera: &Camera,
) {
let transform =
Transform3::identity() * Translation3::from(camera.pos) * camera.orientation;
let data = View {
zoom: camera.scale,
transform,
};
self.layout.view.update(device, encoder, belt, data);
}
pub fn draw<'a>(&'a self, render_pass: &mut wgpu::RenderPass<'a>) {
render_pass.set_pipeline(&self.render_pipeline);
render_pass.set_bind_group(0, &self.render_bind_group, &[]);
render_pass.draw(0..4, 0..1);
}
pub const WORKGROUP_SIZE: u32 = 8;
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 x = (buf.width() - 1) / Self::WORKGROUP_SIZE + 1;
let y = (buf.height() - 1) / Self::WORKGROUP_SIZE + 1;
pass.dispatch_workgroups(x, y, 1);
}
}

282
src/client/render/voxel/ray/voxel/shader/compute.wgsl Normal file
View File

@@ -0,0 +1,282 @@
@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 = 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 dir = view.transform * vec4<f32>(normalize(vec3<f32>(pixel_pos, view.zoom)), 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 = 10000;
// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
// numbers are stored and the bit manipulation used
const MAX_SCALE: u32 = 16;
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;
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_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>();
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 * 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 != prev {
if node != LEAF_BIT {
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 * 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 = node;
}
}
// let fog = min(t / t_mult / 1000.0, 1.0);
// 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 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, 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: {
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: {
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: {
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);
}
}
}
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);
}
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);
}

419
src/client/render/voxel/ray/voxel/shader/compute_shadow.wgsl Normal file
View File

@@ -0,0 +1,419 @@
@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);
}

364
src/client/render/voxel/ray/voxel/shader/compute_struct.wgsl Normal file
View File

@@ -0,0 +1,364 @@
@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);
let start = start_ray(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);
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;
struct Ray {
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();
}
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;
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) / dir;
if outside3f(pos, ZERO3F, dim_f) {
// points of intersection
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_min > ZERO3F);
if !any(hit) {
return RayStart();
}
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;
let t_offset = max(max(t_min.x, t_min.y), t_min.z);
let 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;
let node_start = 1u;
let scale = MAX_SCALE - 1;
let scale_exp2 = 0.5;
let parents = array<u32, MAX_SCALE>();
var child = 0u;
var vox_pos = vec3<f32>(1.0);
let t_center = t_min + scale_exp2 * t_inc;
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;
return RayStart(
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;
let node = voxels[group_offset + node_start + (child ^ inv_dir_bits)];
if node >= LEAF_BIT {
if node != skip {
data = node;
break;
}
// move to next time point and determine which axis to move along
let t_next = t_corner + scale_exp2 * t_inc;
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 * t_inc;
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 RayResult(
Ray(
t,
vox_pos,
t_inc,
scale,
min_adj,
child,
axis,
node_start,
group_offset,
inv_dir_bits,
parents,
),
data
);
}
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);
}

38
src/client/render/voxel/ray/voxel/shader/render.wgsl Normal file
View File

@@ -0,0 +1,38 @@
// Vertex shader
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
@location(0) tex_pos: vec2<f32>,
};
@group(0) @binding(0)
var tex: texture_2d<f32>;
@group(0) @binding(1)
var sample: sampler;
@vertex
fn vs_main(
@builtin(vertex_index) vi: u32,
@builtin(instance_index) ii: u32,
) -> VertexOutput {
var out: VertexOutput;
let pos = vec2<f32>(
f32(vi % 2u),
f32(vi / 2u),
);
out.clip_position = vec4<f32>(pos * 2.0 - 1.0, 0.0, 1.0);
out.tex_pos = pos;
out.tex_pos.y = 1.0 - out.tex_pos.y;
return out;
}
// Fragment shader
@fragment
fn fs_main(
in: VertexOutput,
) -> @location(0) vec4<f32> {
return textureSample(tex, sample, in.tex_pos);
}

19
src/client/render/voxel/ray/voxel/view.rs Normal file
View File

@@ -0,0 +1,19 @@
use nalgebra::Transform3;
#[repr(C, align(16))]
#[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
pub struct View {
pub transform: Transform3<f32>,
pub zoom: f32,
}
unsafe impl bytemuck::Pod for View {}
impl Default for View {
fn default() -> Self {
Self {
zoom: 1.0,
transform: Transform3::identity(),
}
}
}

8
src/client/render/voxel/ray_oct/shader/compute.wgsl
View File

@@ -58,7 +58,7 @@ 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 = 13;
@@ -166,12 +166,8 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
}
var pos = (pos_view + dir_view * real_t).xyz;
pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]);
// if true {return vec4<f32>(floor(pos) / 16.0, 1.0);}
// let pos = (vox_pos - 1.5) * (dir_if) + 0.5 - scale_exp2 * (1.0 - dir_uf);
// let pos = t / t_mult;
// if true {return vec4<f32>(pos, 1.0);}
let vcolor = get_color(node & LEAF_MASK, pos);
let normal = normals[axis];
var normal = normals[axis];
let light_color = vec3<f32>(1.0);
let light_dir = global_lights[0].dir;

82
src/client/render_vulkan/command.rs Normal file
View File

@@ -0,0 +1,82 @@
use crate::{
client::camera::Camera,
common::component::{ChunkMesh, ChunkPos}, util::oct_tree::OctTree,
};
use super::{voxel::VoxelColor, Renderer};
use bevy_ecs::entity::Entity;
use nalgebra::{Rotation3, Vector3};
use ndarray::Array3;
#[derive(Debug, Clone)]
pub enum RenderCommand {
CreateVoxelGrid(CreateVoxelGrid),
AddChunk(AddChunk),
UpdateGridTransform(UpdateGridTransform),
ViewUpdate(Camera),
}
#[derive(Debug, Clone)]
pub struct CreateVoxelGrid {
pub id: Entity,
pub pos: Vector3<f32>,
pub orientation: Rotation3<f32>,
pub dimensions: Vector3<usize>,
pub grid: Array3<VoxelColor>,
}
#[derive(Debug, Clone)]
pub struct AddChunk {
pub id: Entity,
pub pos: ChunkPos,
pub mesh: ChunkMesh,
pub tree: OctTree,
}
#[derive(Debug, Clone)]
pub struct UpdateGridTransform {
pub id: Entity,
pub pos: Vector3<f32>,
pub orientation: Rotation3<f32>,
}
impl Renderer {
pub fn handle_commands(&mut self, commands: Vec<RenderCommand>) {
let mut new_camera = false;
for cmd in commands {
match cmd {
RenderCommand::CreateVoxelGrid(desc) => self.voxel_pipeline.add_group(
&self.device,
&mut self.encoder,
&mut self.staging_belt,
desc,
),
RenderCommand::ViewUpdate(camera) => {
new_camera = true;
self.camera = camera;
}
RenderCommand::UpdateGridTransform(update) => self.voxel_pipeline.update_transform(
&self.device,
&mut self.encoder,
&mut self.staging_belt,
update,
),
RenderCommand::AddChunk(desc) => self.voxel_pipeline.add_chunk(
&self.device,
&mut self.encoder,
&mut self.staging_belt,
desc,
),
}
}
if new_camera {
self.voxel_pipeline.update_view(
&self.device,
&mut self.encoder,
&mut self.staging_belt,
self.size,
&self.camera,
);
}
}
}

79
src/client/render_vulkan/mod.rs Normal file
View File

@@ -0,0 +1,79 @@
mod command;
pub mod voxel;
use std::sync::Arc;
pub use command::*;
use vulkano::{
device::{Device, DeviceCreateInfo, QueueCreateInfo, QueueFlags}, instance::{Instance, InstanceCreateInfo}, memory::allocator::StandardMemoryAllocator, VulkanLibrary
};
use super::camera::Camera;
use crate::client::rsc::CLEAR_COLOR;
use nalgebra::Vector2;
use voxel::VoxelPipeline;
use winit::{dpi::PhysicalSize, window::Window};
pub struct Renderer {
camera: Camera,
}
impl Renderer {
pub fn new(window: Arc<Window>) -> Self {
let library = VulkanLibrary::new().expect("no local Vulkan library/DLL");
let instance = Instance::new(library, InstanceCreateInfo::default())
.expect("failed to create instance");
let physical_device = instance
.enumerate_physical_devices()
.expect("could not enumerate devices")
.next()
.expect("no devices available");
let queue_family_index = physical_device
.queue_family_properties()
.iter()
.enumerate()
.position(|(_queue_family_index, queue_family_properties)| {
queue_family_properties
.queue_flags
.contains(QueueFlags::GRAPHICS)
})
.expect("couldn't find a graphical queue family")
as u32;
let (device, mut queues) = Device::new(
physical_device,
DeviceCreateInfo {
// here we pass the desired queue family to use by index
queue_create_infos: vec![QueueCreateInfo {
queue_family_index,
..Default::default()
}],
..Default::default()
},
)
.expect("failed to create device");
let queue = queues.next().unwrap();
let memory_allocator = Arc::new(StandardMemoryAllocator::new_default(device.clone()));
Self {
camera: Camera::default(),
size: Vector2::new(size.width, size.height),
voxel_pipeline: VoxelPipeline::new(&device, &config),
}
}
pub fn reset_shader(&mut self) {
todo!()
}
pub fn update_shader(&mut self) {
todo!()
}
pub fn draw(&mut self) {}
pub fn resize(&mut self, size: PhysicalSize<u32>) {
self.size = Vector2::new(size.width, size.height);
todo!();
}
}

25
src/client/render_vulkan/voxel/mod.rs Normal file
View File

@@ -0,0 +1,25 @@
use std::sync::Arc;
use vulkano::{buffer::{Buffer, BufferCreateInfo, BufferUsage}, memory::allocator::{AllocationCreateInfo, MemoryAllocator, MemoryTypeFilter}};
pub struct VoxelPipeline {
}
impl VoxelPipeline {
pub fn init(memory_allocator: Arc<impl MemoryAllocator>) {
let buffer = Buffer::from_data(
memory_allocator.clone(),
BufferCreateInfo {
usage: BufferUsage::UNIFORM_BUFFER,
..Default::default()
},
AllocationCreateInfo {
memory_type_filter: MemoryTypeFilter::PREFER_DEVICE
| MemoryTypeFilter::HOST_SEQUENTIAL_WRITE,
..Default::default()
},
data,
);
}
}

15
src/server/chunk/load.rs → src/server/chunk.rs
View File

@@ -3,8 +3,9 @@ use std::collections::{HashMap, HashSet};
use bevy_ecs::{entity::Entity, system::Commands};
use crate::{
common::component::{ChunkBundle, ChunkData, ChunkMesh, ChunkPos}, server::generation::generate_tree, util::
thread::{ExitType, ThreadChannel, ThreadHandle}
common::component::{ChunkBundle, ChunkData, ChunkMesh, ChunkPos},
server::generation::generate_tree,
util::{oct_tree::OctTree, thread::{ExitType, ThreadChannel, ThreadHandle}},
};
pub struct ChunkManager {
@@ -107,7 +108,14 @@ fn chunk_loader_main(channel: ThreadChannel<ServerChunkMsg, ChunkLoaderMsg>) {
let tree = ChunkData::from_tree(generate_tree(pos));
let tree_time = std::time::Instant::now() - start;
println!("gen time: {:<5?}; size: {}", tree_time, tree.raw().len());
// let worst = OctTree::from_fn(f_leaf, f_node, levels);
println!(
"gen time: {:<5?}; size: {} nodes = {} bytes",
tree_time,
tree.raw().len(),
std::mem::size_of_val(tree.raw())
);
channel.send(ServerChunkMsg::ChunkGenerated(GeneratedChunk {
pos,
@@ -121,4 +129,3 @@ fn chunk_loader_main(channel: ThreadChannel<ServerChunkMsg, ChunkLoaderMsg>) {
}
}
}

2
src/server/chunk/mod.rs
View File

@@ -1,2 +0,0 @@
mod load;
pub use load::*;

21
src/server/generation/mod.rs
View File

@@ -1,5 +1,5 @@
use fastnoise_simd::FastNoiseSIMD;
use nalgebra::Vector3;
use simdnoise::NoiseBuilder;
use crate::{
common::component::{chunk, ChunkPos},
@@ -11,11 +11,11 @@ pub fn generate_tree(pos: ChunkPos) -> OctTree {
return OctTree::from_leaf(0, 8);
}
let posf: Vector3<f32> = pos.cast() * chunk::SIDE_LENGTH as f32;
let noise1 = generate_noise_map(0, 1.0, posf, chunk::SCALE, &mut |v: f32| {
let noise1 = generate_noise_map(2, 1.0, posf, chunk::SCALE, &mut |v: f32| {
(v * 2.0).exp2() * TOP * 0.25
});
let noise2 = generate_noise_map(1, 50.0, posf, chunk::SCALE, &mut |v: f32| v * 20.0 + GRASS);
OctTree::from_fn_rec(
let noise2 = generate_noise_map(0, 50.0, posf, chunk::SCALE, &mut |v: f32| v * 20.0 + GRASS);
OctTree::from_fn(
&mut |p| generate_leaf(p, posf, (&noise1.base, &noise2.base)),
&mut |p, lvl| generate_node(p, lvl, posf, (&noise1, &noise2)),
chunk::SCALE,
@@ -89,10 +89,15 @@ fn generate_noise_map(
adjust: &mut impl FnMut(f32) -> f32,
) -> NoiseMap {
let mut size = 2usize.pow(levels);
let (mut base, min, max) = NoiseBuilder::gradient_2d_offset(posf.x, size, posf.z, size)
.with_seed(seed)
.with_freq(freq / (size as f32))
.generate();
let posi = Vector3::new(posf.x as i32, posf.y as i32, posf.z as i32);
let noise = FastNoiseSIMD::new(seed).get_simplex_set(posi.x, posi.y, posi.z, size as i32, size as i32, 1, freq * 150.0 / size as f32);
let mut base = noise.as_vec();
let (mut min, mut max) = (f32::MAX, f32::MIN);
for v in &base {
min = v.min(min);
max = v.max(max);
}
println!("{min}, {max}");
for v in &mut base {
*v = adjust((*v - min) / (max - min));
}

31
src/util/oct_tree.rs
View File

@@ -31,7 +31,13 @@ impl OctNode {
}
}
type OctNodeMap = FxHashMap<[OctNode; 8], OctNode>;
#[derive(Debug, Clone, Copy)]
struct OctNodeEntry {
pub node: OctNode,
pub count: u32,
}
type OctNodeMap = FxHashMap<[OctNode; 8], OctNodeEntry>;
#[derive(Debug, Clone)]
pub struct OctTree {
@@ -61,7 +67,10 @@ impl OctTree {
levels,
}
}
pub fn from_fn_rec(
pub fn from_leaf_fn(f_leaf: &mut impl FnMut(Vector3<usize>) -> u32, levels: u32) -> OctTree {
Self::from_fn(f_leaf, &mut |_, _| None, levels)
}
pub fn from_fn(
f_leaf: &mut impl FnMut(Vector3<usize>) -> u32,
f_node: &mut impl FnMut(Vector3<usize>, u32) -> Option<u32>,
levels: u32,
@@ -102,8 +111,9 @@ impl OctTree {
core::array::from_fn(|i| OctNode::new_leaf(f_leaf(offset + CORNERS[i])));
if leaves[1..].iter().all(|l| *l == leaves[0]) {
data.push(leaves[0]);
} else if let Some(node) = map.get(&leaves) {
data.push(*node);
} else if let Some(entry) = map.get_mut(&leaves) {
data.push(entry.node);
entry.count += 1;
} else {
data.extend_from_slice(&leaves);
}
@@ -128,7 +138,10 @@ impl OctTree {
let node = OctNode::new_node(sub_start as u32);
data[i + j] = node;
data_start += len;
map.insert(data[sub_start..sub_start+8].try_into().unwrap(), node);
map.insert(
data[sub_start..sub_start + 8].try_into().unwrap(),
OctNodeEntry { node, count: 1 },
);
}
}
}
@@ -137,14 +150,15 @@ impl OctTree {
if first.is_leaf() && data[i + 1..i + 8].iter().all(|l| *l == first) {
data.truncate(i);
data.push(first);
} else if let Some(node) = map.get(&data[i..i + 8]) {
} else if let Some(entry) = map.get_mut(&data[i..i + 8]) {
data.truncate(i);
data.push(*node);
data.push(entry.node);
entry.count += 1;
}
}
}
pub fn from_arr(arr: ArrayView3<u32>, levels: u32) -> Self {
Self::from_fn_rec(&mut |p| arr[(p.x, p.y, p.z)], &mut |_, _| None, levels)
Self::from_fn(&mut |p| arr[(p.x, p.y, p.z)], &mut |_, _| None, levels)
}
pub fn get(&self, mut pos: Vector3<usize>) -> u32 {
let mut data_start = 1;
@@ -164,7 +178,6 @@ impl OctTree {
pub fn raw(&self) -> &[OctNode] {
&self.data
}
pub fn mesh(&self) {}
}
pub struct OctTreeIter<'a> {