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

added shader hot reloading, compute shader now uses some of paper technique

Browse Source
This commit is contained in:
shadow cat
2024-09-14 17:09:20 -04:00
parent 244440c38a
commit 1cecc3cdaf
8 changed files with 612 additions and 582 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
src/client/handle_input.rs
View File

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

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

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

169
src/client/render/voxel/ray_oct/compute.wgsl
View File

@@ -37,14 +37,13 @@ fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
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),
(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);
// var color = vec4<f32>(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<f32>(1.0, 1.0, 1.0);
let sky_bg = vec3<f32>(0.3, 0.6, 1.0);
@@ -56,8 +55,8 @@ fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
const ZERO3F = vec3<f32>(0.0);
const ZERO2F = vec2<f32>(0.0);
const DEPTH = 16u;
const FULL_ALPHA = 0.999;
const EPSILON = 0.00000000001;
fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
let gi = 0;
@@ -70,9 +69,13 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
// 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<f32>(0.0));
@@ -85,119 +88,95 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
var next_normal = vec3<f32>(0.0, 0.0, 0.0);
// find where ray intersects with group
let plane_point = (vec3<f32>(1.0) - dir_if) / 2.0 * dim_f;
let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
let pos_max = dir_uf * dim_f;
var pos = pos_start;
// time of intersection; x = td + p, solve for t
let t_min = (pos_min - pos) / dir;
let t_max = (pos_max - pos) / dir;
var t = 0.0;
if outside3f(pos, ZERO3F, dim_f) {
// time of intersection; x = td + p, solve for t
let t_i = (plane_point - pos) / dir;
// points of intersection
let px = pos + t_i.x * dir;
let py = pos + t_i.y * dir;
let pz = pos + t_i.z * dir;
let px = pos + t_min.x * dir;
let py = pos + t_min.y * dir;
let pz = pos + t_min.z * dir;
// check if point is in bounds
let hit = vec3<bool>(
inside2f(px.yz, ZERO2F, dim_f.yz),
inside2f(py.xz, ZERO2F, dim_f.xz),
inside2f(pz.xy, ZERO2F, dim_f.xy),
) && (t_i > ZERO3F);
) && (t_min > ZERO3F);
if !any(hit) {
return vec4<f32>(0.0);
}
pos = select(select(pz, py, hit.y), px, hit.x);
t = select(select(t_i.z, t_i.y, hit.y), t_i.x, hit.x);
t = select(select(t_min.z, t_min.y, hit.y), t_min.x, hit.x);
next_normal = select(select(normals[2], normals[1], hit.y), normals[0], hit.x);
}
// voxel position relative to low_corner
var vox_pos = clamp(vec3<i32>(pos), vec3<i32>(0), dim_i - vec3<i32>(1));
let inc_t = abs(1.0 / dir);
let dir_i = vec3<i32>(dir_if);
let dir_u = ((dir_i + vec3<i32>(1)) / 2);
let dir_bits = u32(dir_u.x * 4 + dir_u.y * 2 + dir_u.z);
// time to move 1 unit using dir
let inc_t = abs(1.0 / dir);
var side_len = 256;
// "unsigned" minimum cube coords of current tree
var low_corner = vec3<i32>(0);
// time of next 1 unit plane hit in each direction
var color = vec4<f32>(0.0);
var data_start = 1u;
var i = 0u;
var axis = 0;
var parents = array<u32, 8>();
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 t_center = (t_max + t_min) / 2.0;
var half_t_span = f32(256 / 2) * inc_t;
var scale = 0;
var dir_idx = vec3<u32>(vec3<f32>(t) > t_center);
var child = vec_to_dir(dir_idx);
var parents = array<u32, 8>();
var color = vec4<f32>(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<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;
}
if color.a > FULL_ALPHA { return color; }
}
// move to next face of cube
let half_len = f32(side_len) / 2.0;
let corner = vec3<f32>(low_corner) + vec3<f32>(half_len) + dir_if * half_len;
let next_t = inc_t * abs(corner - pos_start);
axis = select(select(2, 1, next_t.y < next_t.z), 0, next_t.x < next_t.y && next_t.x < next_t.z);
t = next_t[axis];
next_normal = normals[axis];
pos = pos_start + t * dir;
let old = vox_pos[axis];
vox_pos = vec3<i32>(pos) - low_corner;
vox_pos = 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 center_adj = t_center + half_t_span * vec3<f32>(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<f32>(1.0); }
while (child & move_dir) > 0 {
if scale == 0 { return color; }
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);
node_start = parent_info >> 3;
child = parent_info & 7;
// adjust corner back to parent
let low_corner_adj = vec3<i32>(dir_to_vec(loc)) * i32(side_len);
low_corner -= low_corner_adj;
dir_idx = dir_to_vec(child);
t_center -= half_t_span * (vec3<f32>(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;
// update vox pos to be relative to parent
vox_pos += low_corner_adj;
half_t_span /= 2.0;
t_center += half_t_span * (vec3<f32>(dir_idx * 2) - 1.0);
dir_idx = vec3<u32>(vec3<f32>(t) > t_center);
side_len *= 2;
child = vec_to_dir(dir_idx);
node_start = node_start + 8 + node;
}
}
return vec4<f32>(1.0, 0.0, 1.0, 1.0);
}
@@ -205,31 +184,13 @@ fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
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<u32> {
return vec3<u32>(extractBits(bits, 2u, 1u), extractBits(bits, 1u, 1u), extractBits(bits, 0u, 1u));
}
fn get_voxel(offset: u32, pos_: vec3<u32>) -> u32 {
var data_start = 1u;
var i = 0u;
var pos = pos_;
var side_len: u32 = 256;
var safety = 0;
while voxels[offset + i] < LEAF_BIT {
let node_pos = data_start + voxels[offset + i];
side_len /= 2u;
let corner = pos / side_len;
pos -= corner * side_len;
let j = corner.x * 4 + corner.y * 2 + corner.z;
i = node_pos + j;
data_start = node_pos + 8;
if safety == 10 {
return 10u;
}
safety += 1;
}
return voxels[offset + i] & LEAF_MASK;
fn vec_to_dir(vec: vec3<u32>) -> u32 {
return vec.x * 4 + vec.y * 2 + vec.z * 1;
}
fn get_color(id: u32) -> vec4<f32> {
@@ -241,7 +202,7 @@ fn get_color(id: u32) -> vec4<f32> {
return vec4<f32>(0.5, 0.5, 0.5, 1.0);
}
case 2u: {
return vec4<f32>(0.8, 0.2, 0.2, 1.0);
return vec4<f32>(0.5, 1.0, 0.5, 1.0);
}
case 3u: {
return vec4<f32>(0.5, 0.5, 1.0, 0.5);

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

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

196
src/client/render/voxel/ray_oct/layout.rs Normal file
View File

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

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

@@ -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<View>,
bind_group_layout: wgpu::BindGroupLayout,
bind_group: wgpu::BindGroup,
voxel_groups: Storage<VoxelGroup>,
voxels: Storage<OctNode>,
global_lights: Storage<GlobalLight>,
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!("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<u32>) {
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);

1
src/client/render/voxel/ray_oct/render.wgsl
View File

@@ -41,7 +41,6 @@ fn fs_main(
) -> @location(0) vec4<f32> {
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
var pos = in.clip_position.xy / win_dim;
pos.y = 1.0 - pos.y;
return textureSample(t_diffuse, s_diffuse, pos);
}

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

@@ -1,212 +0,0 @@
// Vertex shader
struct GlobalLight {
dir: vec3<f32>,
};
struct VertexOutput {
@builtin(position) clip_position: vec4<f32>,
};
struct View {
transform: mat4x4<f32>,
width: u32,
height: u32,
zoom: f32,
};
struct VoxelGroup {
transform: mat4x4<f32>,
transform_inv: mat4x4<f32>,
dimensions: vec3<u32>,
offset: u32,
};
@group(0) @binding(0)
var<uniform> view: View;
@group(0) @binding(1)
var<storage, read> voxels: array<u32>;
@group(0) @binding(2)
var<storage, read> voxel_groups: array<VoxelGroup>;
@group(0) @binding(3)
var<storage, read> global_lights: array<GlobalLight>;
@vertex
fn vs_main(
@builtin(vertex_index) vi: u32,
@builtin(instance_index) ii: u32,
) -> VertexOutput {
var out: VertexOutput;
var pos = vec2<f32>(
f32(vi % 2u) * 2.0 - 1.0,
f32(vi / 2u) * 2.0 - 1.0,
) ;
out.clip_position = vec4<f32>(pos.x, pos.y, 0.0, 1.0);
return out;
}
// Fragment shader
@fragment
fn fs_main(
in: VertexOutput,
) -> @location(0) vec4<f32> {
// get position of the pixel; eye at origin, pixel on plane z = 1
let win_dim = vec2<f32>(f32(view.width), f32(view.height));
let aspect = win_dim.y / win_dim.x;
let pixel_pos = vec3<f32>(
(in.clip_position.xy / win_dim - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect),
view.zoom
);
// move to position in world
let pos = view.transform * vec4<f32>(pixel_pos, 1.0);
let dir = view.transform * vec4<f32>(normalize(pixel_pos), 0.0);
var color = trace_full(pos, dir);
let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0);
let sky_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
color.a = 1.0;
return color;
}
const ZERO3F = vec3<f32>(0.0);
const ZERO2F = vec2<f32>(0.0);
const DEPTH = 16u;
const FULL_ALPHA = 0.9999;
fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
let gi = 0;
let group = voxel_groups[gi];
if group.dimensions.x == 0 {
return vec4<f32>(0.0);
}
let dim_f = vec3<f32>(group.dimensions);
let dim_i = vec3<i32>(group.dimensions);
// transform so that group is at 0,0
let pos_start = (group.transform_inv * pos_view).xyz;
let dir = (group.transform_inv * dir_view).xyz;
let dir_if = sign(dir);
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 next_normal = vec3<f32>(0.0, 0.0, 0.0);
// find where ray intersects with group
let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
let pos_max = dir_uf * dim_f;
var pos = pos_start;
// time of intersection; x = td + p, solve for t
let t_min = (pos_min - pos) / dir;
let t_max = (pos_max - pos) / dir;
var t = 0.0;
if outside3f(pos, ZERO3F, dim_f) {
// 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 vec4<f32>(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<i32>(dir_if);
let dir_u = vec3<u32>((dir_i + vec3<i32>(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<f32>(1.0));
// color += vec4<f32>(new_color.xyz * vcolor.a, vcolor.a) * (1.0 - color.a);
if vcolor.a > .999 {
return vcolor;
}
}
return vcolor;
} else {
half_t_span *= 0.5;
let dir_idx = vec3<u32>(vec3<f32>(t) < t_center);
t_center += half_t_span * (1.0 - vec3<f32>(dir_idx * 2));
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<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 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);
}