clean up and change generation to use exp2
This commit is contained in:
@@ -13,8 +13,8 @@ impl DepthTexture {
|
|||||||
label: &str,
|
label: &str,
|
||||||
) -> Self {
|
) -> Self {
|
||||||
let size = wgpu::Extent3d {
|
let size = wgpu::Extent3d {
|
||||||
width: config.width,
|
width: config.width + 1,
|
||||||
height: config.height,
|
height: config.height + 1,
|
||||||
depth_or_array_layers: 1,
|
depth_or_array_layers: 1,
|
||||||
};
|
};
|
||||||
let desc = wgpu::TextureDescriptor {
|
let desc = wgpu::TextureDescriptor {
|
||||||
@@ -88,7 +88,7 @@ impl StorageTexture {
|
|||||||
address_mode_w: wgpu::AddressMode::ClampToEdge,
|
address_mode_w: wgpu::AddressMode::ClampToEdge,
|
||||||
mag_filter: wgpu::FilterMode::Linear,
|
mag_filter: wgpu::FilterMode::Linear,
|
||||||
min_filter: wgpu::FilterMode::Linear,
|
min_filter: wgpu::FilterMode::Linear,
|
||||||
mipmap_filter: wgpu::FilterMode::Nearest,
|
mipmap_filter: wgpu::FilterMode::Linear,
|
||||||
compare: None,
|
compare: None,
|
||||||
lod_min_clamp: 0.0,
|
lod_min_clamp: 0.0,
|
||||||
lod_max_clamp: 100.0,
|
lod_max_clamp: 100.0,
|
||||||
|
|||||||
@@ -1,262 +0,0 @@
|
|||||||
@group(0) @binding(0)
|
|
||||||
var<uniform> view: View;
|
|
||||||
@group(0) @binding(1)
|
|
||||||
var<storage, read> voxels: array<u32>;
|
|
||||||
@group(0) @binding(2)
|
|
||||||
var<storage, read> voxel_groups: array<VoxelGroup>;
|
|
||||||
@group(0) @binding(3)
|
|
||||||
var<storage, read> global_lights: array<GlobalLight>;
|
|
||||||
@group(0) @binding(4)
|
|
||||||
var output: texture_storage_2d<rgba8unorm, write>;
|
|
||||||
|
|
||||||
struct GlobalLight {
|
|
||||||
dir: vec3<f32>,
|
|
||||||
};
|
|
||||||
|
|
||||||
struct View {
|
|
||||||
transform: mat4x4<f32>,
|
|
||||||
width: u32,
|
|
||||||
height: u32,
|
|
||||||
zoom: f32,
|
|
||||||
};
|
|
||||||
|
|
||||||
struct VoxelGroup {
|
|
||||||
transform: mat4x4<f32>,
|
|
||||||
transform_inv: mat4x4<f32>,
|
|
||||||
dimensions: vec3<u32>,
|
|
||||||
offset: u32,
|
|
||||||
};
|
|
||||||
|
|
||||||
@compute
|
|
||||||
@workgroup_size(8, 8, 1)
|
|
||||||
fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
|
|
||||||
// get position of the pixel; eye at origin, pixel on plane z = 1
|
|
||||||
if cell.x >= view.width || cell.y >= view.height {
|
|
||||||
return;
|
|
||||||
}
|
|
||||||
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
|
|
||||||
let aspect = win_dim.y / win_dim.x;
|
|
||||||
let pixel_pos = vec3<f32>(
|
|
||||||
(vec2<f32>(cell.xy) / win_dim - 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 ZERO3F = vec3<f32>(0.0);
|
|
||||||
const ZERO2F = vec2<f32>(0.0);
|
|
||||||
const DEPTH = 16u;
|
|
||||||
const FULL_ALPHA = 0.999;
|
|
||||||
|
|
||||||
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);
|
|
||||||
}
|
|
||||||
// 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;
|
|
||||||
var i = 0u;
|
|
||||||
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
|
|
||||||
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;
|
|
||||||
}
|
|
||||||
}
|
|
||||||
|
|
||||||
// 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);
|
|
||||||
let 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;
|
|
||||||
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
|
|
||||||
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;
|
|
||||||
let child_pos = u32(vcorner.x * 4 + vcorner.y * 2 + vcorner.z);
|
|
||||||
i = children_pos + child_pos;
|
|
||||||
data_start = children_pos + 8;
|
|
||||||
|
|
||||||
vox_pos -= vcorner * side_len;
|
|
||||||
low_corner += vec3<i32>(dir_to_vec(child_pos)) * i32(side_len);
|
|
||||||
|
|
||||||
continue;
|
|
||||||
}
|
|
||||||
|
|
||||||
// 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
|
|
||||||
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);
|
|
||||||
low_corner -= low_corner_adj;
|
|
||||||
|
|
||||||
// update vox pos to be relative to parent
|
|
||||||
vox_pos += low_corner_adj;
|
|
||||||
|
|
||||||
side_len *= 2;
|
|
||||||
}
|
|
||||||
return vec4<f32>(1.0, 0.0, 1.0, 1.0);
|
|
||||||
}
|
|
||||||
|
|
||||||
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.8, 0.2, 0.2, 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);
|
|
||||||
}
|
|
||||||
@@ -5,7 +5,7 @@ use nalgebra::{Projective3, Vector3};
|
|||||||
pub struct VoxelGroup {
|
pub struct VoxelGroup {
|
||||||
pub transform: Projective3<f32>,
|
pub transform: Projective3<f32>,
|
||||||
pub transform_inv: Projective3<f32>,
|
pub transform_inv: Projective3<f32>,
|
||||||
pub dimensions: Vector3<u32>,
|
pub scale: u32,
|
||||||
pub offset: u32,
|
pub offset: u32,
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
@@ -31,7 +31,7 @@ impl Layout {
|
|||||||
"global lights",
|
"global lights",
|
||||||
3,
|
3,
|
||||||
&[GlobalLight {
|
&[GlobalLight {
|
||||||
direction: Vector3::new(-1.0, -2.0, 2.0).normalize(),
|
direction: Vector3::new(-1.0, -2.3, 2.0).normalize(),
|
||||||
}],
|
}],
|
||||||
);
|
);
|
||||||
let texture = StorageTexture::init(
|
let texture = StorageTexture::init(
|
||||||
@@ -48,10 +48,9 @@ impl Layout {
|
|||||||
let render_bind_layout =
|
let render_bind_layout =
|
||||||
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
|
device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
|
||||||
entries: &[
|
entries: &[
|
||||||
view.bind_group_layout_entry(),
|
|
||||||
wgpu::BindGroupLayoutEntry {
|
wgpu::BindGroupLayoutEntry {
|
||||||
binding: 1,
|
binding: 0,
|
||||||
visibility: wgpu::ShaderStages::FRAGMENT,
|
visibility: wgpu::ShaderStages::FRAGMENT | wgpu::ShaderStages::VERTEX,
|
||||||
ty: wgpu::BindingType::Texture {
|
ty: wgpu::BindingType::Texture {
|
||||||
multisampled: false,
|
multisampled: false,
|
||||||
view_dimension: wgpu::TextureViewDimension::D2,
|
view_dimension: wgpu::TextureViewDimension::D2,
|
||||||
@@ -60,7 +59,7 @@ impl Layout {
|
|||||||
count: None,
|
count: None,
|
||||||
},
|
},
|
||||||
wgpu::BindGroupLayoutEntry {
|
wgpu::BindGroupLayoutEntry {
|
||||||
binding: 2,
|
binding: 1,
|
||||||
visibility: wgpu::ShaderStages::FRAGMENT,
|
visibility: wgpu::ShaderStages::FRAGMENT,
|
||||||
// This should match the filterable field of the
|
// This should match the filterable field of the
|
||||||
// corresponding Texture entry above.
|
// corresponding Texture entry above.
|
||||||
@@ -106,23 +105,24 @@ impl Layout {
|
|||||||
format: config.format,
|
format: config.format,
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn render_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
|
pub fn render_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
|
||||||
device.create_bind_group(&wgpu::BindGroupDescriptor {
|
device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||||
layout: &self.render_bind_layout,
|
layout: &self.render_bind_layout,
|
||||||
entries: &[
|
entries: &[
|
||||||
self.view.bind_group_entry(),
|
|
||||||
wgpu::BindGroupEntry {
|
wgpu::BindGroupEntry {
|
||||||
binding: 1,
|
binding: 0,
|
||||||
resource: wgpu::BindingResource::TextureView(&self.texture.view),
|
resource: wgpu::BindingResource::TextureView(&self.texture.view),
|
||||||
},
|
},
|
||||||
wgpu::BindGroupEntry {
|
wgpu::BindGroupEntry {
|
||||||
binding: 2,
|
binding: 1,
|
||||||
resource: wgpu::BindingResource::Sampler(&self.texture.sampler),
|
resource: wgpu::BindingResource::Sampler(&self.texture.sampler),
|
||||||
},
|
},
|
||||||
],
|
],
|
||||||
label: Some("tile_bind_group"),
|
label: Some("tile_bind_group"),
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn compute_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
|
pub fn compute_bind_group(&self, device: &wgpu::Device) -> wgpu::BindGroup {
|
||||||
device.create_bind_group(&wgpu::BindGroupDescriptor {
|
device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||||
layout: &self.compute_bind_layout,
|
layout: &self.compute_bind_layout,
|
||||||
@@ -136,6 +136,7 @@ impl Layout {
|
|||||||
label: Some("voxel compute"),
|
label: Some("voxel compute"),
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn render_pipeline(
|
pub fn render_pipeline(
|
||||||
&self,
|
&self,
|
||||||
device: &wgpu::Device,
|
device: &wgpu::Device,
|
||||||
@@ -179,6 +180,7 @@ impl Layout {
|
|||||||
cache: None,
|
cache: None,
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
|
|
||||||
pub fn compute_pipeline(
|
pub fn compute_pipeline(
|
||||||
&self,
|
&self,
|
||||||
device: &wgpu::Device,
|
device: &wgpu::Device,
|
||||||
|
|||||||
@@ -19,7 +19,7 @@ use crate::{
|
|||||||
use bevy_ecs::entity::Entity;
|
use bevy_ecs::entity::Entity;
|
||||||
pub use color::*;
|
pub use color::*;
|
||||||
use layout::Layout;
|
use layout::Layout;
|
||||||
use nalgebra::{Projective3, Transform3, Translation3, Vector2};
|
use nalgebra::{Projective3, Transform3, Translation3, Vector2, Vector3};
|
||||||
use std::{collections::HashMap, ops::Deref};
|
use std::{collections::HashMap, ops::Deref};
|
||||||
use wgpu::include_wgsl;
|
use wgpu::include_wgsl;
|
||||||
use {group::VoxelGroup, view::View};
|
use {group::VoxelGroup, view::View};
|
||||||
@@ -33,8 +33,8 @@ pub struct VoxelPipeline {
|
|||||||
id_map: HashMap<Entity, (usize, VoxelGroup)>,
|
id_map: HashMap<Entity, (usize, VoxelGroup)>,
|
||||||
}
|
}
|
||||||
|
|
||||||
const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("render.wgsl");
|
const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/render.wgsl");
|
||||||
const COMPUTE_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("compute.wgsl");
|
const COMPUTE_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/compute.wgsl");
|
||||||
|
|
||||||
impl VoxelPipeline {
|
impl VoxelPipeline {
|
||||||
pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self {
|
pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self {
|
||||||
@@ -82,7 +82,7 @@ impl VoxelPipeline {
|
|||||||
|
|
||||||
pub fn update_shader(&mut self, device: &wgpu::Device) {
|
pub fn update_shader(&mut self, device: &wgpu::Device) {
|
||||||
let Ok(shader) = std::fs::read_to_string(
|
let Ok(shader) = std::fs::read_to_string(
|
||||||
env!("CARGO_MANIFEST_DIR").to_owned() + "/src/client/render/voxel/ray_oct/compute.wgsl",
|
env!("CARGO_MANIFEST_DIR").to_owned() + "/src/client/render/voxel/ray_oct/shader/compute.wgsl",
|
||||||
) else {
|
) else {
|
||||||
println!("Failed to reload shader!");
|
println!("Failed to reload shader!");
|
||||||
return;
|
return;
|
||||||
@@ -125,7 +125,7 @@ impl VoxelPipeline {
|
|||||||
let group = VoxelGroup {
|
let group = VoxelGroup {
|
||||||
transform: proj,
|
transform: proj,
|
||||||
transform_inv: proj.inverse(),
|
transform_inv: proj.inverse(),
|
||||||
dimensions: chunk::DIMENSIONS.cast(),
|
scale: chunk::SCALE,
|
||||||
offset: offset as u32,
|
offset: offset as u32,
|
||||||
};
|
};
|
||||||
let updates = [ArrBufUpdate {
|
let updates = [ArrBufUpdate {
|
||||||
@@ -166,10 +166,11 @@ impl VoxelPipeline {
|
|||||||
update: UpdateGridTransform,
|
update: UpdateGridTransform,
|
||||||
) {
|
) {
|
||||||
if let Some((i, group)) = self.id_map.get_mut(&update.id) {
|
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()
|
let proj = Projective3::identity()
|
||||||
* Translation3::from(update.pos)
|
* Translation3::from(update.pos)
|
||||||
* update.orientation
|
* update.orientation
|
||||||
* Translation3::from(-group.dimensions.cast() / 2.0);
|
* Translation3::from(offset);
|
||||||
group.transform = proj;
|
group.transform = proj;
|
||||||
group.transform_inv = proj.inverse();
|
group.transform_inv = proj.inverse();
|
||||||
let updates = [ArrBufUpdate {
|
let updates = [ArrBufUpdate {
|
||||||
@@ -194,8 +195,6 @@ impl VoxelPipeline {
|
|||||||
let transform =
|
let transform =
|
||||||
Transform3::identity() * Translation3::from(camera.pos) * camera.orientation;
|
Transform3::identity() * Translation3::from(camera.pos) * camera.orientation;
|
||||||
let data = View {
|
let data = View {
|
||||||
width: size.x,
|
|
||||||
height: size.y,
|
|
||||||
zoom: camera.scale,
|
zoom: camera.scale,
|
||||||
transform,
|
transform,
|
||||||
};
|
};
|
||||||
|
|||||||
@@ -1,46 +0,0 @@
|
|||||||
// 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;
|
|
||||||
return textureSample(t_diffuse, s_diffuse, pos);
|
|
||||||
}
|
|
||||||
|
|
||||||
@@ -15,29 +15,28 @@ struct GlobalLight {
|
|||||||
|
|
||||||
struct View {
|
struct View {
|
||||||
transform: mat4x4<f32>,
|
transform: mat4x4<f32>,
|
||||||
width: u32,
|
|
||||||
height: u32,
|
|
||||||
zoom: f32,
|
zoom: f32,
|
||||||
};
|
};
|
||||||
|
|
||||||
struct VoxelGroup {
|
struct VoxelGroup {
|
||||||
transform: mat4x4<f32>,
|
transform: mat4x4<f32>,
|
||||||
transform_inv: mat4x4<f32>,
|
transform_inv: mat4x4<f32>,
|
||||||
dimensions: vec3<u32>,
|
scale: u32,
|
||||||
offset: u32,
|
offset: u32,
|
||||||
};
|
};
|
||||||
|
|
||||||
@compute
|
@compute
|
||||||
@workgroup_size(8, 8, 1)
|
@workgroup_size(8, 8, 1)
|
||||||
fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
|
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
|
// get position of the pixel; eye at origin, pixel on plane z = 1
|
||||||
if cell.x >= view.width || cell.y >= view.height {
|
if cell.x >= view_dim.x || cell.y >= view_dim.y {
|
||||||
return;
|
return;
|
||||||
}
|
}
|
||||||
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
|
let view_dim_f = vec2<f32>(view_dim);
|
||||||
let aspect = win_dim.y / win_dim.x;
|
let aspect = view_dim_f.y / view_dim_f.x;
|
||||||
let pixel_pos = vec3<f32>(
|
let pixel_pos = vec3<f32>(
|
||||||
(vec2<f32>(cell.xy) / win_dim - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
|
(vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
|
||||||
view.zoom
|
view.zoom
|
||||||
);
|
);
|
||||||
let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
|
let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
|
||||||
@@ -58,18 +57,17 @@ const ZERO2F = vec2<f32>(0.0);
|
|||||||
const FULL_ALPHA = 0.999;
|
const FULL_ALPHA = 0.999;
|
||||||
const EPSILON = 0.00000000001;
|
const EPSILON = 0.00000000001;
|
||||||
const MAX_ITERS = 1000;
|
const MAX_ITERS = 1000;
|
||||||
|
const MAX_SCALE: u32 = 10;
|
||||||
const MAX_DEPTH: u32 = 10;
|
|
||||||
const MAX_LENGTH: u32 = (1u << MAX_DEPTH);
|
|
||||||
|
|
||||||
fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
||||||
let gi = 0;
|
let gi = 0;
|
||||||
let group = voxel_groups[gi];
|
let group = voxel_groups[gi];
|
||||||
if group.dimensions.x == 0 {
|
if group.scale == 0 {
|
||||||
return vec4<f32>(0.0);
|
return vec4<f32>(0.0);
|
||||||
}
|
}
|
||||||
let dim_f = vec3<f32>(group.dimensions);
|
let dimensions = vec3<u32>(1u << group.scale);
|
||||||
let dim_i = vec3<i32>(group.dimensions);
|
let dim_f = vec3<f32>(dimensions);
|
||||||
|
let dim_i = vec3<i32>(dimensions);
|
||||||
|
|
||||||
// transform so that group is at 0,0
|
// transform so that group is at 0,0
|
||||||
let pos_start = (group.transform_inv * pos_view).xyz;
|
let pos_start = (group.transform_inv * pos_view).xyz;
|
||||||
@@ -127,11 +125,11 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
|||||||
let inv_dir_bits = 7 - dir_bits;
|
let inv_dir_bits = 7 - dir_bits;
|
||||||
|
|
||||||
var node_start = 1u;
|
var node_start = 1u;
|
||||||
var scale = MAX_DEPTH - 1;
|
var scale = group.scale - 1;
|
||||||
var half_t_span = f32(1u << scale) * inc_t;
|
var half_t_span = f32(1u << scale) * inc_t;
|
||||||
var t_center = t_min + half_t_span;
|
var t_center = t_min + half_t_span;
|
||||||
var color = vec4<f32>(0.0);
|
var color = vec4<f32>(0.0);
|
||||||
var parents = array<u32, MAX_DEPTH>();
|
var parents = array<u32, MAX_SCALE>();
|
||||||
|
|
||||||
var child = (u32(t > t_center.x) << 2) + (u32(t > t_center.y) << 1) + u32(t > t_center.z);
|
var child = (u32(t > t_center.x) << 2) + (u32(t > t_center.y) << 1) + u32(t > t_center.z);
|
||||||
var child_pos = dir_to_vec(child);
|
var child_pos = dir_to_vec(child);
|
||||||
@@ -145,8 +143,8 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
|||||||
iters += 1;
|
iters += 1;
|
||||||
let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
|
let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
|
||||||
if node >= LEAF_BIT {
|
if node >= LEAF_BIT {
|
||||||
|
if node != LEAF_BIT {
|
||||||
let vcolor = get_color(node & LEAF_MASK);
|
let vcolor = get_color(node & LEAF_MASK);
|
||||||
if vcolor.a > 0.0 {
|
|
||||||
let diffuse = max(dot(global_lights[0].dir, normals[axis]) + 0.1, 0.0);
|
let diffuse = max(dot(global_lights[0].dir, normals[axis]) + 0.1, 0.0);
|
||||||
let ambient = 0.2;
|
let ambient = 0.2;
|
||||||
let lighting = max(diffuse, ambient);
|
let lighting = max(diffuse, ambient);
|
||||||
@@ -164,11 +162,10 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
|||||||
// check if need to pop stack
|
// check if need to pop stack
|
||||||
if (child & move_dir) > 0 {
|
if (child & move_dir) > 0 {
|
||||||
// calculate new scale; first differing bit after adding
|
// calculate new scale; first differing bit after adding
|
||||||
let new_pos = vox_pos[axis] + (1u << scale);
|
let axis_pos = vox_pos[axis];
|
||||||
if new_pos == MAX_LENGTH { break; }
|
let differing = axis_pos ^ (axis_pos + (1u << scale));
|
||||||
let differing = vox_pos[axis] ^ new_pos;
|
|
||||||
vox_pos[axis] = new_pos;
|
|
||||||
scale = firstLeadingBit(differing);
|
scale = firstLeadingBit(differing);
|
||||||
|
if scale == group.scale { break; }
|
||||||
|
|
||||||
// restore & recalculate parent
|
// restore & recalculate parent
|
||||||
let parent_info = parents[scale];
|
let parent_info = parents[scale];
|
||||||
@@ -194,7 +191,7 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
|
|||||||
child_pos = vec3<u32>(vec3<f32>(t) > t_center);
|
child_pos = vec3<u32>(vec3<f32>(t) > t_center);
|
||||||
child = (child_pos.x << 2) + (child_pos.y << 1) + child_pos.z;
|
child = (child_pos.x << 2) + (child_pos.y << 1) + child_pos.z;
|
||||||
vox_pos += child_pos * (1u << scale);
|
vox_pos += child_pos * (1u << scale);
|
||||||
node_start = node_start + 8 + node;
|
node_start += 8 + node;
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
// return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
|
// return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
|
||||||
38
src/client/render/voxel/ray_oct/shader/render.wgsl
Normal file
38
src/client/render/voxel/ray_oct/shader/render.wgsl
Normal 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);
|
||||||
|
}
|
||||||
|
|
||||||
@@ -4,8 +4,6 @@ use nalgebra::Transform3;
|
|||||||
#[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
|
#[derive(Clone, Copy, PartialEq, bytemuck::Zeroable)]
|
||||||
pub struct View {
|
pub struct View {
|
||||||
pub transform: Transform3<f32>,
|
pub transform: Transform3<f32>,
|
||||||
pub width: u32,
|
|
||||||
pub height: u32,
|
|
||||||
pub zoom: f32,
|
pub zoom: f32,
|
||||||
}
|
}
|
||||||
|
|
||||||
@@ -14,8 +12,6 @@ unsafe impl bytemuck::Pod for View {}
|
|||||||
impl Default for View {
|
impl Default for View {
|
||||||
fn default() -> Self {
|
fn default() -> Self {
|
||||||
Self {
|
Self {
|
||||||
width: 1,
|
|
||||||
height: 1,
|
|
||||||
zoom: 1.0,
|
zoom: 1.0,
|
||||||
transform: Transform3::identity(),
|
transform: Transform3::identity(),
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -8,8 +8,8 @@ use bevy_derive::{Deref, DerefMut};
|
|||||||
use bevy_ecs::{bundle::Bundle, component::Component, entity::Entity, system::Resource};
|
use bevy_ecs::{bundle::Bundle, component::Component, entity::Entity, system::Resource};
|
||||||
use nalgebra::Vector3;
|
use nalgebra::Vector3;
|
||||||
|
|
||||||
pub const SIDE_POW: u32 = 10;
|
pub const SCALE: u32 = 8;
|
||||||
pub const SIDE_LENGTH: usize = 2usize.pow(SIDE_POW);
|
pub const SIDE_LENGTH: usize = 2usize.pow(SCALE);
|
||||||
pub const SHAPE: (usize, usize, usize) = (SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);
|
pub const SHAPE: (usize, usize, usize) = (SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);
|
||||||
pub const DIMENSIONS: Vector3<usize> = Vector3::new(SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);
|
pub const DIMENSIONS: Vector3<usize> = Vector3::new(SIDE_LENGTH, SIDE_LENGTH, SIDE_LENGTH);
|
||||||
pub const LEN: usize = SHAPE.0 * SHAPE.1 * SHAPE.2;
|
pub const LEN: usize = SHAPE.0 * SHAPE.1 * SHAPE.2;
|
||||||
@@ -26,7 +26,7 @@ impl ChunkData {
|
|||||||
}
|
}
|
||||||
pub fn empty() -> Self {
|
pub fn empty() -> Self {
|
||||||
Self {
|
Self {
|
||||||
data: OctTree::from_leaf(0, SIDE_POW),
|
data: OctTree::from_leaf(0, SCALE),
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|||||||
@@ -160,7 +160,7 @@ fn generate(pos: ChunkPos) -> Array3<u32> {
|
|||||||
return Array3::from_elem(shape, 0);
|
return Array3::from_elem(shape, 0);
|
||||||
}
|
}
|
||||||
let posf: Vector3<f32> = (pos.cast() * chunk::SIDE_LENGTH as f32) - Vector3::from_element(1.0);
|
let posf: Vector3<f32> = (pos.cast() * chunk::SIDE_LENGTH as f32) - Vector3::from_element(1.0);
|
||||||
let (noise, ..) = NoiseBuilder::gradient_2d_offset(
|
let (noise, min, max) = NoiseBuilder::gradient_2d_offset(
|
||||||
posf.x,
|
posf.x,
|
||||||
chunk::SIDE_LENGTH + 2,
|
chunk::SIDE_LENGTH + 2,
|
||||||
posf.z,
|
posf.z,
|
||||||
@@ -170,7 +170,7 @@ fn generate(pos: ChunkPos) -> Array3<u32> {
|
|||||||
.with_freq(0.005)
|
.with_freq(0.005)
|
||||||
.generate();
|
.generate();
|
||||||
Array3::from_shape_fn(shape, |(x, y, z)| {
|
Array3::from_shape_fn(shape, |(x, y, z)| {
|
||||||
generate_at(Vector3::new(x, y, z), posf, &noise)
|
generate_at(Vector3::new(x, y, z), posf, &noise, min, max)
|
||||||
})
|
})
|
||||||
}
|
}
|
||||||
|
|
||||||
@@ -179,31 +179,29 @@ fn generate_tree(pos: ChunkPos) -> OctTree {
|
|||||||
return OctTree::from_leaf(0, 8);
|
return OctTree::from_leaf(0, 8);
|
||||||
}
|
}
|
||||||
let posf: Vector3<f32> = pos.cast() * chunk::SIDE_LENGTH as f32;
|
let posf: Vector3<f32> = pos.cast() * chunk::SIDE_LENGTH as f32;
|
||||||
let (noise, ..) =
|
let (noise, min, max) =
|
||||||
NoiseBuilder::gradient_2d_offset(posf.x, chunk::SIDE_LENGTH, posf.z, chunk::SIDE_LENGTH)
|
NoiseBuilder::gradient_2d_offset(posf.x, chunk::SIDE_LENGTH, posf.z, chunk::SIDE_LENGTH)
|
||||||
.with_seed(0)
|
.with_seed(0)
|
||||||
.with_freq(0.01 / (chunk::SIDE_POW as f32))
|
.with_freq(1.0 / (chunk::SIDE_LENGTH as f32))
|
||||||
.generate();
|
.generate();
|
||||||
OctTree::from_fn_rec(&mut |p| generate_at(p, posf, &noise), chunk::SIDE_POW)
|
OctTree::from_fn_rec(&mut |p| generate_at(p, posf, &noise, min, max), chunk::SCALE)
|
||||||
}
|
}
|
||||||
|
|
||||||
fn generate_at(p: Vector3<usize>, posf: Vector3<f32>, noise: &[f32]) -> u32 {
|
fn generate_at(p: Vector3<usize>, posf: Vector3<f32>, noise: &[f32], min: f32, max: f32) -> u32 {
|
||||||
// 0 air 1 stone 2 "sand" 3 water
|
// 0 air 1 stone 2 "sand" 3 water
|
||||||
let y = p.y as f32 + posf.y;
|
let y = p.y as f32 + posf.y;
|
||||||
let [a, b, c] = [0.20, 0.35, 0.5].map(|f| chunk::SIDE_LENGTH as f32 * f);
|
let [a, b, c] = [0.18, 0.35, 0.5].map(|f| chunk::SIDE_LENGTH as f32 * f);
|
||||||
let n = (noise[p.x + p.z * chunk::SIDE_LENGTH] + 0.022) * (1.0 / 0.044) * c;
|
let n = ((noise[p.x + p.z * chunk::SIDE_LENGTH] - min) / (max - min) * 2.0).exp2() * c * 0.25;
|
||||||
if y < n.max(a) {
|
if y < n {
|
||||||
if y < a {
|
if y < a {
|
||||||
if y > n {
|
|
||||||
3
|
|
||||||
} else {
|
|
||||||
1
|
1
|
||||||
}
|
|
||||||
} else if y < b {
|
} else if y < b {
|
||||||
2
|
2
|
||||||
} else {
|
} else {
|
||||||
1
|
1
|
||||||
}
|
}
|
||||||
|
} else if y < a {
|
||||||
|
3
|
||||||
} else {
|
} else {
|
||||||
0
|
0
|
||||||
}
|
}
|
||||||
|
|||||||
Reference in New Issue
Block a user