From 1cecc3cdaf31adb617579f0ee76c7f606e9c6c23 Mon Sep 17 00:00:00 2001 From: shadow cat Date: Sat, 14 Sep 2024 17:09:20 -0400 Subject: [PATCH] added shader hot reloading, compute shader now uses some of paper technique --- src/client/handle_input.rs | 7 + src/client/render/mod.rs | 8 + src/client/render/voxel/ray_oct/compute.wgsl | 173 ++++----- .../render/voxel/ray_oct/compute_old.wgsl | 262 ++++++++++++++ src/client/render/voxel/ray_oct/layout.rs | 196 ++++++++++ src/client/render/voxel/ray_oct/mod.rs | 335 ++++-------------- src/client/render/voxel/ray_oct/render.wgsl | 1 - src/client/render/voxel/ray_oct/shader.wgsl | 212 ----------- 8 files changed, 612 insertions(+), 582 deletions(-) create mode 100644 src/client/render/voxel/ray_oct/compute_old.wgsl create mode 100644 src/client/render/voxel/ray_oct/layout.rs delete mode 100644 src/client/render/voxel/ray_oct/shader.wgsl diff --git a/src/client/handle_input.rs b/src/client/handle_input.rs index 132327f..35a1692 100644 --- a/src/client/handle_input.rs +++ b/src/client/handle_input.rs @@ -123,5 +123,12 @@ impl Client<'_> { })), })); } + + if input.just_pressed(Key::KeyR) { + self.renderer.update_shader(); + } + if input.just_pressed(Key::KeyT) { + self.renderer.reset_shader(); + } } } diff --git a/src/client/render/mod.rs b/src/client/render/mod.rs index c316a87..446048c 100644 --- a/src/client/render/mod.rs +++ b/src/client/render/mod.rs @@ -94,6 +94,14 @@ impl<'a> Renderer<'a> { } } + pub fn reset_shader(&mut self) { + self.voxel_pipeline.reset_shader(&self.device); + } + + pub fn update_shader(&mut self) { + self.voxel_pipeline.update_shader(&self.device); + } + fn create_encoder(device: &wgpu::Device) -> wgpu::CommandEncoder { device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("Render Encoder"), diff --git a/src/client/render/voxel/ray_oct/compute.wgsl b/src/client/render/voxel/ray_oct/compute.wgsl index a31111f..dfccc15 100644 --- a/src/client/render/voxel/ray_oct/compute.wgsl +++ b/src/client/render/voxel/ray_oct/compute.wgsl @@ -37,14 +37,13 @@ fn main(@builtin(global_invocation_id) cell: vec3) { let win_dim = vec2(f32(view.width), f32(view.height)); let aspect = win_dim.y / win_dim.x; let pixel_pos = vec3( - (vec2(cell.xy) / win_dim - vec2(0.5)) * vec2(2.0, 2.0 * aspect), + (vec2(cell.xy) / win_dim - vec2(0.5)) * vec2(2.0, -2.0 * aspect), view.zoom ); let pos = view.transform * vec4(pixel_pos, 1.0); let dir = view.transform * vec4(normalize(pixel_pos), 0.0); var color = trace_full(pos, dir); - // var color = vec4(0.0); 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(1.0, 1.0, 1.0); let sky_bg = vec3(0.3, 0.6, 1.0); @@ -56,8 +55,8 @@ fn main(@builtin(global_invocation_id) cell: vec3) { const ZERO3F = vec3(0.0); const ZERO2F = vec2(0.0); -const DEPTH = 16u; const FULL_ALPHA = 0.999; +const EPSILON = 0.00000000001; fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { let gi = 0; @@ -70,9 +69,13 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { // 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; + 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(0.0)); @@ -85,119 +88,95 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { var next_normal = vec3(0.0, 0.0, 0.0); // find where ray intersects with group - let plane_point = (vec3(1.0) - dir_if) / 2.0 * dim_f; + let pos_min = (vec3(1.0) - dir_uf) * dim_f; + let pos_max = dir_uf * dim_f; var pos = pos_start; + // time of intersection; x = td + p, solve for t + let t_min = (pos_min - pos) / dir; + let t_max = (pos_max - pos) / dir; var t = 0.0; if outside3f(pos, ZERO3F, dim_f) { - // time of intersection; x = td + p, solve for t - let t_i = (plane_point - pos) / dir; // points of intersection - let px = pos + t_i.x * dir; - let py = pos + t_i.y * dir; - let pz = pos + t_i.z * dir; + let px = pos + t_min.x * dir; + let py = pos + t_min.y * dir; + let pz = pos + t_min.z * dir; // check if point is in bounds let hit = vec3( inside2f(px.yz, ZERO2F, dim_f.yz), inside2f(py.xz, ZERO2F, dim_f.xz), inside2f(pz.xy, ZERO2F, dim_f.xy), - ) && (t_i > ZERO3F); + ) && (t_min > ZERO3F); if !any(hit) { return vec4(0.0); } pos = select(select(pz, py, hit.y), px, hit.x); - t = select(select(t_i.z, t_i.y, hit.y), t_i.x, hit.x); + t = select(select(t_min.z, t_min.y, hit.y), t_min.x, hit.x); next_normal = select(select(normals[2], normals[1], hit.y), normals[0], hit.x); } - // voxel position relative to low_corner var vox_pos = clamp(vec3(pos), vec3(0), dim_i - vec3(1)); - - + let inc_t = abs(1.0 / dir); let dir_i = vec3(dir_if); - let dir_u = ((dir_i + vec3(1)) / 2); - let dir_bits = u32(dir_u.x * 4 + dir_u.y * 2 + dir_u.z); - // time to move 1 unit using dir - let inc_t = abs(1.0 / dir); - var side_len = 256; - // "unsigned" minimum cube coords of current tree - var low_corner = vec3(0); - // time of next 1 unit plane hit in each direction - var color = vec4(0.0); - var data_start = 1u; - var i = 0u; - var axis = 0; - var parents = array(); + let dir_u = vec3((dir_i + vec3(1)) / 2); + let dir_bits = vec_to_dir(dir_u); + let inv_dir_bits = 7 - dir_bits; + + var node_start = 1u; + var t_center = (t_max + t_min) / 2.0; + var half_t_span = f32(256 / 2) * inc_t; var scale = 0; + + var dir_idx = vec3(vec3(t) > t_center); + var child = vec_to_dir(dir_idx); + var parents = array(); + var color = vec4(0.0); + for (var safety = 0; safety < 1000; safety += 1) { - let node = voxels[group.offset + i]; + let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)]; if node >= LEAF_BIT { - // leaf - let leaf = node & LEAF_MASK; - if leaf != 0 { - let vcolor = get_color(leaf); + let vcolor = get_color(node & LEAF_MASK); + if vcolor.a > 0.0 { let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); let ambient = 0.2; let lighting = max(diffuse, ambient); let new_color = min(vcolor.xyz * lighting, vec3(1.0)); color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); - if color.a > .999 { - // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); - return color; - } + if color.a > FULL_ALPHA { return color; } } - // move to next face of cube - let half_len = f32(side_len) / 2.0; - let corner = vec3(low_corner) + vec3(half_len) + dir_if * half_len; - let next_t = inc_t * abs(corner - pos_start); - axis = select(select(2, 1, next_t.y < next_t.z), 0, next_t.x < next_t.y && next_t.x < next_t.z); - t = next_t[axis]; - next_normal = normals[axis]; - pos = pos_start + t * dir; - let old = vox_pos[axis]; - vox_pos = vec3(pos) - low_corner; - vox_pos = clamp(vox_pos, vec3(0), side_len - vec3(1)); - vox_pos[axis] += dir_i[axis]; - } else if inside3i(vox_pos, vec3(0), vec3(side_len - 1)) { - // node - parents[scale] = (data_start << 3) + (data_start - i - 1); + let center_adj = t_center + half_t_span * vec3(dir_idx); + t = min(min(center_adj.x, center_adj.y), center_adj.z); + var move_dir = 0u; + if t == center_adj.x { move_dir = 4u; next_normal = normals[0]; } + if t == center_adj.y { move_dir = 2u; next_normal = normals[1]; } + if t == center_adj.z { move_dir = 1u; next_normal = normals[2]; } + if move_dir == 0 { return vec4(1.0); } + while (child & move_dir) > 0 { + if scale == 0 { return color; } + + scale -= 1; + let parent_info = parents[scale]; + node_start = parent_info >> 3; + child = parent_info & 7; + + dir_idx = dir_to_vec(child); + t_center -= half_t_span * (vec3(dir_idx * 2) - 1.0); + half_t_span *= 2.0; + } + child = child ^ move_dir; + dir_idx = dir_to_vec(child); + } else { + parents[scale] = (node_start << 3) + child; 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; + half_t_span /= 2.0; + t_center += half_t_span * (vec3(dir_idx * 2) - 1.0); + dir_idx = vec3(vec3(t) > t_center); - vox_pos -= vcorner * side_len; - low_corner += vec3(dir_to_vec(child_pos)) * i32(side_len); - - continue; + child = vec_to_dir(dir_idx); + node_start = node_start + 8 + node; } - - // exit if highest node - if scale == 0 { - // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); - return color; - } - - // get parent info and reset "pointers" to parent - scale -= 1; - let parent_info = parents[scale]; - let loc = 8 - (data_start - i); - data_start = parent_info >> 3; - i = data_start - ((parent_info & 7) + 1); - - // adjust corner back to parent - let low_corner_adj = vec3(dir_to_vec(loc)) * i32(side_len); - low_corner -= low_corner_adj; - - // update vox pos to be relative to parent - vox_pos += low_corner_adj; - - side_len *= 2; } return vec4(1.0, 0.0, 1.0, 1.0); } @@ -205,31 +184,13 @@ fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { const LEAF_BIT = 1u << 31u; const LEAF_MASK = ~LEAF_BIT; -// there's no way this is efficient, mod is faster for all I know +// no clue if this is efficient, mod is faster for all I know fn dir_to_vec(bits: u32) -> vec3 { return vec3(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u)); } -fn get_voxel(offset: u32, pos_: vec3) -> u32 { - var data_start = 1u; - var i = 0u; - var pos = pos_; - var side_len: u32 = 256; - var safety = 0; - while voxels[offset + i] < LEAF_BIT { - let node_pos = data_start + voxels[offset + i]; - side_len /= 2u; - let corner = pos / side_len; - pos -= corner * side_len; - let j = corner.x * 4 + corner.y * 2 + corner.z; - i = node_pos + j; - data_start = node_pos + 8; - if safety == 10 { - return 10u; - } - safety += 1; - } - return voxels[offset + i] & LEAF_MASK; +fn vec_to_dir(vec: vec3) -> u32 { + return vec.x * 4 + vec.y * 2 + vec.z * 1; } fn get_color(id: u32) -> vec4 { @@ -241,7 +202,7 @@ fn get_color(id: u32) -> vec4 { return vec4(0.5, 0.5, 0.5, 1.0); } case 2u: { - return vec4(0.8, 0.2, 0.2, 1.0); + return vec4(0.5, 1.0, 0.5, 1.0); } case 3u: { return vec4(0.5, 0.5, 1.0, 0.5); diff --git a/src/client/render/voxel/ray_oct/compute_old.wgsl b/src/client/render/voxel/ray_oct/compute_old.wgsl new file mode 100644 index 0000000..59b749f --- /dev/null +++ b/src/client/render/voxel/ray_oct/compute_old.wgsl @@ -0,0 +1,262 @@ +@group(0) @binding(0) +var view: View; +@group(0) @binding(1) +var voxels: array; +@group(0) @binding(2) +var voxel_groups: array; +@group(0) @binding(3) +var global_lights: array; +@group(0) @binding(4) +var output: texture_storage_2d; + +struct GlobalLight { + dir: vec3, +}; + +struct View { + transform: mat4x4, + width: u32, + height: u32, + zoom: f32, +}; + +struct VoxelGroup { + transform: mat4x4, + transform_inv: mat4x4, + dimensions: vec3, + offset: u32, +}; + +@compute +@workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) cell: vec3) { + // 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(view.width), f32(view.height)); + let aspect = win_dim.y / win_dim.x; + let pixel_pos = vec3( + (vec2(cell.xy) / win_dim - vec2(0.5)) * vec2(2.0, -2.0 * aspect), + view.zoom + ); + let pos = view.transform * vec4(pixel_pos, 1.0); + let dir = view.transform * vec4(normalize(pixel_pos), 0.0); + + var color = trace_full(pos, dir); + let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0); + let sun_color = light_mult * vec3(1.0, 1.0, 1.0); + let sky_bg = vec3(0.3, 0.6, 1.0); + let sky_color = sun_color + sky_bg * (1.0 - light_mult); + color += vec4(sky_color * (1.0 - color.a), 1.0 - color.a); + color.a = 1.0; + textureStore(output, cell.xy, color); +} + +const ZERO3F = vec3(0.0); +const ZERO2F = vec2(0.0); +const DEPTH = 16u; +const FULL_ALPHA = 0.999; + +fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { + let gi = 0; + let group = voxel_groups[gi]; + if group.dimensions.x == 0 { + return vec4(0.0); + } + let dim_f = vec3(group.dimensions); + let dim_i = vec3(group.dimensions); + + // transform so that group is at 0,0 + let pos_start = (group.transform_inv * pos_view).xyz; + let dir = (group.transform_inv * dir_view).xyz; + + let dir_if = sign(dir); + + + + // calculate normals + var normals = mat3x3( + (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, + (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, + (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, + ); + var next_normal = vec3(0.0, 0.0, 0.0); + + // find where ray intersects with group + let plane_point = (vec3(1.0) - dir_if) / 2.0 * dim_f; + var pos = pos_start; + var t = 0.0; + if outside3f(pos, ZERO3F, dim_f) { + // time of intersection; x = td + p, solve for t + let t_i = (plane_point - pos) / dir; + // points of intersection + let px = pos + t_i.x * dir; + let py = pos + t_i.y * dir; + let pz = pos + t_i.z * dir; + + // check if point is in bounds + let hit = vec3( + inside2f(px.yz, ZERO2F, dim_f.yz), + inside2f(py.xz, ZERO2F, dim_f.xz), + inside2f(pz.xy, ZERO2F, dim_f.xy), + ) && (t_i > ZERO3F); + if !any(hit) { + return vec4(0.0); + } + pos = select(select(pz, py, hit.y), px, hit.x); + t = select(select(t_i.z, t_i.y, hit.y), t_i.x, hit.x); + next_normal = select(select(normals[2], normals[1], hit.y), normals[0], hit.x); + } + // voxel position relative to low_corner + var vox_pos = clamp(vec3(pos), vec3(0), dim_i - vec3(1)); + + + + let dir_i = vec3(dir_if); + let dir_u = ((dir_i + vec3(1)) / 2); + let dir_bits = u32(dir_u.x * 4 + dir_u.y * 2 + dir_u.z); + // time to move 1 unit using dir + let inc_t = abs(1.0 / dir); + var side_len = 256; + // "unsigned" minimum cube coords of current tree + var low_corner = vec3(0); + // time of next 1 unit plane hit in each direction + var color = vec4(0.0); + var data_start = 1u; + var i = 0u; + var parents = array(); + var scale = 0; + for (var safety = 0; safety < 1000; safety += 1) { + let node = voxels[group.offset + i]; + if node >= LEAF_BIT { + // leaf + let leaf = node & LEAF_MASK; + if leaf != 0 { + let vcolor = get_color(leaf); + let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); + let ambient = 0.2; + let lighting = max(diffuse, ambient); + let new_color = min(vcolor.xyz * lighting, vec3(1.0)); + color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); + if color.a > .999 { + // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); + return color; + } + } + + // move to next face of cube + let half_len = f32(side_len) / 2.0; + let corner = vec3(low_corner) + vec3(half_len) + dir_if * half_len; + let next_t = inc_t * abs(corner - pos_start); + 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(pos) - low_corner; + vox_pos = clamp(vox_pos, vec3(0), side_len - vec3(1)); + vox_pos[axis] += dir_i[axis]; + } else if inside3i(vox_pos, vec3(0), vec3(side_len - 1)) { + // node + 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(dir_to_vec(child_pos)) * i32(side_len); + + continue; + } + + // exit if highest node + if scale == 0 { + // return vec4(f32(safety) / 1000.0, 0.0, 0.0, 1.0); + return color; + } + + // get parent info and reset "pointers" to parent + scale -= 1; + let parent_info = parents[scale]; + let loc = 8 - (data_start - i); + data_start = parent_info >> 3; + i = data_start - ((parent_info & 7) + 1); + + // adjust corner back to parent + let low_corner_adj = vec3(dir_to_vec(loc)) * i32(side_len); + low_corner -= low_corner_adj; + + // update vox pos to be relative to parent + vox_pos += low_corner_adj; + + side_len *= 2; + } + return vec4(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 { + return vec3(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u)); +} + +fn get_voxel(offset: u32, pos_: vec3) -> u32 { + var data_start = 1u; + var i = 0u; + var pos = pos_; + var side_len: u32 = 256; + var safety = 0; + while voxels[offset + i] < LEAF_BIT { + let node_pos = data_start + voxels[offset + i]; + side_len /= 2u; + let corner = pos / side_len; + pos -= corner * side_len; + let j = corner.x * 4 + corner.y * 2 + corner.z; + i = node_pos + j; + data_start = node_pos + 8; + if safety == 10 { + return 10u; + } + safety += 1; + } + return voxels[offset + i] & LEAF_MASK; +} + +fn get_color(id: u32) -> vec4 { + switch id { + case 0u: { + return vec4(0.0); + } + case 1u: { + return vec4(0.5, 0.5, 0.5, 1.0); + } + case 2u: { + return vec4(0.8, 0.2, 0.2, 1.0); + } + case 3u: { + return vec4(0.5, 0.5, 1.0, 0.5); + } + default: { + return vec4(1.0, 0.0, 0.0, 1.0); + } + } +} + +fn outside3f(v: vec3, low: vec3, high: vec3) -> bool { + return any(v < low) || any(v > high); +} + +fn inside2f(v: vec2, low: vec2, high: vec2) -> bool { + return all(v >= low) && all(v <= high); +} + +fn inside3i(v: vec3, low: vec3, high: vec3) -> bool { + return all(v >= low) && all(v <= high); +} diff --git a/src/client/render/voxel/ray_oct/layout.rs b/src/client/render/voxel/ray_oct/layout.rs new file mode 100644 index 0000000..1243cbc --- /dev/null +++ b/src/client/render/voxel/ray_oct/layout.rs @@ -0,0 +1,196 @@ +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, + pub voxel_groups: Storage, + pub voxels: Storage, + pub global_lights: Storage, + 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(-0.5, -4.0, 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: &[ + view.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("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: &[ + 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"), + }) + } + 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, + }) + } +} diff --git a/src/client/render/voxel/ray_oct/mod.rs b/src/client/render/voxel/ray_oct/mod.rs index 5cc9926..fa3ff75 100644 --- a/src/client/render/voxel/ray_oct/mod.rs +++ b/src/client/render/voxel/ray_oct/mod.rs @@ -1,220 +1,70 @@ mod color; mod grid; mod group; +mod layout; mod light; mod view; -pub use color::*; -use wgpu::include_wgsl; - use super::super::UpdateGridTransform; use crate::{ client::{ camera::Camera, render::{ - util::{ArrBufUpdate, DepthTexture, Storage, StorageTexture, Uniform}, + util::{ArrBufUpdate, StorageTexture}, AddChunk, CreateVoxelGrid, }, }, common::component::chunk, - util::oct_tree::OctNode, }; use bevy_ecs::entity::Entity; -use light::GlobalLight; -use nalgebra::{Projective3, Transform3, Translation3, Vector2, Vector3}; +pub use color::*; +use layout::Layout; +use nalgebra::{Projective3, Transform3, Translation3, Vector2}; use std::{collections::HashMap, ops::Deref}; - +use wgpu::include_wgsl; use {group::VoxelGroup, view::View}; pub struct VoxelPipeline { + layout: Layout, compute_pipeline: wgpu::ComputePipeline, - texture: StorageTexture, - cbind_group_layout: wgpu::BindGroupLayout, - cbind_group: wgpu::BindGroup, - - pipeline: wgpu::RenderPipeline, - view: Uniform, - bind_group_layout: wgpu::BindGroupLayout, - bind_group: wgpu::BindGroup, - voxel_groups: Storage, - voxels: Storage, - global_lights: Storage, + compute_bind_group: wgpu::BindGroup, + render_pipeline: wgpu::RenderPipeline, + render_bind_group: wgpu::BindGroup, id_map: HashMap, } +const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("render.wgsl"); +const COMPUTE_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("compute.wgsl"); + impl VoxelPipeline { pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self { // shaders - let shader = device.create_shader_module(include_wgsl!("render.wgsl")); - 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(-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, - ); + let layout = Layout::init(device, config); - // bind groups - let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { - entries: &[ - view.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"), - }); + 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 bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - layout: &bind_group_layout, - entries: &[ - view.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"), - }); - - // pipeline - let render_pipeline_layout = - device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { - label: Some("Tile Pipeline Layout"), - bind_group_layouts: &[&bind_group_layout], - push_constant_ranges: &[], - }); - - let render_pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor { - label: Some("Voxel Pipeline"), - layout: Some(&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: config.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, - }); - - 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, - }); + 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, - texture, - cbind_group_layout, - cbind_group, - pipeline: render_pipeline, - view, - bind_group, - bind_group_layout, - voxels, - voxel_groups, - global_lights, + 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, @@ -228,46 +78,29 @@ impl VoxelPipeline { grid, }: CreateVoxelGrid, ) { - // let offset = self.voxels.len(); - // - // let updates = [ArrBufUpdate { - // offset, - // data: &grid.as_slice().unwrap(), - // }]; - // let size = offset + grid.len(); - // self.voxels.update(device, encoder, belt, size, &updates); - // - // let proj = Projective3::identity() - // * Translation3::from(pos) - // * orientation - // * Translation3::from(-dimensions.cast() / 2.0); - // let group = VoxelGroup { - // transform: proj, - // transform_inv: proj.inverse(), - // dimensions: dimensions.cast(), - // offset: offset as u32, - // }; - // let updates = [ArrBufUpdate { - // offset: self.voxel_groups.len(), - // data: &[group], - // }]; - // let i = self.voxel_groups.len(); - // let size = i + 1; - // self.voxel_groups - // .update(device, encoder, belt, size, &updates); - // - // self.id_map.insert(id, (i, group)); - // - // self.bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { - // layout: &self.bind_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(), - // ], - // label: Some("tile_bind_group"), - // }); + } + + 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/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( @@ -277,12 +110,14 @@ impl VoxelPipeline { belt: &mut wgpu::util::StagingBelt, AddChunk { id, pos, tree, .. }: AddChunk, ) { - let offset = self.voxels.len(); + let offset = self.layout.voxels.len(); let data = tree.raw(); let updates = [ArrBufUpdate { offset, data }]; let size = offset + data.len(); - self.voxels.update(device, encoder, belt, size, &updates); + self.layout + .voxels + .update(device, encoder, belt, size, &updates); let proj = Projective3::identity() * Translation3::from((pos.deref() * chunk::SIDE_LENGTH as i32).cast()) @@ -294,34 +129,21 @@ impl VoxelPipeline { offset: offset as u32, }; let updates = [ArrBufUpdate { - offset: self.voxel_groups.len(), + offset: self.layout.voxel_groups.len(), data: &[group], }]; - let i = self.voxel_groups.len(); + let i = self.layout.voxel_groups.len(); let size = i + 1; - self.voxel_groups + self.layout + .voxel_groups .update(device, encoder, belt, size, &updates); self.id_map.insert(id, (i, group)); - self.update_cbind_group(device); - } - - pub fn update_cbind_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"), - }); + self.compute_bind_group = self.layout.compute_bind_group(device); } pub fn resize(&mut self, device: &wgpu::Device, size: Vector2) { - self.texture = StorageTexture::init( + self.layout.texture = StorageTexture::init( device, wgpu::Extent3d { width: size.x, @@ -332,22 +154,8 @@ impl VoxelPipeline { wgpu::ShaderStages::COMPUTE | wgpu::ShaderStages::FRAGMENT, 4, ); - self.update_cbind_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"), - }); + self.compute_bind_group = self.layout.compute_bind_group(device); + self.render_bind_group = self.layout.render_bind_group(device); } pub fn update_transform( @@ -368,8 +176,9 @@ impl VoxelPipeline { offset: *i, data: &[*group], }]; - let size = self.voxel_groups.len(); - self.voxel_groups + let size = self.layout.voxel_groups.len(); + self.layout + .voxel_groups .update(device, encoder, belt, size, &updates); } } @@ -390,12 +199,12 @@ impl VoxelPipeline { zoom: camera.scale, transform, }; - self.view.update(device, encoder, belt, data) + 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.pipeline); - render_pass.set_bind_group(0, &self.bind_group, &[]); + render_pass.set_pipeline(&self.render_pipeline); + render_pass.set_bind_group(0, &self.render_bind_group, &[]); render_pass.draw(0..4, 0..1); } @@ -403,8 +212,8 @@ impl VoxelPipeline { pub fn compute(&self, pass: &mut wgpu::ComputePass) { pass.set_pipeline(&self.compute_pipeline); - pass.set_bind_group(0, &self.cbind_group, &[]); - let buf = &self.texture.buf; + 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); diff --git a/src/client/render/voxel/ray_oct/render.wgsl b/src/client/render/voxel/ray_oct/render.wgsl index d325d67..53c39ba 100644 --- a/src/client/render/voxel/ray_oct/render.wgsl +++ b/src/client/render/voxel/ray_oct/render.wgsl @@ -41,7 +41,6 @@ fn fs_main( ) -> @location(0) vec4 { let win_dim = vec2(f32(view.width), f32(view.height)); var pos = in.clip_position.xy / win_dim; - pos.y = 1.0 - pos.y; return textureSample(t_diffuse, s_diffuse, pos); } diff --git a/src/client/render/voxel/ray_oct/shader.wgsl b/src/client/render/voxel/ray_oct/shader.wgsl deleted file mode 100644 index bee2c87..0000000 --- a/src/client/render/voxel/ray_oct/shader.wgsl +++ /dev/null @@ -1,212 +0,0 @@ -// Vertex shader - -struct GlobalLight { - dir: vec3, -}; - -struct VertexOutput { - @builtin(position) clip_position: vec4, -}; - -struct View { - transform: mat4x4, - width: u32, - height: u32, - zoom: f32, -}; - -struct VoxelGroup { - transform: mat4x4, - transform_inv: mat4x4, - dimensions: vec3, - offset: u32, -}; - -@group(0) @binding(0) -var view: View; -@group(0) @binding(1) -var voxels: array; -@group(0) @binding(2) -var voxel_groups: array; -@group(0) @binding(3) -var global_lights: array; - -@vertex -fn vs_main( - @builtin(vertex_index) vi: u32, - @builtin(instance_index) ii: u32, -) -> VertexOutput { - var out: VertexOutput; - - var pos = vec2( - f32(vi % 2u) * 2.0 - 1.0, - f32(vi / 2u) * 2.0 - 1.0, - ) ; - out.clip_position = vec4(pos.x, pos.y, 0.0, 1.0); - return out; -} - -// Fragment shader - -@fragment -fn fs_main( - in: VertexOutput, -) -> @location(0) vec4 { - // get position of the pixel; eye at origin, pixel on plane z = 1 - let win_dim = vec2(f32(view.width), f32(view.height)); - let aspect = win_dim.y / win_dim.x; - let pixel_pos = vec3( - (in.clip_position.xy / win_dim - vec2(0.5)) * vec2(2.0, -2.0 * aspect), - view.zoom - ); - - // move to position in world - let pos = view.transform * vec4(pixel_pos, 1.0); - let dir = view.transform * vec4(normalize(pixel_pos), 0.0); - - var color = trace_full(pos, dir); - let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0); - let sky_color = light_mult * vec3(1.0, 1.0, 1.0); - color += vec4(sky_color * (1.0 - color.a), 1.0 - color.a); - color.a = 1.0; - return color; -} - -const ZERO3F = vec3(0.0); -const ZERO2F = vec2(0.0); -const DEPTH = 16u; -const FULL_ALPHA = 0.9999; - -fn trace_full(pos_view: vec4, dir_view: vec4) -> vec4 { - let gi = 0; - let group = voxel_groups[gi]; - if group.dimensions.x == 0 { - return vec4(0.0); - } - let dim_f = vec3(group.dimensions); - let dim_i = vec3(group.dimensions); - - // transform so that group is at 0,0 - let pos_start = (group.transform_inv * pos_view).xyz; - let dir = (group.transform_inv * dir_view).xyz; - - let dir_if = sign(dir); - let dir_uf = max(dir_if, vec3(0.0)); - - - - // calculate normals - var normals = mat3x3( - (group.transform * vec4(dir_if.x, 0.0, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, dir_if.y, 0.0, 0.0)).xyz, - (group.transform * vec4(0.0, 0.0, dir_if.z, 0.0)).xyz, - ); - var next_normal = vec3(0.0, 0.0, 0.0); - - // find where ray intersects with group - let pos_min = (vec3(1.0) - dir_uf) * dim_f; - let pos_max = dir_uf * dim_f; - var pos = pos_start; - // time of intersection; x = td + p, solve for t - let t_min = (pos_min - pos) / dir; - let t_max = (pos_max - pos) / dir; - var t = 0.0; - if outside3f(pos, ZERO3F, dim_f) { - // 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( - 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(0.0); - } - pos = select(select(pz, py, hit.y), px, 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); - } - let inc_t = abs(1.0 / dir); - let dir_i = vec3(dir_if); - let dir_u = vec3((dir_i + vec3(1)) / 2); - var i = 0u; - 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 { - let vcolor = get_color(node & LEAF_MASK); - if vcolor.a > 0.0 { - // let diffuse = max(dot(global_lights[0].dir, next_normal) + 0.1, 0.0); - // let ambient = 0.2; - // let lighting = max(diffuse, ambient); - // let new_color = min(vcolor.xyz * lighting, vec3(1.0)); - // color += vec4(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a); - if vcolor.a > .999 { - return vcolor; - } - } - return vcolor; - } else { - half_t_span *= 0.5; - let dir_idx = vec3(vec3(t) < t_center); - t_center += half_t_span * (1.0 - vec3(dir_idx * 2)); - - let child_i = vec_to_dir(dir_idx ^ dir_u); - let node_pos = data_start + node; - i = node_pos + child_i; - data_start = node_pos + 8; - continue; - } - } - return vec4(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 { - return vec3(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u)); -} - -fn vec_to_dir(vec: vec3) -> u32 { - return vec.x * 4 + vec.y * 2 + vec.z * 1; -} - -fn get_color(id: u32) -> vec4 { - switch id { - case 0u: { - return vec4(0.0); - } - case 1u: { - return vec4(0.5, 0.5, 0.5, 1.0); - } - case 2u: { - return vec4(0.5, 1.0, 0.5, 1.0); - } - case 3u: { - return vec4(0.5, 0.5, 1.0, 0.5); - } - default: { - return vec4(1.0, 0.0, 0.0, 1.0); - } - } -} - -fn outside3f(v: vec3, low: vec3, high: vec3) -> bool { - return any(v < low) || any(v > high); -} - -fn inside2f(v: vec2, low: vec2, high: vec2) -> bool { - return all(v >= low) && all(v <= high); -} - -fn inside3i(v: vec3, low: vec3, high: vec3) -> bool { - return all(v >= low) && all(v <= high); -}