rustgpu!!

.gitignore

@@ -1,3 +1,4 @@
 /target
+shader/target
 flamegraph.svg
 perf.data*

Cargo.toml

@@ -15,7 +15,7 @@ ndarray = "0.15.6"
 pollster = "0.3"
 rand = "0.8.5"
 simba = "0.8.1"
-wgpu = "22.0.0"
+wgpu = { version = "22.0.0", features = ["spirv"] }
 bevy_ecs = "0.13.2"
 bevy_derive = "0.13.2"
 winit = {version="0.30.1", features=["serde"]}
@@ -23,8 +23,13 @@ block-mesh = "0.2.0"
 rustc-hash = "2.0.0"
 simdnoise = { git = "https://github.com/valadaptive/rust-simd-noise", rev = "965175f" }
 
-[profile.dev]
-opt-level = 1
-[profile.dev.package."*"]
-opt-level = 3
+[build-dependencies]
+spirv-builder = { git = "https://github.com/Rust-GPU/rust-gpu", rev = "9a533a3" }
 
+# Compile build-dependencies in release mode with
+# the same settings as regular dependencies.
+[profile.release.build-override]
+opt-level = 3
+codegen-units = 16
+[profile.dev.build-override]
+opt-level = 3

build.rs

@@ -0,0 +1,9 @@
+use spirv_builder::{Capability, MetadataPrintout, SpirvBuilder};
+
+fn main() -> Result<(), Box<dyn std::error::Error>> {
+    SpirvBuilder::new("./shader", "spirv-unknown-spv1.5")
+        .capability(Capability::ImageQuery)
+        .print_metadata(MetadataPrintout::Full)
+        .build()?;
+    Ok(())
+}

readme.md

@@ -1,12 +0,0 @@
-# Voxel Game
-
-except right now it's just a raytraced sparse voxel octree renderer
-
-paper used: https://research.nvidia.com/publication/2010-02_efficient-sparse-voxel-octrees
-
-can render a lot of voxels right now, but a lot of work needs to be done for effects & multiple trees
-
-eg. (2^14)^3 voxels:
-![screenshot image here](screenshot.png)
-
-might take a bit to generate the terrain

rust-toolchain.toml

@@ -0,0 +1,7 @@
+[toolchain]
+channel = "nightly-2024-11-22"
+components = ["rust-src", "rustc-dev", "llvm-tools"]
+# commit_hash = b19329a37cedf2027517ae22c87cf201f93d776e
+
+# Whenever changing the nightly channel, update the commit hash above, and make
+# sure to change `REQUIRED_TOOLCHAIN` in `crates/rustc_codegen_spirv/build.rs` also.

shader/Cargo.lock

@@ -0,0 +1,113 @@
+# This file is automatically @generated by Cargo.
+# It is not intended for manual editing.
+version = 4
+
+[[package]]
+name = "autocfg"
+version = "1.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ace50bade8e6234aa140d9a2f552bbee1db4d353f69b8217bc503490fc1a9f26"
+
+[[package]]
+name = "bitflags"
+version = "1.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a"
+
+[[package]]
+name = "glam"
+version = "0.24.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b5418c17512bdf42730f9032c74e1ae39afc408745ebb2acf72fbc4691c17945"
+dependencies = [
+ "libm",
+]
+
+[[package]]
+name = "libm"
+version = "0.2.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8355be11b20d696c8f18f6cc018c4e372165b1fa8126cef092399c9951984ffa"
+
+[[package]]
+name = "num-traits"
+version = "0.2.19"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841"
+dependencies = [
+ "autocfg",
+ "libm",
+]
+
+[[package]]
+name = "proc-macro2"
+version = "1.0.94"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a31971752e70b8b2686d7e46ec17fb38dad4051d94024c88df49b667caea9c84"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "quote"
+version = "1.0.40"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1885c039570dc00dcb4ff087a89e185fd56bae234ddc7f056a945bf36467248d"
+dependencies = [
+ "proc-macro2",
+]
+
+[[package]]
+name = "shader"
+version = "0.1.0"
+dependencies = [
+ "spirv-std",
+]
+
+[[package]]
+name = "spirv-std"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68c3c0972a2df79abe2c8af2fe7f7937a9aa558b6a1f78fc5edf93f4d480d757"
+dependencies = [
+ "bitflags",
+ "glam",
+ "num-traits",
+ "spirv-std-macros",
+ "spirv-std-types",
+]
+
+[[package]]
+name = "spirv-std-macros"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "73f776bf9f2897ea7acff15d7753711fdf1693592bd7459a01c394262b1df45c"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "spirv-std-types",
+ "syn",
+]
+
+[[package]]
+name = "spirv-std-types"
+version = "0.9.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a73417b7d72d95b4995c840dceb4e3b4bcbad4ff7f35df9c1655b6826c18d3a9"
+
+[[package]]
+name = "syn"
+version = "1.0.109"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "unicode-ident"
+version = "1.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5a5f39404a5da50712a4c1eecf25e90dd62b613502b7e925fd4e4d19b5c96512"

shader/Cargo.toml

@@ -0,0 +1,10 @@
+[package]
+name = "shader"
+version = "0.1.0"
+edition = "2021"
+
+[lib]
+crate-type = ["dylib"]
+
+[dependencies]
+spirv-std = { version = "0.9" }

shader/src/compute.rs

@@ -0,0 +1,407 @@
+use core::mem::transmute;
+
+use super::fns::*;
+use spirv_std::glam::{
+    vec2, vec3, vec4, Mat3, Mat4, UVec2, UVec3, Vec3, Vec3Swizzles, Vec4, Vec4Swizzles,
+};
+use spirv_std::num_traits::{clamp, Pow};
+use spirv_std::{spirv, Image};
+
+#[repr(C, align(16))]
+#[derive(Copy, Clone)]
+pub struct View {
+    transform: Mat4,
+    zoom: f32,
+    chunk_scale: u32,
+    chunk_dist: u32,
+}
+
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct GlobalLight {
+    dir: Vec3,
+}
+
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct Chunk {
+    offset: u32,
+}
+
+#[repr(C)]
+#[derive(Copy, Clone)]
+pub struct PushConstants {
+    time: u32,
+}
+
+#[spirv(compute(threads(8, 8)))]
+pub fn main(
+    #[spirv(uniform, descriptor_set = 0, binding = 0)] view: &View,
+    #[spirv(push_constant)] consts: &PushConstants,
+    #[spirv(storage_buffer, descriptor_set = 0, binding = 1)] chunks: &[Chunk],
+    #[spirv(storage_buffer, descriptor_set = 0, binding = 2)] voxel_data: &[u32],
+    #[spirv(storage_buffer, descriptor_set = 0, binding = 3)] global_lights: &[GlobalLight],
+    #[spirv(descriptor_set = 0, binding = 4)] output: &Image!(2D, format = rgba8, sampled = false),
+    #[spirv(global_invocation_id)] cell: UVec3,
+) {
+    let view_dim: UVec2 = output.query_size();
+    // get position of the pixel; eye at origin, pixel on plane z = 1
+    if cell.x >= view_dim.x || cell.y >= view_dim.y {
+        return;
+    }
+    let view_dim_f = view_dim.as_vec2();
+    let aspect = view_dim_f.y / view_dim_f.x;
+    let pixel_pos = (cell.xy().as_vec2() / view_dim_f - vec2(0.5, 0.5)) * vec2(2.0, -2.0 * aspect);
+    let offset = Vec3::from_array([(1 << (view.chunk_scale - 1)) as f32; 3]);
+    let pos = view.transform * vec4(pixel_pos.x, pixel_pos.y, 1.0, 1.0) + offset.extend(0.0);
+    let dir = view.transform * pixel_pos.extend(view.zoom).normalize().extend(0.0);
+
+    let mut color = trace_full(
+        global_lights,
+        consts.time,
+        view,
+        voxel_data,
+        chunks,
+        pos,
+        dir,
+    );
+    let light_mult = clamp(
+        (-dir.xyz().dot(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 += (sky_color * (1.0 - color.w)).extend(1.0 - color.w);
+    color.w = 1.0;
+    unsafe {
+        output.write(cell.xy(), color);
+    }
+}
+
+const LEAF_BIT: u32 = 1u32 << 31;
+const LEAF_MASK: u32 = !LEAF_BIT;
+const MAX_HITS: u32 = 10;
+
+const FULL_ALPHA: f32 = 0.999;
+const EPSILON: f32 = 0.0000000001;
+const MAX_ITERS: u32 = 10000;
+// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
+// numbers are stored and the bit manipulation used
+const MAX_SCALE: u32 = 23;
+
+fn trace_full(
+    global_lights: &[GlobalLight],
+    time: u32,
+    view: &View,
+    voxel_data: &[u32],
+    chunks: &[Chunk],
+    pos_view: Vec4,
+    dir_view: Vec4,
+) -> Vec4 {
+    if voxel_data.len() == 1 {
+        return Vec4::ZERO;
+    }
+    let gi = 0;
+    let chunk = chunks[gi];
+    let side_len = 1u32 << view.chunk_scale;
+    let dimensions = UVec3::from_array([side_len; 3]);
+    let dim_f = dimensions.as_vec3();
+
+    let pos_start = pos_view.xyz();
+    let mut dir = dir_view.xyz();
+    if dir.x == 0.0 {
+        dir.x = EPSILON;
+    }
+    if dir.y == 0.0 {
+        dir.y = EPSILON;
+    }
+    if dir.z == 0.0 {
+        dir.z = EPSILON;
+    }
+
+    let dir_if = sign(dir);
+    let dir_uf = dir_if.max(Vec3::ZERO);
+
+    // find where ray intersects with group
+    // closest (min) and furthest (max) corners of cube relative to direction
+    let pos_min = (Vec3::ONE - dir_uf) * dim_f;
+    let pos_max = dir_uf * dim_f;
+    // time of intersection; x = td + p, solve for t
+    let t_min = (pos_min - pos_start) / dir;
+    let t_max = (pos_max - pos_start) / dir;
+    // time of entrance and exit of the cube
+    let t_start = t_min.max_element();
+    let t_end = t_max.min_element();
+    if t_end < t_start {
+        return Vec4::ZERO;
+    }
+    // axis of intersection
+    let axis = if t_start == t_min.x {
+        0
+    } else if t_start == t_min.y {
+        1
+    } else {
+        2
+    };
+    // time to move entire side length in each direction
+    let inc_t = (1.0 / dir).abs() * (side_len as f32);
+    let t = t_start.max(0.0);
+
+    let inv_dir_bits = 7 - vec_to_dir(dir_uf.as_uvec3());
+    let corner_adj = t_min - inc_t;
+
+    // calculate normals
+    let normals = Mat3::from_diagonal(dir_if);
+
+    let result = cast_ray(
+        voxel_data,
+        chunk.offset,
+        t,
+        axis,
+        inv_dir_bits,
+        inc_t,
+        corner_adj,
+    );
+
+    shade_ray(
+        global_lights,
+        time,
+        result,
+        pos_start,
+        dir_view.xyz(),
+        t_end,
+        normals,
+    )
+}
+
+fn shade_ray(
+    global_lights: &[GlobalLight],
+    time: u32,
+    result: RayResult,
+    pos_start: Vec3,
+    dir: Vec3,
+    t_end: f32,
+    normals: Mat3,
+) -> Vec4 {
+    let hits = result.hits;
+
+    let mut color = Vec4::ZERO;
+    for i in 0..result.len {
+        let hit = hits[i];
+        let id = hit.id;
+        let t = hit.t;
+        let axis = hit.axis;
+
+        let next_t = if i == result.len - 1 {
+            t_end
+        } else {
+            hits[i + 1].t
+        };
+
+        let mut pos = pos_start + dir * t;
+        let val = round(idx_vec(pos, axis)) - (idx_vec(dir, axis) < 0.0) as u32 as f32;
+        set_idx_vec(&mut pos, axis, val);
+        let normal = idx_mat(normals, axis);
+        let vcolor = shade(global_lights, time, id, pos, normal, dir, next_t - t);
+        color += vcolor * (1.0 - color.w);
+        if color.w > FULL_ALPHA {
+            break;
+        }
+    }
+    color
+}
+
+const AMBIENT: f32 = 0.2;
+const SPECULAR: f32 = 0.5;
+
+// returns premultiplied
+fn shade(
+    global_lights: &[GlobalLight],
+    time: u32,
+    id: u32,
+    pos: Vec3,
+    normal: Vec3,
+    dir_view: Vec3,
+    dist: f32,
+) -> Vec4 {
+    let mut color = Vec4::ZERO;
+    if id == 0 {
+        return color;
+    }
+    let random = randomf(pos.floor());
+    let random2 = randomf(pos.floor() + vec3(0.0001, 0.0001, 0.0001));
+    match id {
+        0 => {
+            color = Vec4::ZERO;
+        }
+        1 => {
+            color = (vec3(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8)).extend(1.0);
+        }
+        2 => {
+            color = (vec3(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9))
+                .extend(1.0);
+        }
+        3 => {
+            let fog = (dist / 64.0).min(1.0);
+            let a = 0.5;
+            let rand =
+                (sin(((time as f32) / 2000.0 + random) * core::f32::consts::TAU) + 1.0) / 2.0;
+            let rgb = vec3(0.5, 0.5, 1.0) * (rand * 0.2 + 0.8);
+            color = (rgb * (1.0 - fog * a)).extend(a + fog * (1.0 - a));
+        }
+        _ => (),
+    }
+    let light_color = Vec3::ONE;
+    let light_dir = global_lights[0].dir;
+
+    let diffuse = light_dir.dot(normal).max(0.0) * light_color;
+    let ambient = AMBIENT * light_color;
+    let spec_val = dir_view
+        .xyz()
+        .dot(reflect(-light_dir, normal))
+        .max(0.0)
+        .pow(32.0)
+        * SPECULAR;
+    let specular = spec_val * light_color;
+    let new_color = (ambient + diffuse + specular) * color.xyz();
+    let new_a = (color.w + spec_val).min(1.0);
+    (new_color * new_a).extend(new_a)
+}
+
+fn randomf(pos: Vec3) -> f32 {
+    fract(sin(pos.dot(vec3(12.9898, 78.233, 25.1279))) * 43758.545)
+}
+
+#[repr(C)]
+#[derive(Copy, Clone, Default)]
+struct RayHit {
+    t: f32,
+    id: u32,
+    axis: u32,
+}
+
+#[repr(C)]
+#[derive(Copy, Clone)]
+struct RayResult {
+    hits: [RayHit; MAX_HITS as usize],
+    len: usize,
+}
+
+fn vec_to_dir(vec: UVec3) -> u32 {
+    vec.x * 4 + vec.y * 2 + vec.z
+}
+
+fn cast_ray(
+    voxel_data: &[u32],
+    data_offset: u32,
+    t_start: f32,
+    axis_start: u32,
+    inv_dir_bits: u32,
+    inc_t: Vec3,
+    corner_adj: Vec3,
+) -> RayResult {
+    let mut hits: [RayHit; MAX_HITS as usize] = Default::default();
+    let mut depth = 0;
+    let mut min_alpha = 0.0;
+
+    let mut t = t_start;
+    let mut axis = axis_start;
+    let mut node_start = 0;
+    let mut scale = MAX_SCALE;
+    let mut scale_exp2 = 1.0;
+    let mut parents = [0u32; MAX_SCALE as usize];
+    let mut child = inv_dir_bits;
+    let mut vox_pos = Vec3::ONE;
+    let mut prev = 0;
+
+    for _ in 0..MAX_ITERS {
+        let t_corner = vox_pos * inc_t + corner_adj;
+        let node = voxel_data[(data_offset + node_start + (child ^ inv_dir_bits)) as usize];
+        if node >= LEAF_BIT {
+            // ignore consecutive identical leaves
+            if node != prev {
+                let id = node & LEAF_MASK;
+                hits[depth] = RayHit { t, id, axis };
+                min_alpha += min_alpha_for(id) * (1.0 - min_alpha);
+                depth += 1;
+                prev = node;
+                if depth == 10 || min_alpha >= FULL_ALPHA {
+                    break;
+                }
+            }
+
+            // move to next time point and determine which axis to move along
+            let t_next = t_corner + scale_exp2 * inc_t;
+            t = t_next.min_element();
+            axis = if t == t_next.x {
+                0
+            } else if t == t_next.y {
+                1
+            } else {
+                2
+            };
+            let move_dir = 4 >> axis;
+
+            // check if need to pop stack
+            if (child & move_dir) > 0 {
+                // calculate new scale; first differing bit after adding
+                let axis_pos = idx_vec(vox_pos, axis);
+                // AWARE
+                let differing = axis_pos.to_bits() ^ (axis_pos + scale_exp2).to_bits();
+                scale = ((differing as f32).to_bits() >> 23) - 127 - (23 - MAX_SCALE);
+                scale_exp2 = f32::from_bits((scale + 127 - MAX_SCALE) << 23);
+                if scale >= MAX_SCALE {
+                    break;
+                }
+
+                // restore & recalculate parent
+                let parent_info = parents[scale as usize - 1];
+                node_start = parent_info >> 3;
+                child = parent_info & 7;
+                let scale_vec = UVec3::from_array([scale + 23 - MAX_SCALE; 3]);
+                // remove bits lower than current scale
+                unsafe {
+                    vox_pos = transmute::<UVec3, Vec3>(
+                        (transmute::<Vec3, UVec3>(vox_pos) >> scale_vec) << scale_vec,
+                    );
+                }
+            }
+            // move to next child and voxel position
+            child += move_dir;
+            add_idx_vec(&mut vox_pos, axis, scale_exp2);
+        } else {
+            // push current node to stack
+            parents[scale as usize - 1] = (node_start << 3) + child;
+            scale -= 1;
+
+            // calculate child node vars
+            scale_exp2 *= 0.5;
+            child = 0;
+            let t_center = t_corner + scale_exp2 * inc_t;
+            if t > t_center.x {
+                vox_pos.x += scale_exp2;
+                child |= 4;
+            }
+            if t > t_center.y {
+                vox_pos.y += scale_exp2;
+                child |= 2;
+            }
+            if t > t_center.z {
+                vox_pos.z += scale_exp2;
+                child |= 1;
+            }
+            node_start = node;
+        }
+    }
+    RayResult { hits, len: depth }
+}
+
+fn min_alpha_for(id: u32) -> f32 {
+    match id {
+        0 => 0.0,
+        3 => 0.5,
+        _ => 1.0,
+    }
+}

shader/src/fns.rs

@@ -0,0 +1,94 @@
+#![allow(asm_sub_register)]
+
+use spirv_std::glam::{vec3, Mat3, Vec3};
+
+pub fn sign(vec: Vec3) -> Vec3 {
+    vec3(
+        if vec.x > 0.0 { 1.0 } else { -1.0 },
+        if vec.y > 0.0 { 1.0 } else { -1.0 },
+        if vec.z > 0.0 { 1.0 } else { -1.0 },
+    )
+}
+
+pub fn idx_vec(vec: Vec3, i: u32) -> f32 {
+    match i {
+        0 => vec.x,
+        1 => vec.y,
+        2 => vec.z,
+        _ => 0.0,
+    }
+}
+
+pub fn idx_mat(mat: Mat3, i: u32) -> Vec3 {
+    match i {
+        0 => mat.x_axis,
+        1 => mat.y_axis,
+        2 => mat.z_axis,
+        _ => Vec3::ZERO,
+    }
+}
+
+pub fn add_idx_vec(vec: &mut Vec3, i: u32, v: f32) {
+    match i {
+        0 => vec.x += v,
+        1 => vec.y += v,
+        2 => vec.z += v,
+        _ => (),
+    }
+}
+
+pub fn set_idx_vec(vec: &mut Vec3, i: u32, v: f32) {
+    match i {
+        0 => vec.x = v,
+        1 => vec.y = v,
+        2 => vec.z = v,
+        _ => (),
+    }
+}
+
+macro_rules! ext_import {
+    () => {
+        "%glsl = OpExtInstImport \"GLSL.std.450\""
+    };
+}
+
+macro_rules! float_op {
+    ($name: ident, $id: expr) => {
+        pub fn $name(x: f32) -> f32 {
+            let mut res: f32;
+            unsafe {
+                core::arch::asm!(
+                    ext_import!(),
+                    "%float = OpTypeFloat 32",
+                    concat!("{res} = OpExtInst %float %glsl ", $id, " {x}"),
+                    x = in(reg) x,
+                    res = out(reg) res,
+                );
+            }
+            res
+        }
+    };
+}
+
+float_op!(fract, 10);
+float_op!(sin, 13);
+float_op!(round, 1);
+
+pub fn reflect(x: Vec3, y: Vec3) -> Vec3 {
+    let mut o: Vec3 = Default::default();
+    unsafe {
+        core::arch::asm!(
+            ext_import!(),
+            "%float = OpTypeFloat 32",
+            "%vec = OpTypeVector %float 3",
+            "%x = OpLoad _ {x}",
+            "%y = OpLoad _ {y}",
+            "%o = OpExtInst %vec %glsl 71 %x %y",
+            "OpStore {o} %o",
+            x = in(reg) &x,
+            y = in(reg) &y,
+            o = in(reg) &mut o,
+        );
+    }
+    o
+}

shader/src/lib.rs

@@ -0,0 +1,8 @@
+#![no_std]
+#![allow(unexpected_cfgs)]
+#![feature(asm_experimental_arch)]
+
+pub mod render;
+pub mod compute;
+pub mod fns;
+

shader/src/render.rs

@@ -0,0 +1,24 @@
+use spirv_std::glam::{vec2, Vec2, Vec4};
+use spirv_std::{spirv, Image, Sampler};
+
+#[spirv(vertex)]
+pub fn vs_main(
+    #[spirv(vertex_index)] vi: u32,
+    #[spirv(position)] clip_position: &mut Vec4,
+    tex_pos: &mut Vec2,
+) {
+    let pos = vec2((vi % 2) as f32, (vi / 2) as f32);
+    *clip_position = (pos * 2.0 - 1.0).extend(0.0).extend(1.0);
+    *tex_pos = pos;
+    tex_pos.y = 1.0 - tex_pos.y;
+}
+
+#[spirv(fragment)]
+pub fn fs_main(
+    #[spirv(descriptor_set = 0, binding = 0)] texture: &Image!(2D, type=f32, sampled=true),
+    #[spirv(descriptor_set = 0, binding = 1)] sampler: &Sampler,
+    tex_pos: Vec2,
+    output: &mut Vec4,
+) {
+    *output = texture.sample(*sampler, tex_pos);
+}

src/client/render/mod.rs

@@ -53,7 +53,8 @@ impl<'a> Renderer<'a> {
                 required_features: wgpu::Features::PUSH_CONSTANTS
                     | wgpu::Features::TIMESTAMP_QUERY
                     | wgpu::Features::TIMESTAMP_QUERY_INSIDE_ENCODERS
-                    | wgpu::Features::TIMESTAMP_QUERY_INSIDE_PASSES,
+                    | wgpu::Features::TIMESTAMP_QUERY_INSIDE_PASSES
+                    | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
                 required_limits: wgpu::Limits {
                     max_storage_buffer_binding_size: buf_size,
                     max_buffer_size: buf_size as u64,

src/client/render/voxel/ray_oct/layout.rs

@@ -105,7 +105,7 @@ impl Layout {
                         binding: 4,
                         visibility: ShaderStages::COMPUTE,
                         ty: wgpu::BindingType::StorageTexture {
-                            access: wgpu::StorageTextureAccess::WriteOnly,
+                            access: wgpu::StorageTextureAccess::ReadWrite,
                             format: texture.format(),
                             view_dimension: wgpu::TextureViewDimension::D2,
                         },
@@ -171,20 +171,20 @@ impl Layout {
     pub fn render_pipeline(
         &self,
         device: &wgpu::Device,
-        shader: wgpu::ShaderModule,
+        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",
+                module: shader,
+                entry_point: "render::vs_main",
                 buffers: &[],
                 compilation_options: wgpu::PipelineCompilationOptions::default(),
             },
             fragment: Some(wgpu::FragmentState {
-                module: &shader,
-                entry_point: "fs_main",
+                module: shader,
+                entry_point: "render::fs_main",
                 targets: &[Some(wgpu::ColorTargetState {
                     format: self.format,
                     blend: Some(wgpu::BlendState::REPLACE),
@@ -221,7 +221,7 @@ impl Layout {
             label: Some("voxel"),
             layout: Some(&self.compute_pipeline_layout),
             module: shader,
-            entry_point: "main",
+            entry_point: "compute::main",
             compilation_options: Default::default(),
             cache: None,
         })

src/client/render/voxel/ray_oct/mod.rs

@@ -15,7 +15,6 @@ pub use color::*;
 use layout::Layout;
 use nalgebra::{Transform3, Translation3, Vector2};
 use std::collections::HashMap;
-use wgpu::include_wgsl;
 use {chunk::Chunk, view::View};
 
 pub struct VoxelPipeline {
@@ -27,21 +26,27 @@ pub struct VoxelPipeline {
     id_map: HashMap<Entity, (usize, Chunk)>,
 }
 
-const RENDER_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/render.wgsl");
-const COMPUTE_SHADER: wgpu::ShaderModuleDescriptor<'_> = include_wgsl!("shader/compute.wgsl");
-
 impl VoxelPipeline {
     pub fn new(device: &wgpu::Device, config: &wgpu::SurfaceConfiguration) -> Self {
         // shaders
 
         let layout = Layout::init(device, config);
 
+        let shader_source: Vec<u32> = include_bytes!(env!("shader.spv"))
+            .as_chunks::<4>()
+            .0
+            .iter()
+            .map(|bytes: &[u8; 4]| u32::from_le_bytes(*bytes))
+            .collect();
+        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("shaders"),
+            source: wgpu::ShaderSource::SpirV(std::borrow::Cow::from(shader_source)),
+        });
+
         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 render_pipeline = layout.render_pipeline(device, &shader);
 
         let compute_bind_group = layout.compute_bind_group(device);
-        let shader = device.create_shader_module(COMPUTE_SHADER);
         let compute_pipeline = layout.compute_pipeline(device, &shader);
 
         Self {
@@ -55,8 +60,9 @@ impl VoxelPipeline {
     }
 
     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);
+        // let shader = device.create_shader_module(COMPUTE_SHADER);
+        // self.compute_pipeline = self.layout.compute_pipeline(device, &shader);
+        println!("removed for now cause of rustgpu!");
     }
 
     pub fn add_group(
@@ -75,27 +81,28 @@ impl VoxelPipeline {
     }
 
     pub fn update_shader(&mut self, device: &wgpu::Device) {
-        let Ok(shader) = std::fs::read_to_string(
-            env!("CARGO_MANIFEST_DIR").to_owned()
-                + "/src/client/render/voxel/ray_oct/shader/compute.wgsl",
-        ) else {
-            println!("Failed to reload shader!");
-            return;
-        };
-        device.push_error_scope(wgpu::ErrorFilter::Validation);
-        let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
-            label: Some("Shader"),
-            source: wgpu::ShaderSource::Wgsl(shader.into()),
-        });
-        if pollster::block_on(device.pop_error_scope()).is_some() {
-            let comp_info = pollster::block_on(shader.get_compilation_info());
-            println!("Failed to compile shaders:");
-            for msg in comp_info.messages {
-                println!("{}", msg.message);
-            }
-        } else {
-            self.compute_pipeline = self.layout.compute_pipeline(device, &shader);
-        }
+        // let Ok(shader) = std::fs::read_to_string(
+        //     env!("CARGO_MANIFEST_DIR").to_owned()
+        //         + "/src/client/render/voxel/ray_oct/shader/compute.wgsl",
+        // ) else {
+        //     println!("Failed to reload shader!");
+        //     return;
+        // };
+        // device.push_error_scope(wgpu::ErrorFilter::Validation);
+        // let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+        //     label: Some("Shader"),
+        //     source: wgpu::ShaderSource::Wgsl(shader.into()),
+        // });
+        // if pollster::block_on(device.pop_error_scope()).is_some() {
+        //     let comp_info = pollster::block_on(shader.get_compilation_info());
+        //     println!("Failed to compile shaders:");
+        //     for msg in comp_info.messages {
+        //         println!("{}", msg.message);
+        //     }
+        // } else {
+        //     self.compute_pipeline = self.layout.compute_pipeline(device, &shader);
+        // }
+        println!("removed for now cause of rustgpu!");
     }
 
     pub fn add_chunk(

src/client/render/voxel/ray_oct/shader/compute.wgsl

@@ -1,303 +0,0 @@
-@group(0) @binding(0)
-var<uniform> view: View;
-@group(0) @binding(1)
-var<storage, read> chunks: array<Chunk>;
-@group(0) @binding(2)
-var<storage, read> voxel_data: array<u32>;
-@group(0) @binding(3)
-var<storage, read> global_lights: array<GlobalLight>;
-@group(0) @binding(4)
-var output: texture_storage_2d<rgba8unorm, write>;
-
-var<push_constant> time: u32;
-
-struct GlobalLight {
-    dir: vec3<f32>,
-};
-
-struct View {
-    transform: mat4x4<f32>,
-    zoom: f32,
-    chunk_scale: u32,
-    chunk_dist: u32,
-};
-
-struct Chunk {
-    offset: u32,
-};
-
-@compute
-@workgroup_size(8, 8, 1)
-fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
-    let view_dim = textureDimensions(output);
-    // get position of the pixel; eye at origin, pixel on plane z = 1
-    if cell.x >= view_dim.x || cell.y >= view_dim.y {
-        return;
-    }
-    let view_dim_f = vec2<f32>(view_dim);
-    let aspect = view_dim_f.y / view_dim_f.x;
-    let pixel_pos = vec2<f32>(
-        (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect)
-    );
-    let offset = vec3<f32>(f32(1u << (view.chunk_scale - 1)));
-    let pos = view.transform * vec4<f32>(pixel_pos, 1.0, 1.0) + vec4<f32>(offset, 0.0);
-    let dir = view.transform * vec4<f32>(normalize(vec3<f32>(pixel_pos, view.zoom)), 0.0);
-
-    var color = trace_full(pos, dir);
-    let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0);
-    let sun_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
-    let sky_bg = vec3<f32>(0.3, 0.6, 1.0);
-    let sky_color = sun_color + sky_bg * (1.0 - light_mult);
-    color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
-    color.a = 1.0;
-    textureStore(output, cell.xy, color);
-}
-
-const LEAF_BIT = 1u << 31u;
-const LEAF_MASK = ~LEAF_BIT;
-const MAX_HITS = 10;
-
-const ZERO3F = vec3<f32>(0.0);
-const ZERO2F = vec2<f32>(0.0);
-const FULL_ALPHA = 0.999;
-const EPSILON = 0.00000000001;
-const MAX_ITERS = 10000;
-// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
-// numbers are stored and the bit manipulation used
-const MAX_SCALE: u32 = 23;
-fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    if arrayLength(&voxel_data) == 1 {
-        return vec4<f32>(0.0);
-    }
-    let gi = 0;
-    let chunk = chunks[gi];
-    let side_len = 1u << view.chunk_scale;
-    let dimensions = vec3<u32>(side_len);
-    let dim_f = vec3<f32>(dimensions);
-
-    let pos_start = pos_view.xyz;
-    var dir = 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));
-
-    // find where ray intersects with group
-    // closest (min) and furthest (max) corners of cube relative to direction
-    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
-    let pos_max = dir_uf * dim_f;
-    // time of intersection; x = td + p, solve for t
-    let t_min = (pos_min - pos_start) / dir;
-    let t_max = (pos_max - pos_start) / dir;
-    // time of entrance and exit of the cube
-    let t_start = max(max(t_min.x, t_min.y), t_min.z);
-    let t_end = min(min(t_max.x, t_max.y), t_max.z);
-    if t_end < t_start { return vec4<f32>(0.0); }
-    // axis of intersection
-    let axis = select(select(2u, 1u, t_start == t_min.y), 0u, t_start == t_min.x);
-    // time to move entire side length in each direction
-    let inc_t = abs(1.0 / dir) * f32(side_len);
-    let t = max(0.0, t_start);
-
-    let inv_dir_bits = 7 - vec_to_dir(vec3<u32>(dir_uf));
-    let corner_adj = t_min - inc_t;
-
-    // calculate normals
-    var normals = mat3x3<f32>(
-        vec3<f32>(dir_if.x, 0.0, 0.0),
-        vec3<f32>(0.0, dir_if.y, 0.0),
-        vec3<f32>(0.0, 0.0, dir_if.z),
-    );
-
-    let result = cast_ray(chunk.offset, t, axis, inv_dir_bits, inc_t, corner_adj);
-    return shade_ray(result, pos_start, dir_view.xyz, t_end, normals);
-}
-
-fn shade_ray(result: RayResult, pos_start: vec3<f32>, dir: vec3<f32>, t_end: f32, normals: mat3x3<f32>) -> vec4<f32> {
-    var hits = result.hits;
-
-    var color = vec4<f32>(0.0);
-    for (var i = 0u; i < result.len; i += 1u) {
-        let hit = hits[i];
-        let id = hit.id;
-        let t = hit.t;
-        let axis = hit.axis;
-
-        let next_t = select(hits[i + 1].t, t_end, i == result.len - 1);
-
-        var pos = pos_start + dir * t;
-        pos[axis] = round(pos[axis]) - f32(dir[axis] < 0.0);
-        let normal = select(select(normals[0], normals[1], axis == 1), normals[2], axis == 2);
-        let vcolor = shade(id, pos, normal, dir, next_t - t);
-        color += vcolor * (1.0 - color.a);
-        if color.a > FULL_ALPHA { break; }
-    }
-    return color;
-}
-
-struct RayHit {
-    t: f32,
-    id: u32,
-    axis: u32,
-}
-
-struct RayResult {
-    hits: array<RayHit, MAX_HITS>,
-    len: u32,
-}
-
-fn cast_ray(
-    data_offset: u32, t_start: f32, axis_start: u32,
-    inv_dir_bits: u32, inc_t: vec3<f32>, corner_adj: vec3<f32>
-) -> RayResult {
-    var hits = array<RayHit, MAX_HITS>();
-    var depth = 0u;
-    var min_alpha = 0.0;
-
-    var t = t_start;
-    var axis = axis_start;
-    var node_start = 0u;
-    var scale = MAX_SCALE;
-    var scale_exp2 = 1.0;
-    var parents = array<u32, MAX_SCALE>();
-    var child = inv_dir_bits;
-    var vox_pos = vec3<f32>(1.0);
-    var prev = 0u;
-
-    var iters = 0;
-    loop {
-        if iters == MAX_ITERS { break; }
-        iters += 1;
-        let t_corner = vox_pos * inc_t + corner_adj;
-        let node = voxel_data[data_offset + node_start + (child ^ inv_dir_bits)];
-        if node >= LEAF_BIT {
-            // ignore consecutive identical leaves
-            if node != prev {
-                let id = node & LEAF_MASK;
-                hits[depth] = RayHit(t, id, axis);
-                min_alpha += min_alpha(id) * (1.0 - min_alpha);
-                depth += 1u;
-                prev = node;
-                if depth == 10 || min_alpha >= FULL_ALPHA { break; }
-            }
-
-            // move to next time point and determine which axis to move along
-            let t_next = t_corner + scale_exp2 * inc_t;
-            t = min(min(t_next.x, t_next.y), t_next.z);
-            axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z);
-            let move_dir = 4u >> axis;
-
-            // check if need to pop stack
-            if (child & move_dir) > 0 {
-                // calculate new scale; first differing bit after adding
-                let axis_pos = vox_pos[axis];
-                // AWARE
-                let differing = bitcast<u32>(axis_pos) ^ bitcast<u32>(axis_pos + scale_exp2);
-                scale = (bitcast<u32>(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE);
-                scale_exp2 = bitcast<f32>((scale + 127 - MAX_SCALE) << 23);
-                if scale >= MAX_SCALE { break; }
-
-                // restore & recalculate parent
-                let parent_info = parents[scale];
-                node_start = parent_info >> 3;
-                child = parent_info & 7;
-                let scale_vec = vec3<u32>(scale + 23 - MAX_SCALE);
-                // remove bits lower than current scale
-                vox_pos = bitcast<vec3<f32>>((bitcast<vec3<u32>>(vox_pos) >> scale_vec) << scale_vec);
-            }
-            // move to next child and voxel position
-            child += move_dir;
-            vox_pos[axis] += scale_exp2;
-        } else {
-            // push current node to stack
-            parents[scale] = (node_start << 3) + child;
-            scale -= 1u;
-
-            // calculate child node vars
-            scale_exp2 *= 0.5;
-            child = 0u;
-            let t_center = t_corner + scale_exp2 * inc_t;
-            if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; }
-            if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; }
-            if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; }
-            node_start = node;
-        }
-    }
-    return RayResult(hits, depth);
-}
-
-fn min_alpha(id: u32) -> f32 {
-    switch id {
-        case 0u: {return 0.0;}
-        case 3u: {return 0.5;}
-        default: {return 1.0;}
-    }
-}
-
-const AMBIENT: f32 = 0.2;
-const SPECULAR: f32 = 0.5;
-
-// returns premultiplied
-fn shade(id: u32, pos: vec3<f32>, normal: vec3<f32>, dir_view: vec3<f32>, dist: f32) -> vec4<f32> {
-    var color = vec4<f32>(0.0);
-    if id == 0 {
-        return color;
-    }
-    let random = random(floor(pos));
-    let random2 = random(floor(pos) + vec3<f32>(0.0001));
-    switch id {
-        case 0u: {
-            color = vec4<f32>(0.0);
-        }
-        case 1u: {
-            color = vec4<f32>(vec3<f32>(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8), 1.0);
-        }
-        case 2u: {
-            color = vec4<f32>(vec3<f32>(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9), 1.0);
-        }
-        case 3u: {
-            let fog = min(dist / 64.0, 1.0);
-            let a = 0.5;
-            let rand = (sin((f32(time) / 2000.0 + random) * 6.28) + 1.0) / 2.0;
-            let rgb = vec3<f32>(0.5, 0.5, 1.0) * (rand * 0.2 + 0.8);
-            color = vec4<f32>(rgb * (1.0 - fog * a), a + fog * (1.0 - a));
-        }
-        default: {}
-    }
-    let light_color = vec3<f32>(1.0);
-    let light_dir = global_lights[0].dir;
-
-    let diffuse = max(dot(light_dir, normal), 0.0) * light_color;
-    let ambient = AMBIENT * light_color;
-    let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR;
-    let specular = spec_val * light_color;
-    let new_color = (ambient + diffuse + specular) * color.xyz;
-    let new_a = min(color.a + spec_val, 1.0);
-    return vec4<f32>(new_color * new_a, new_a);
-}
-
-fn dir_to_vec(bits: u32) -> vec3<u32> {
-    return vec3<u32>(bits >> 2, (bits & 2) >> 1, bits & 1);
-}
-
-fn vec_to_dir(vec: vec3<u32>) -> u32 {
-    return vec.x * 4 + vec.y * 2 + vec.z * 1;
-}
-
-fn random(pos: vec3<f32>) -> f32 {
-    return fract(sin(dot(pos, vec3<f32>(12.9898, 78.233, 25.1279))) * 43758.5453123);
-}
-
-fn outside3f(v: vec3<f32>, low: vec3<f32>, high: vec3<f32>) -> bool {
-    return any(v < low) || any(v > high);
-}
-
-fn inside2f(v: vec2<f32>, low: vec2<f32>, high: vec2<f32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}
-
-fn inside3i(v: vec3<i32>, low: vec3<i32>, high: vec3<i32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}

src/client/render/voxel/ray_oct/shader/compute_single.wgsl

@@ -1,282 +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>,
-    zoom: f32,
-};
-
-struct VoxelGroup {
-    transform: mat4x4<f32>,
-    transform_inv: mat4x4<f32>,
-    scale: u32,
-    offset: u32,
-};
-
-@compute
-@workgroup_size(8, 8, 1)
-fn main(@builtin(global_invocation_id) cell: vec3<u32>) {
-    let view_dim = textureDimensions(output);
-    // get position of the pixel; eye at origin, pixel on plane z = 1
-    if cell.x >= view_dim.x || cell.y >= view_dim.y {
-        return;
-    }
-    let view_dim_f = vec2<f32>(view_dim);
-    let aspect = view_dim_f.y / view_dim_f.x;
-    let pixel_pos = vec2<f32>(
-        (vec2<f32>(cell.xy) / view_dim_f - vec2<f32>(0.5)) * vec2<f32>(2.0, -2.0 * aspect)
-    );
-    let pos = view.transform * vec4<f32>(pixel_pos, 1.0, 1.0);
-    let dir = view.transform * vec4<f32>(normalize(vec3<f32>(pixel_pos, view.zoom)), 0.0);
-
-    var color = trace_full(pos, dir);
-    let light_mult = clamp((-dot(dir.xyz, global_lights[0].dir) - 0.99) * 200.0, 0.0, 1.0);
-    let sun_color = light_mult * vec3<f32>(1.0, 1.0, 1.0);
-    let sky_bg = vec3<f32>(0.3, 0.6, 1.0);
-    let sky_color = sun_color + sky_bg * (1.0 - light_mult);
-    color += vec4<f32>(sky_color * (1.0 - color.a), 1.0 - color.a);
-    color.a = 1.0;
-    textureStore(output, cell.xy, color);
-}
-
-const LEAF_BIT = 1u << 31u;
-const LEAF_MASK = ~LEAF_BIT;
-
-const ZERO3F = vec3<f32>(0.0);
-const ZERO2F = vec2<f32>(0.0);
-const FULL_ALPHA = 0.999;
-const EPSILON = 0.00000000001;
-const MAX_ITERS = 10000;
-// NOTE: CANNOT GO HIGHER THAN 23 due to how floating point
-// numbers are stored and the bit manipulation used
-const MAX_SCALE: u32 = 13;
-const AMBIENT: f32 = 0.2;
-const SPECULAR: f32 = 0.5;
-
-fn trace_full(pos_view: vec4<f32>, dir_view: vec4<f32>) -> vec4<f32> {
-    let gi = 0;
-    let group = voxel_groups[gi];
-    if group.scale == 0 {
-        return vec4<f32>(0.0);
-    }
-    let dimensions = vec3<u32>(1u << group.scale);
-    let dim_f = vec3<f32>(dimensions);
-    let dim_i = vec3<i32>(dimensions);
-
-    // transform so that group is at 0,0
-    let pos_start = (group.transform_inv * pos_view).xyz;
-    var dir = (group.transform_inv * dir_view).xyz;
-    if dir.x == 0 {dir.x = EPSILON;}
-    if dir.y == 0 {dir.y = EPSILON;}
-    if dir.z == 0 {dir.z = EPSILON;}
-
-    let dir_if = sign(dir);
-    let dir_uf = max(dir_if, vec3<f32>(0.0));
-
-
-
-    // calculate normals
-    var normals = mat3x3<f32>(
-        (group.transform * vec4<f32>(dir_if.x, 0.0, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, dir_if.y, 0.0, 0.0)).xyz,
-        (group.transform * vec4<f32>(0.0, 0.0, dir_if.z, 0.0)).xyz,
-    );
-    var axis = 0u;
-
-    // find where ray intersects with group
-    let pos_min = (vec3<f32>(1.0) - dir_uf) * dim_f;
-    // time of intersection; x = td + p, solve for t
-    var t_min = (pos_min - pos_start) / dir;
-    if outside3f(pos_start, ZERO3F, dim_f) {
-        // points of intersection
-        let px = pos_start + t_min.x * dir;
-        let py = pos_start + t_min.y * dir;
-        let pz = pos_start + t_min.z * dir;
-
-        // check if point is in bounds
-        let hit = vec3<bool>(
-            inside2f(px.yz, ZERO2F, dim_f.yz),
-            inside2f(py.xz, ZERO2F, dim_f.xz),
-            inside2f(pz.xy, ZERO2F, dim_f.xy),
-        ) && (t_min > ZERO3F);
-        if !any(hit) {
-            return vec4<f32>(0.0);
-        }
-        axis = select(select(2u, 1u, hit.y), 0u, hit.x);
-    }
-    let t_mult = f32(1u << (MAX_SCALE - group.scale));
-    t_min *= t_mult;
-    let full = f32(1u << MAX_SCALE);
-    // time to move entire side length in each direction
-    let inc_t = abs(1.0 / dir) * full;
-    let t_offset = max(max(t_min.x, t_min.y), t_min.z);
-    var t = max(0.0, t_offset);
-
-    let dir_i = vec3<i32>(dir_if);
-    let dir_u = vec3<u32>(dir_uf);
-    let dir_bits = vec_to_dir(dir_u);
-    let inv_dir_bits = 7 - dir_bits;
-
-    var node_start = 1u;
-    var scale = MAX_SCALE - 1;
-    var scale_exp2 = 0.5;
-    var color = vec4<f32>(0.0);
-    var parents = array<u32, MAX_SCALE>();
-    var prev = LEAF_BIT;
-    var old_t = t / t_mult;
-
-    var child = 0u;
-    var vox_pos = vec3<f32>(1.0);
-    let t_center = t_min + scale_exp2 * inc_t;
-    if t > t_center.x { vox_pos.x = 1.5; child |= 4u; }
-    if t > t_center.y { vox_pos.y = 1.5; child |= 2u; }
-    if t > t_center.z { vox_pos.z = 1.5; child |= 1u; }
-    let min_adj = t_min - inc_t;
-
-    var iters = 0;
-    loop {
-        if iters == MAX_ITERS {
-            return vec4<f32>(1.0, 0.0, 1.0, 1.0);
-        }
-        iters += 1;
-        let t_corner = vox_pos * inc_t + min_adj;
-        let node = voxels[group.offset + node_start + (child ^ inv_dir_bits)];
-        if node >= LEAF_BIT {
-            if node != prev {
-                if node != LEAF_BIT {
-                    let real_t = t / t_mult;
-                    let dist = real_t - old_t;
-                    old_t = real_t;
-                    let filt = min(dist / 64.0, 1.0);
-                    if prev == LEAF_BIT + 3 {
-                        color.a += filt * (1.0 - color.a);
-                        if color.a > FULL_ALPHA { break; }
-                    }
-                    var pos = (pos_view + dir_view * real_t).xyz;
-                    pos[axis] = round(pos[axis]) - (1.0 - dir_uf[axis]);
-                    let vcolor = get_color(node & LEAF_MASK, pos);
-                    var normal = normals[axis];
-                    let light_color = vec3<f32>(1.0);
-                    let light_dir = global_lights[0].dir;
-
-                    let diffuse = max(dot(light_dir, normal), 0.0) * light_color;
-                    let ambient = AMBIENT * light_color;
-                    let spec_val = pow(max(dot(dir_view.xyz, reflect(-light_dir, normal)), 0.0), 32.0) * SPECULAR;
-                    let specular = spec_val * light_color;
-                    let new_color = (ambient + diffuse + specular) * vcolor.xyz;
-                    let new_a = min(vcolor.a + spec_val, 1.0);
-
-                    color += vec4<f32>(new_color.xyz * new_a, new_a) * (1.0 - color.a);
-                    if color.a > FULL_ALPHA { break; }
-                }
-                prev = node;
-            }
-
-            // move to next time point and determine which axis to move along
-            let t_next = t_corner + scale_exp2 * inc_t;
-            t = min(min(t_next.x, t_next.y), t_next.z);
-            axis = select(select(0u, 1u, t == t_next.y), 2u, t == t_next.z);
-            let move_dir = 4u >> axis;
-
-            // check if need to pop stack
-            if (child & move_dir) > 0 {
-                // calculate new scale; first differing bit after adding
-                let axis_pos = vox_pos[axis];
-                // AWARE
-                let differing = bitcast<u32>(axis_pos) ^ bitcast<u32>(axis_pos + scale_exp2);
-                scale = (bitcast<u32>(f32(differing)) >> 23) - 127 - (23 - MAX_SCALE);
-                scale_exp2 = bitcast<f32>((scale + 127 - MAX_SCALE) << 23);
-                if scale >= MAX_SCALE { break; }
-
-                // restore & recalculate parent
-                let parent_info = parents[scale];
-                node_start = parent_info >> 3;
-                child = parent_info & 7;
-                let scale_vec = vec3<u32>(scale + 23 - MAX_SCALE);
-                // remove bits lower than current scale
-                vox_pos = bitcast<vec3<f32>>((bitcast<vec3<u32>>(vox_pos) >> scale_vec) << scale_vec);
-            }
-            // move to next child and voxel position
-            child += move_dir;
-            vox_pos[axis] += scale_exp2;
-        } else {
-            // push current node to stack
-            parents[scale] = (node_start << 3) + child;
-            scale -= 1u;
-
-            // calculate child node vars
-            scale_exp2 *= 0.5;
-            child = 0u;
-            let t_center = t_corner + scale_exp2 * inc_t;
-            if t > t_center.x { vox_pos.x += scale_exp2; child |= 4u; }
-            if t > t_center.y { vox_pos.y += scale_exp2; child |= 2u; }
-            if t > t_center.z { vox_pos.z += scale_exp2; child |= 1u; }
-            node_start = node;
-        }
-    }
-    // let fog = min(t / t_mult / 1000.0, 1.0);
-    // return vec4<f32>(color.xyz * (1.0 - fog) + vec3<f32>(fog), color.a * (1.0 - fog) + fog);
-    // return vec4<f32>(f32(iters) / f32(MAX_ITERS), 0.0, 0.0, 1.0);
-    return color;
-}
-
-fn dir_to_vec(bits: u32) -> vec3<u32> {
-    return vec3<u32>(bits >> 2, (bits & 2) >> 1, bits & 1);
-}
-
-fn vec_to_dir(vec: vec3<u32>) -> u32 {
-    return vec.x * 4 + vec.y * 2 + vec.z * 1;
-}
-
-fn get_color(id: u32, pos: vec3<f32>) -> vec4<f32> {
-    let random = random(floor(pos));
-    let random2 = random(floor(pos) + vec3<f32>(0.0001));
-    switch id {
-        case 0u: {
-            return vec4<f32>(0.0);
-        }
-        case 1u: {
-            let color = vec3<f32>(0.5, 0.5, 0.5 + random * 0.2) * (random2 * 0.4 + 0.8);
-            return vec4<f32>(color, 1.0);
-        }
-        case 2u: {
-            let color = vec3<f32>(0.4 + random * 0.2, 0.9, 0.4 + random * 0.2) * (random2 * 0.2 + 0.9);
-            return vec4<f32>(color, 1.0);
-        }
-        case 3u: {
-            let color = vec3<f32>(0.5, 0.5, 1.0) * (random2 * 0.2 + 0.8);
-            return vec4<f32>(color, 0.5);
-        }
-        default: {
-            return vec4<f32>(1.0, 0.0, 0.0, 1.0);
-        }
-    }
-}
-
-fn random(pos: vec3<f32>) -> f32 {
-    return fract(sin(dot(pos, vec3<f32>(12.9898, 78.233, 25.1279))) * 43758.5453123);
-}
-
-fn outside3f(v: vec3<f32>, low: vec3<f32>, high: vec3<f32>) -> bool {
-    return any(v < low) || any(v > high);
-}
-
-fn inside2f(v: vec2<f32>, low: vec2<f32>, high: vec2<f32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}
-
-fn inside3i(v: vec3<i32>, low: vec3<i32>, high: vec3<i32>) -> bool {
-    return all(v >= low) && all(v <= high);
-}

src/client/render/voxel/ray_oct/shader/render.wgsl

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

src/common/component/chunk/mod.rs

@@ -8,7 +8,7 @@ use bevy_derive::{Deref, DerefMut};
 use bevy_ecs::{bundle::Bundle, component::Component, entity::Entity, system::Resource};
 use nalgebra::Vector3;
 
-pub const SCALE: u32 = 14;
+pub const SCALE: u32 = 10;
 pub const SIDE_LENGTH: usize = 2usize.pow(SCALE);
 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);

src/main.rs

@@ -1,4 +1,5 @@
 #![allow(clippy::type_complexity)]
+#![feature(slice_as_chunks)]
 
 use client::ClientApp;
 use winit::event_loop::EventLoop;