diff --git a/Cargo.lock b/Cargo.lock index 222e6a1a..2543d7c7 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -41,6 +41,15 @@ dependencies = [ "version_check", ] +[[package]] +name = "aho-corasick" +version = "0.7.18" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "1e37cfd5e7657ada45f742d6e99ca5788580b5c529dc78faf11ece6dc702656f" +dependencies = [ + "memchr", +] + [[package]] name = "ansi_term" version = "0.12.1" @@ -144,12 +153,48 @@ version = "0.1.6" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" +[[package]] +name = "block-buffer" +version = "0.7.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "c0940dc441f31689269e10ac70eb1002a3a1d3ad1390e030043662eb7fe4688b" +dependencies = [ + "block-padding", + "byte-tools", + "byteorder", + "generic-array", +] + +[[package]] +name = "block-padding" +version = "0.1.5" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "fa79dedbb091f449f1f39e53edf88d5dbe95f895dae6135a8d7b881fb5af73f5" +dependencies = [ + "byte-tools", +] + +[[package]] +name = "bstr" +version = "0.2.17" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "ba3569f383e8f1598449f1a423e72e99569137b47740b1da11ef19af3d5c3223" +dependencies = [ + "memchr", +] + [[package]] name = "bumpalo" version = "3.8.0" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "8f1e260c3a9040a7c19a12468758f4c16f31a81a1fe087482be9570ec864bb6c" +[[package]] +name = "byte-tools" +version = "0.3.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "e3b5ca7a04898ad4bcd41c90c5285445ff5b791899bb1b0abdd2a2aa791211d7" + [[package]] name = "bytemuck" version = "1.7.3" @@ -233,6 +278,28 @@ dependencies = [ "winapi", ] +[[package]] +name = "chrono-tz" +version = "0.6.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "58549f1842da3080ce63002102d5bc954c7bc843d4f47818e642abdc36253552" +dependencies = [ + "chrono", + "chrono-tz-build", + "phf", +] + +[[package]] +name = "chrono-tz-build" +version = "0.0.2" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "db058d493fb2f65f41861bfed7e3fe6335264a9f0f92710cab5bdf01fef09069" +dependencies = [ + "parse-zoneinfo", + "phf", + "phf_codegen", +] + [[package]] name = "cloudabi" version = "0.0.3" @@ -512,6 +579,21 @@ dependencies = [ "syn", ] +[[package]] +name = "deunicode" +version = "0.4.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "850878694b7933ca4c9569d30a34b55031b9b139ee1fc7b94a527c4ef960d690" + +[[package]] +name = "digest" +version = "0.8.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "f3d0c8c8752312f9713efd397ff63acb9f85585afbf179282e720e7704954dd5" +dependencies = [ + "generic-array", +] + [[package]] name = "dispatch" version = "0.2.0" @@ -540,6 +622,7 @@ dependencies = [ "dotrix_primitives", "dotrix_sky", "dotrix_terrain", + "dotrix_voxel", "noise", "rand 0.8.4", ] @@ -629,6 +712,23 @@ dependencies = [ "wgpu", ] +[[package]] +name = "dotrix_voxel" +version = "0.5.0" +dependencies = [ + "bytemuck", + "dotrix_core", + "dotrix_math", + "dotrix_pbr", + "dotrix_primitives", + "futures", + "num", + "rand 0.8.4", + "rayon", + "tera", + "wgpu", +] + [[package]] name = "downcast-rs" version = "1.2.0" @@ -671,6 +771,12 @@ dependencies = [ "nohash-hasher", ] +[[package]] +name = "fake-simd" +version = "0.1.2" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "e88a8acf291dafb59c2d96e8f59828f3838bb1a70398823ade51a84de6a6deed" + [[package]] name = "fnv" version = "1.0.7" @@ -784,6 +890,15 @@ dependencies = [ "byteorder", ] +[[package]] +name = "generic-array" +version = "0.12.4" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "ffdf9f34f1447443d37393cc6c2b8313aebddcd96906caf34e54c68d8e57d7bd" +dependencies = [ + "typenum", +] + [[package]] name = "getrandom" version = "0.1.16" @@ -816,6 +931,30 @@ dependencies = [ "weezl", ] +[[package]] +name = "globset" +version = "0.4.8" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "10463d9ff00a2a068db14231982f5132edebad0d7660cd956a1c30292dbcbfbd" +dependencies = [ + "aho-corasick", + "bstr", + "fnv", + "log", + "regex", +] + +[[package]] +name = "globwalk" +version = "0.8.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "93e3af942408868f6934a7b85134a3230832b9977cf66125df2f9edcfce4ddcc" +dependencies = [ + "bitflags", + "ignore", + "walkdir", +] + [[package]] name = "glow" version = "0.11.1" @@ -928,12 +1067,36 @@ version = "0.2.1" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" +[[package]] +name = "humansize" +version = "1.1.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "02296996cb8796d7c6e3bc2d9211b7802812d36999a51bb754123ead7d37d026" + [[package]] name = "ident_case" version = "1.0.1" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39" +[[package]] +name = "ignore" +version = "0.4.18" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "713f1b139373f96a2e0ce3ac931cd01ee973c3c5dd7c40c0c2efe96ad2b6751d" +dependencies = [ + "crossbeam-utils", + "globset", + "lazy_static", + "log", + "memchr", + "regex", + "same-file", + "thread_local", + "walkdir", + "winapi-util", +] + [[package]] name = "image" version = "0.23.14" @@ -1077,6 +1240,12 @@ dependencies = [ "libc", ] +[[package]] +name = "maplit" +version = "1.0.2" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "3e2e65a1a2e43cfcb47a895c4c8b10d1f4a61097f9f254f183aee60cad9c651d" + [[package]] name = "matchers" version = "0.0.1" @@ -1419,6 +1588,12 @@ version = "1.9.0" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "da32515d9f6e6e489d7bc9d84c71b060db7247dc035bbe44eac88cf87486d8d5" +[[package]] +name = "opaque-debug" +version = "0.2.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "2839e79665f131bdb5782e51f2c6c9599c133c6098982a54c794358bf432529c" + [[package]] name = "owned_ttf_parser" version = "0.13.2" @@ -1453,12 +1628,103 @@ dependencies = [ "winapi", ] +[[package]] +name = "parse-zoneinfo" +version = "0.3.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "c705f256449c60da65e11ff6626e0c16a0a0b96aaa348de61376b249bc340f41" +dependencies = [ + "regex", +] + [[package]] name = "percent-encoding" version = "2.1.0" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "d4fd5641d01c8f18a23da7b6fe29298ff4b55afcccdf78973b24cf3175fee32e" +[[package]] +name = "pest" +version = "2.1.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "10f4872ae94d7b90ae48754df22fd42ad52ce740b8f370b03da4835417403e53" +dependencies = [ + "ucd-trie", +] + +[[package]] +name = "pest_derive" +version = "2.1.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "833d1ae558dc601e9a60366421196a8d94bc0ac980476d0b67e1d0988d72b2d0" +dependencies = [ + "pest", + "pest_generator", +] + +[[package]] +name = "pest_generator" +version = "2.1.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "99b8db626e31e5b81787b9783425769681b347011cc59471e33ea46d2ea0cf55" +dependencies = [ + "pest", + "pest_meta", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "pest_meta" +version = "2.1.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "54be6e404f5317079812fc8f9f5279de376d8856929e21c184ecf6bbd692a11d" +dependencies = [ + "maplit", + "pest", + "sha-1", +] + +[[package]] +name = "phf" +version = "0.10.1" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "fabbf1ead8a5bcbc20f5f8b939ee3f5b0f6f281b6ad3468b84656b658b455259" +dependencies = [ + "phf_shared", +] + +[[package]] +name = "phf_codegen" +version = "0.10.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "4fb1c3a8bc4dd4e5cfce29b44ffc14bedd2ee294559a294e2a4d4c9e9a6a13cd" +dependencies = [ + "phf_generator", + "phf_shared", +] + +[[package]] +name = "phf_generator" +version = "0.10.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "5d5285893bb5eb82e6aaf5d59ee909a06a16737a8970984dd7746ba9283498d6" +dependencies = [ + "phf_shared", + "rand 0.8.4", +] + +[[package]] +name = "phf_shared" +version = "0.10.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "b6796ad771acdc0123d2a88dc428b5e38ef24456743ddb1744ed628f9815c096" +dependencies = [ + "siphasher", + "uncased", +] + [[package]] name = "pin-project-lite" version = "0.2.7" @@ -1795,6 +2061,8 @@ version = "1.5.4" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "d07a8629359eb56f1e2fb1652bb04212c072a87ba68546a04065d525673ac461" dependencies = [ + "aho-corasick", + "memchr", "regex-syntax", ] @@ -1842,6 +2110,15 @@ version = "1.0.9" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "73b4b750c782965c211b42f022f59af1fbceabdd026623714f104152f1ec149f" +[[package]] +name = "same-file" +version = "1.0.6" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502" +dependencies = [ + "winapi-util", +] + [[package]] name = "scoped-tls" version = "1.0.0" @@ -1891,6 +2168,18 @@ dependencies = [ "serde", ] +[[package]] +name = "sha-1" +version = "0.8.2" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "f7d94d0bede923b3cea61f3f1ff57ff8cdfd77b400fb8f9998949e0cf04163df" +dependencies = [ + "block-buffer", + "digest", + "fake-simd", + "opaque-debug", +] + [[package]] name = "shaderc" version = "0.7.3" @@ -1920,6 +2209,12 @@ dependencies = [ "lazy_static", ] +[[package]] +name = "siphasher" +version = "0.3.9" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "a86232ab60fa71287d7f2ddae4a7073f6b7aac33631c3015abb556f08c6d0a3e" + [[package]] name = "slab" version = "0.4.5" @@ -1935,6 +2230,15 @@ dependencies = [ "version_check", ] +[[package]] +name = "slug" +version = "0.1.4" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "b3bc762e6a4b6c6fcaade73e77f9ebc6991b676f88bb2358bddb56560f073373" +dependencies = [ + "deunicode", +] + [[package]] name = "smallvec" version = "1.7.0" @@ -1987,6 +2291,28 @@ dependencies = [ "unicode-xid", ] +[[package]] +name = "tera" +version = "1.15.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "d3cac831b615c25bcef632d1cabf864fa05813baad3d526829db18eb70e8b58d" +dependencies = [ + "chrono", + "chrono-tz", + "globwalk", + "humansize", + "lazy_static", + "percent-encoding", + "pest", + "pest_derive", + "rand 0.8.4", + "regex", + "serde", + "serde_json", + "slug", + "unic-segment", +] + [[package]] name = "termcolor" version = "1.1.2" @@ -2125,6 +2451,77 @@ version = "0.13.4" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "76dacc724328b3d5e2ed67f9e30cdb56893a34ab239032502cc8f19f8dae4bbc" +[[package]] +name = "typenum" +version = "1.15.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "dcf81ac59edc17cc8697ff311e8f5ef2d99fcbd9817b34cec66f90b6c3dfd987" + +[[package]] +name = "ucd-trie" +version = "0.1.3" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "56dee185309b50d1f11bfedef0fe6d036842e3fb77413abef29f8f8d1c5d4c1c" + +[[package]] +name = "uncased" +version = "0.9.6" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "5baeed7327e25054889b9bd4f975f32e5f4c5d434042d59ab6cd4142c0a76ed0" +dependencies = [ + "version_check", +] + +[[package]] +name = "unic-char-property" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "a8c57a407d9b6fa02b4795eb81c5b6652060a15a7903ea981f3d723e6c0be221" +dependencies = [ + "unic-char-range", +] + +[[package]] +name = "unic-char-range" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "0398022d5f700414f6b899e10b8348231abf9173fa93144cbc1a43b9793c1fbc" + +[[package]] +name = "unic-common" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "80d7ff825a6a654ee85a63e80f92f054f904f21e7d12da4e22f9834a4aaa35bc" + +[[package]] +name = "unic-segment" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "e4ed5d26be57f84f176157270c112ef57b86debac9cd21daaabbe56db0f88f23" +dependencies = [ + "unic-ucd-segment", +] + +[[package]] +name = "unic-ucd-segment" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "2079c122a62205b421f499da10f3ee0f7697f012f55b675e002483c73ea34700" +dependencies = [ + "unic-char-property", + "unic-char-range", + "unic-ucd-version", +] + +[[package]] +name = "unic-ucd-version" +version = "0.9.0" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "96bd2f2237fe450fcd0a1d2f5f4e91711124f7857ba2e964247776ebeeb7b0c4" +dependencies = [ + "unic-common", +] + [[package]] name = "unicode-width" version = "0.1.9" @@ -2143,6 +2540,17 @@ version = "0.9.3" source = "registry+/~https://github.com/rust-lang/crates.io-index" checksum = "5fecdca9a5291cc2b8dcf7dc02453fee791a280f3743cb0905f8822ae463b3fe" +[[package]] +name = "walkdir" +version = "2.3.2" +source = "registry+/~https://github.com/rust-lang/crates.io-index" +checksum = "808cf2735cd4b6866113f648b791c6adc5714537bc222d9347bb203386ffda56" +dependencies = [ + "same-file", + "winapi", + "winapi-util", +] + [[package]] name = "wasi" version = "0.9.0+wasi-snapshot-preview1" diff --git a/Cargo.toml b/Cargo.toml index b5fe3e19..e65d9a35 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -16,7 +16,7 @@ name = "dotrix" path = "src/lib.rs" [features] -default = [ "egui", "overlay", "pbr", "primitives", "sky", "terrain" ] +default = [ "egui", "overlay", "pbr", "primitives", "sky", "terrain", "voxel" ] overlay = ["dotrix_overlay"] sky = ["dotrix_sky"] @@ -24,6 +24,7 @@ pbr = ["dotrix_pbr"] primitives = ["dotrix_primitives"] egui = ["dotrix_egui"] terrain = ["dotrix_terrain"] +voxel = ["dotrix_voxel"] [workspace] members = [ @@ -34,6 +35,7 @@ members = [ "dotrix_primitives", "dotrix_sky", "dotrix_terrain", + "dotrix_voxel", ] [dependencies.dotrix_core] @@ -76,6 +78,11 @@ version = "0.5" path = "./dotrix_terrain" optional = true +[dependencies.dotrix_voxel] +version = "0.5" +path = "./dotrix_voxel" +optional = true + [dev-dependencies.rand] version = "0.8" @@ -125,3 +132,8 @@ path = "examples/msaa/main.rs" [[example]] name = "normal_map" path = "examples/normal_map/main.rs" + +[[example]] +name = "voxel_sdf" +path = "examples/voxel_sdf/main.rs" +required-features = ["voxel"] diff --git a/assets/textures/Bricks075B/Bricks075B_1K_AmbientOcclusion.jpg b/assets/textures/Bricks075B/Bricks075B_1K_AmbientOcclusion.jpg new file mode 100644 index 00000000..b6fdb7c7 Binary files /dev/null and b/assets/textures/Bricks075B/Bricks075B_1K_AmbientOcclusion.jpg differ diff --git a/assets/textures/Bricks075B/Bricks075B_1K_Color.jpg b/assets/textures/Bricks075B/Bricks075B_1K_Color.jpg new file mode 100644 index 00000000..36d6f273 Binary files /dev/null and b/assets/textures/Bricks075B/Bricks075B_1K_Color.jpg differ diff --git a/assets/textures/Bricks075B/Bricks075B_1K_NormalDX.jpg b/assets/textures/Bricks075B/Bricks075B_1K_NormalDX.jpg new file mode 100644 index 00000000..8b70ea08 Binary files /dev/null and b/assets/textures/Bricks075B/Bricks075B_1K_NormalDX.jpg differ diff --git a/assets/textures/Bricks075B/Bricks075B_1K_Roughness.jpg b/assets/textures/Bricks075B/Bricks075B_1K_Roughness.jpg new file mode 100644 index 00000000..e5bf8169 Binary files /dev/null and b/assets/textures/Bricks075B/Bricks075B_1K_Roughness.jpg differ diff --git a/assets/textures/Bricks075B/license.md b/assets/textures/Bricks075B/license.md new file mode 100644 index 00000000..d0c4f566 --- /dev/null +++ b/assets/textures/Bricks075B/license.md @@ -0,0 +1,27 @@ +# Source + +https://ambientcg.com/view?id=Bricks075B + +# Licensing + +All the PBR materials, brushes, photos and 3D models which are offered for download on ambientCG are provided under the Creative Commons CC0 1.0 Universal License. + +Creative Commons - CC0 1.0 Universal + +# Summary + +The Creative Commons CC0 license gives you the freedom to use these assets… + + … in your own creations of any kind for free. + … even in situations that require them to be redistributed as individual files (for example as part of an open source video game or tool). + … for commercial projects. + +There is no requirement to give credit but it would of course be appreciated. + +# Giving credit + +If you wish to give credit you can do so using this text: + +``` +Contains assets from ambientCG.com, licensed under CC0 1.0 Universal. +``` diff --git a/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_AmbientOcclusion.png b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_AmbientOcclusion.png new file mode 100644 index 00000000..9b0eae7d Binary files /dev/null and b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_AmbientOcclusion.png differ diff --git a/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Color.png b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Color.png new file mode 100644 index 00000000..a4aa4e23 Binary files /dev/null and b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Color.png differ diff --git a/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_NormalDX.png b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_NormalDX.png new file mode 100644 index 00000000..b892ed29 Binary files /dev/null and b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_NormalDX.png differ diff --git a/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Roughness.png b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Roughness.png new file mode 100644 index 00000000..5788a7ab Binary files /dev/null and b/assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Roughness.png differ diff --git a/assets/textures/PaintedPlaster010/license.md b/assets/textures/PaintedPlaster010/license.md new file mode 100644 index 00000000..280c0cac --- /dev/null +++ b/assets/textures/PaintedPlaster010/license.md @@ -0,0 +1,27 @@ +# Source + +https://ambientcg.com/view?id=PaintedPlaster010 + +# Licensing + +All the PBR materials, brushes, photos and 3D models which are offered for download on ambientCG are provided under the Creative Commons CC0 1.0 Universal License. + +Creative Commons - CC0 1.0 Universal + +# Summary + +The Creative Commons CC0 license gives you the freedom to use these assets… + + … in your own creations of any kind for free. + … even in situations that require them to be redistributed as individual files (for example as part of an open source video game or tool). + … for commercial projects. + +There is no requirement to give credit but it would of course be appreciated. + +# Giving credit + +If you wish to give credit you can do so using this text: + +``` +Contains assets from ambientCG.com, licensed under CC0 1.0 Universal. +``` diff --git a/assets/textures/PavingStones113/PavingStones113_1K_AmbientOcclusion.jpg b/assets/textures/PavingStones113/PavingStones113_1K_AmbientOcclusion.jpg new file mode 100644 index 00000000..2e34f155 Binary files /dev/null and b/assets/textures/PavingStones113/PavingStones113_1K_AmbientOcclusion.jpg differ diff --git a/assets/textures/PavingStones113/PavingStones113_1K_Color.jpg b/assets/textures/PavingStones113/PavingStones113_1K_Color.jpg new file mode 100644 index 00000000..737c9a09 Binary files /dev/null and b/assets/textures/PavingStones113/PavingStones113_1K_Color.jpg differ diff --git a/assets/textures/PavingStones113/PavingStones113_1K_NormalDX.jpg b/assets/textures/PavingStones113/PavingStones113_1K_NormalDX.jpg new file mode 100644 index 00000000..19803966 Binary files /dev/null and b/assets/textures/PavingStones113/PavingStones113_1K_NormalDX.jpg differ diff --git a/assets/textures/PavingStones113/PavingStones113_1K_Roughness.jpg b/assets/textures/PavingStones113/PavingStones113_1K_Roughness.jpg new file mode 100644 index 00000000..f4ec6d34 Binary files /dev/null and b/assets/textures/PavingStones113/PavingStones113_1K_Roughness.jpg differ diff --git a/assets/textures/PavingStones113/license.md b/assets/textures/PavingStones113/license.md new file mode 100644 index 00000000..e21a2678 --- /dev/null +++ b/assets/textures/PavingStones113/license.md @@ -0,0 +1,27 @@ +# Source + +https://ambientcg.com/view?id=Bricks076C + +# Licensing + +All the PBR materials, brushes, photos and 3D models which are offered for download on ambientCG are provided under the Creative Commons CC0 1.0 Universal License. + +Creative Commons - CC0 1.0 Universal + +# Summary + +The Creative Commons CC0 license gives you the freedom to use these assets… + + … in your own creations of any kind for free. + … even in situations that require them to be redistributed as individual files (for example as part of an open source video game or tool). + … for commercial projects. + +There is no requirement to give credit but it would of course be appreciated. + +# Giving credit + +If you wish to give credit you can do so using this text: + +``` +Contains assets from ambientCG.com, licensed under CC0 1.0 Universal. +``` diff --git a/dotrix_core/src/assets/mesh.rs b/dotrix_core/src/assets/mesh.rs index c12c8d1c..035e7234 100644 --- a/dotrix_core/src/assets/mesh.rs +++ b/dotrix_core/src/assets/mesh.rs @@ -354,7 +354,7 @@ mod tests { [-width, width, width], ]; - width = width / 2.0; + width /= 2.0; let verticies_test_original_2: Vec<[f32; 3]> = vec![ [-width, -width, -width], @@ -367,7 +367,7 @@ mod tests { [-width, width, width], ]; - width = width / 2.0; + width /= 2.0; let verticies_test_original_3: Vec<[u32; 3]> = vec![ [-width as u32, -width as u32, -width as u32], diff --git a/dotrix_core/src/assets/shader.rs b/dotrix_core/src/assets/shader.rs index c4e31f32..bbe61c2d 100644 --- a/dotrix_core/src/assets/shader.rs +++ b/dotrix_core/src/assets/shader.rs @@ -16,6 +16,7 @@ impl Shader { /// Loads the shader to GPU pub fn load(&mut self, renderer: &Renderer) { if !self.module.loaded() { + self.module.label = self.name.clone(); renderer.load_shader(&mut self.module, &self.code); } } diff --git a/dotrix_core/src/assets/texture.rs b/dotrix_core/src/assets/texture.rs index cee7ab0d..3315271e 100644 --- a/dotrix_core/src/assets/texture.rs +++ b/dotrix_core/src/assets/texture.rs @@ -45,4 +45,17 @@ impl Texture { pub fn unload(&mut self) { self.buffer.unload(); } + + /// Fetch data from the gpu + /// + /// This is useful textures that are altered on the gpu + /// + /// This operation is slow and should mostly be + /// used for debugging + pub fn fetch_from_gpu( + &mut self, + renderer: &mut Renderer, + ) -> impl std::future::Future, wgpu::BufferAsyncError>> { + renderer.fetch_texture(&self.buffer, [self.width, self.height, self.depth]) + } } diff --git a/dotrix_core/src/cubemap.rs b/dotrix_core/src/cubemap.rs index 39d0b075..80d0348a 100644 --- a/dotrix_core/src/cubemap.rs +++ b/dotrix_core/src/cubemap.rs @@ -32,7 +32,7 @@ impl Default for CubeMap { bottom: Id::default(), back: Id::default(), front: Id::default(), - buffer: TextureBuffer::new("CubeMap Texture Buffer"), + buffer: TextureBuffer::new_cube("CubeMap Texture Buffer"), } } } diff --git a/dotrix_core/src/renderer.rs b/dotrix_core/src/renderer.rs index 43f20d19..c454dce9 100644 --- a/dotrix_core/src/renderer.rs +++ b/dotrix_core/src/renderer.rs @@ -84,7 +84,7 @@ impl Renderer { buffer.load(self.context(), attributes, indices, count as u32); }*/ - /// Loads the texture buffer to GPU + /// Loads the texture buffer to GPU, this will recreate the texture and will need to be rebound pub fn load_texture<'a>( &self, texture: &mut Texture, @@ -95,11 +95,42 @@ impl Renderer { texture.load(self.context(), width, height, layers); } + /// Load data from cpu to a texture buffer on GPU + /// This will fail if texture has not been loaded with `load_texture` + /// possible unexpected results occur if the dimensions differs from it dimensions at load time + pub fn update_texture<'a>( + &self, + texture: &mut Texture, + width: u32, + height: u32, + layers: &'a [&'a [u8]], + ) { + texture.update(self.context(), width, height, layers); + } + + /// This will `[update_texture]` if texture has been loaded or `[load_texture]` if not + /// the same cavets of `[update_texture]` apply in that care must be taken not to change + /// the dimensions between `load` and `update` + pub fn update_or_load_texture<'a>( + &self, + texture: &mut Texture, + width: u32, + height: u32, + layers: &'a [&'a [u8]], + ) { + texture.update_or_load(self.context(), width, height, layers); + } + /// Loads the buffer to GPU pub fn load_buffer<'a>(&self, buffer: &mut Buffer, data: &'a [u8]) { buffer.load(self.context(), data); } + /// Create a buffer on GPU without data + pub fn create_buffer(&self, buffer: &mut Buffer, size: u32, mapped: bool) { + buffer.create(self.context(), size, mapped); + } + /// Loads the sampler to GPU pub fn load_sampler(&self, sampler: &mut Sampler) { sampler.load(self.context()); @@ -164,6 +195,26 @@ impl Renderer { .run_compute_pipeline(pipeline.shader, &pipeline.bindings, args); } + /// Copy a texture to a buffer + pub fn copy_texture_to_buffer( + &mut self, + texture: &Texture, + buffer: &Buffer, + extent: [u32; 3], + bytes_per_pixel: u32, + ) { + self.context_mut() + .run_copy_texture_to_buffer(texture, buffer, extent, bytes_per_pixel); + } + + /// Fetch texture from GPU + pub fn fetch_texture( + &mut self, + texture: &Texture, + dimensions: [u32; 3], + ) -> impl std::future::Future, wgpu::BufferAsyncError>> { + texture.fetch_from_gpu(dimensions, self.context_mut()) + } /// Returns surface size pub fn surface_size(&self) -> Vec2 { let ctx = self.context(); @@ -263,6 +314,10 @@ pub fn release(mut renderer: Mut) { if renderer.cycle == 0 { renderer.cycle = 1; } + // Check for resource cleanups and mapping callbacks + if let Some(context) = renderer.context.as_ref() { + context.device.poll(wgpu::Maintain::Poll); + } } /// Resize handling system diff --git a/dotrix_core/src/renderer/backend.rs b/dotrix_core/src/renderer/backend.rs new file mode 100644 index 00000000..906973ca --- /dev/null +++ b/dotrix_core/src/renderer/backend.rs @@ -0,0 +1,1056 @@ +/// WGPU backend wrapper module +use std::{borrow::Cow, collections::HashMap}; +use wgpu; +use wgpu::util::DeviceExt; +use winit; + +use crate::{assets::Shader, color::Color, id::Id}; + +use super::{AttributeFormat, BindGroup, Binding, DepthBufferMode, Options, PipelineLayout, Stage}; + +pub(crate) struct Context { + #[allow(dead_code)] + adapter: wgpu::Adapter, + device: wgpu::Device, + queue: wgpu::Queue, + surface: wgpu::Surface, + sur_desc: wgpu::SurfaceConfiguration, + depth_buffer: wgpu::TextureView, + frame: Option, + encoder: Option, + pipelines: HashMap, PipelineBackend>, +} + +impl Context { + pub(crate) fn bind_frame(&mut self, clear_color: &Color) { + let frame = match self.surface.get_current_texture() { + Ok(frame) => frame, + Err(_) => { + self.surface.configure(&self.device, &self.sur_desc); + self.surface + .get_current_texture() + .expect("Failed to acquire next surface texture") + } + }; + + let command_encoder_descriptor = wgpu::CommandEncoderDescriptor { label: None }; + let view = frame + .texture + .create_view(&wgpu::TextureViewDescriptor::default()); + let mut encoder = self + .device + .create_command_encoder(&command_encoder_descriptor); + { + encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: None, + color_attachments: &[wgpu::RenderPassColorAttachment { + view: &view, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Clear(wgpu::Color { + r: clear_color.r as f64, + g: clear_color.g as f64, + b: clear_color.b as f64, + a: clear_color.a as f64, + }), + store: true, + }, + }], + depth_stencil_attachment: Some(wgpu::RenderPassDepthStencilAttachment { + view: &self.depth_buffer, + depth_ops: Some(wgpu::Operations { + load: wgpu::LoadOp::Clear(1.0), + store: true, + }), + stencil_ops: None, + }), + }); + } + self.encoder = Some(encoder); + self.frame = Some(frame); + } + + pub(crate) fn release_frame(&mut self) { + if let Some(encoder) = self.encoder.take() { + self.queue.submit(Some(encoder.finish())); + } + if let Some(frame) = self.frame.take() { + frame.present(); + } + } + + pub(crate) fn resize(&mut self, width: u32, height: u32) { + if width > 0 && height > 0 { + self.sur_desc.width = width; + self.sur_desc.height = height; + + self.surface.configure(&self.device, &self.sur_desc); + self.depth_buffer = create_depth_buffer(&self.device, width, height); + } + } + + pub(crate) fn drop_pipeline(&mut self, shader: Id) { + self.pipelines.remove(&shader); + } + + pub(crate) fn drop_all_pipelines(&mut self) { + self.pipelines.clear(); + } + + pub(crate) fn add_pipeline(&mut self, shader: Id, pipeline_backend: PipelineBackend) { + self.pipelines.insert(shader, pipeline_backend); + } + + pub(crate) fn has_pipeline(&self, shader: Id) -> bool { + self.pipelines.contains_key(&shader) + } + + pub(crate) fn pipeline(&self, shader: Id) -> Option<&PipelineBackend> { + self.pipelines.get(&shader) + } + + pub(crate) fn run_render_pipeline( + &mut self, + shader: Id, + vertex_buffer: &VertexBuffer, + bindings: &Bindings, + options: &Options, + ) { + if let Some(pipeline) = self.pipelines.get(&shader) { + let pipeline_backend = pipeline.instance.render(); + let depth_buffer_mode = pipeline_backend.depth_buffer_mode; + let encoder = self.encoder.as_mut().expect("WGPU encoder must be set"); + + let frame = self.frame.as_ref().expect("WGPU frame must be set"); + let view = frame + .texture + .create_view(&wgpu::TextureViewDescriptor::default()); + let mut rpass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { + label: None, + color_attachments: &[wgpu::RenderPassColorAttachment { + view: &view, + resolve_target: None, + ops: wgpu::Operations { + load: wgpu::LoadOp::Load, + store: true, + }, + }], + depth_stencil_attachment: if depth_buffer_mode != DepthBufferMode::Disabled { + Some(wgpu::RenderPassDepthStencilAttachment { + view: &self.depth_buffer, + depth_ops: Some(wgpu::Operations { + load: wgpu::LoadOp::Load, + store: true, + }), + stencil_ops: None, + }) + } else { + None + }, + }); + + rpass.push_debug_group("Prepare to run pipeline"); + rpass.set_pipeline(&pipeline_backend.wgpu_pipeline); + + if let Some(scissors_rect) = options.scissors_rect.as_ref() { + rpass.set_scissor_rect( + scissors_rect.clip_min_x, + scissors_rect.clip_min_y, + scissors_rect.width, + scissors_rect.height, + ); + } + + for (index, wgpu_bind_group) in bindings.wgpu_bind_groups.iter().enumerate() { + rpass.set_bind_group(index as u32, wgpu_bind_group, &[]); + } + rpass.set_vertex_buffer(0, vertex_buffer.get().slice(..)); + rpass.pop_debug_group(); + + let count = vertex_buffer.count; + + if let Some(indices_buffer) = vertex_buffer.indices().as_ref() { + rpass.insert_debug_marker("Draw indexed"); + rpass.set_index_buffer(indices_buffer.slice(..), wgpu::IndexFormat::Uint32); + rpass.draw_indexed(0..count, 0, options.start_index..options.end_index); + } else { + rpass.insert_debug_marker("Draw"); + rpass.draw(0..count, options.start_index..options.end_index); + } + } + } + + pub(crate) fn run_compute_pipeline( + &mut self, + shader: Id, + bindings: &Bindings, + work_groups: &WorkGroups, + ) { + if let Some(pipeline) = self.pipelines.get(&shader) { + let pipeline_backend = pipeline.instance.compute(); + let encoder = self.encoder.as_mut().expect("WGPU encoder must be set"); + + // compute pass + let mut cpass = + encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { label: None }); + cpass.set_pipeline(&pipeline_backend.wgpu_pipeline); + for (index, wgpu_bind_group) in bindings.wgpu_bind_groups.iter().enumerate() { + cpass.set_bind_group(index as u32, wgpu_bind_group, &[]); + } + cpass.dispatch(work_groups.x, work_groups.y, work_groups.z); + } + } +} + +pub(crate) async fn init(window: &winit::window::Window) -> Context { + let instance = wgpu::Instance::new(wgpu::Backends::PRIMARY); + let surface = unsafe { instance.create_surface(window) }; + let adapter = instance + .request_adapter(&wgpu::RequestAdapterOptions { + power_preference: wgpu::PowerPreference::HighPerformance, + // Request an adapter which can render to our surface + compatible_surface: Some(&surface), + force_fallback_adapter: false, + }) + .await + .expect("Failed to find an appropiate adapter"); + + // Create the logical device and command queue + let (device, queue) = adapter + .request_device( + &wgpu::DeviceDescriptor { + label: None, + features: wgpu::Features::VERTEX_WRITABLE_STORAGE, + limits: wgpu::Limits::default(), + }, + None, // Some(&std::path::Path::new("./wgpu-trace/")), + ) + .await + .expect("Failed to create device"); + + let size = window.inner_size(); + + let sur_desc = wgpu::SurfaceConfiguration { + usage: wgpu::TextureUsages::RENDER_ATTACHMENT, + format: surface.get_preferred_format(&adapter).unwrap(), + width: size.width, + height: size.height, + present_mode: wgpu::PresentMode::Mailbox, + }; + + surface.configure(&device, &sur_desc); + let depth_buffer = create_depth_buffer(&device, size.width, size.height); + + Context { + adapter, + device, + queue, + surface, + sur_desc, + depth_buffer, + frame: None, + encoder: None, + pipelines: std::collections::HashMap::new(), + } +} + +fn create_depth_buffer(device: &wgpu::Device, width: u32, height: u32) -> wgpu::TextureView { + let buffer_extent = wgpu::Extent3d { + width, + height, + depth_or_array_layers: 1, + }; + + let texture = wgpu::TextureDescriptor { + label: Some("Depth Buffer"), + size: buffer_extent, + mip_level_count: 1, + sample_count: 1, + dimension: wgpu::TextureDimension::D2, + format: wgpu::TextureFormat::Depth32Float, + usage: wgpu::TextureUsages::RENDER_ATTACHMENT + | wgpu::TextureUsages::TEXTURE_BINDING + | wgpu::TextureUsages::COPY_DST, + }; + + device + .create_texture(&texture) + .create_view(&wgpu::TextureViewDescriptor::default()) +} + +/// Buffer for vertices attributes +#[derive(Default)] +pub struct VertexBuffer { + /// Packed vertex attributes + attributes: Option, + /// Optional Indices buffer + indices: Option, + count: u32, +} + +impl VertexBuffer { + /// Loads data into the vertex buffer + pub(crate) fn load<'a>( + &mut self, + ctx: &Context, + attributes: &'a [u8], + indices: Option<&'a [u8]>, + count: u32, + ) { + if let Some(buffer) = self.attributes.as_ref() { + ctx.queue.write_buffer(buffer, 0, attributes); + } else { + self.attributes = Some(ctx.device.create_buffer_init( + &wgpu::util::BufferInitDescriptor { + label: Some("VertexBuffer"), + contents: attributes, + usage: wgpu::BufferUsages::VERTEX, + }, + )); + } + + if let Some(buffer) = self.indices.as_ref() { + let indices = indices.expect("Indexed meshed can't be reloaded without indices"); + ctx.queue.write_buffer(buffer, 0, indices); + } else { + self.indices = indices.map(|contents| { + ctx.device + .create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("IndexBuffer"), + contents, + usage: wgpu::BufferUsages::INDEX, + }) + }); + } + + self.count = count; + } + + /// Checks if buffer is empty + pub fn is_empty(&self) -> bool { + self.attributes.is_none() + } + + /// Release all resources used by the buffer + pub fn empty(&mut self) { + self.attributes.take(); + self.indices.take(); + } + + fn get(&self) -> &wgpu::Buffer { + self.attributes + .as_ref() + .expect("Attributes buffer must be loaded") + } + + fn indices(&self) -> Option<&wgpu::Buffer> { + self.indices.as_ref() + } +} + +/// Texture Buffer +pub struct TextureBuffer { + wgpu_texture_view: Option, + mode: super::StorageTextureAccess, + format: super::TextureFormat, +} + +impl Default for TextureBuffer { + fn default() -> Self { + Self { + mode: super::StorageTextureAccess::Read, + format: super::TextureFormat::rgba_u8norm_srgb(), + wgpu_texture_view: None, + } + } +} + +impl TextureBuffer { + /// Create a texture with alternate storage texture access/formats + /// + /// * mode: Storage texture access mode (only used when binding a storage texture) + /// * format: Pixel format + pub fn new(mode: super::StorageTextureAccess, format: super::TextureFormat) -> Self { + Self { + mode, + format, + wgpu_texture_view: Default::default(), + } + } + + /// Loads data into the texture buffer + pub(crate) fn load<'a>( + &mut self, + ctx: &Context, + width: u32, + height: u32, + depth: u32, + layers: &[&'a [u8]], + usage: wgpu::TextureUsages, + ) { + let depth_or_array_layers = layers.len() as u32; + + let size = wgpu::Extent3d { + width, + height, + depth_or_array_layers, + }; + + let layer_size = wgpu::Extent3d { + depth_or_array_layers: 1, + ..size + }; + + let max_mips = 1; //layer_size.max_mips(); + let format: wgpu::TextureFormat = self.format.into(); + + let texture = ctx.device.create_texture(&wgpu::TextureDescriptor { + label: Some("TextureBuffer"), + size, + mip_level_count: max_mips as u32, + sample_count: 1, + dimension: if depth > 1 { + wgpu::TextureDimension::D3 + } else { + wgpu::TextureDimension::D2 + }, + format, + usage, + }); + + self.wgpu_texture_view = Some(texture.create_view(&wgpu::TextureViewDescriptor { + label: None, + format: Some(format), + dimension: Some(if layers.len() == 6 && depth == 1 { + wgpu::TextureViewDimension::Cube + } else if depth > 1 { + wgpu::TextureViewDimension::D3 + } else { + wgpu::TextureViewDimension::D2 + }), + ..wgpu::TextureViewDescriptor::default() + })); + + for (i, data) in layers.iter().enumerate() { + let bytes_per_row = std::num::NonZeroU32::new(data.len() as u32 / height).unwrap(); + + ctx.queue.write_texture( + wgpu::ImageCopyTexture { + texture: &texture, + mip_level: 0, + origin: wgpu::Origin3d { + x: 0, + y: 0, + z: i as u32, + }, + aspect: wgpu::TextureAspect::All, + }, + data, + wgpu::ImageDataLayout { + offset: 0, + bytes_per_row: Some(bytes_per_row), + rows_per_image: Some(std::num::NonZeroU32::new(height).unwrap()), + }, + layer_size, + ); + } + } + + /// Checks if buffer is empty + pub fn loaded(&self) -> bool { + self.wgpu_texture_view.is_some() + } + + /// Release all resources used by the buffer + pub fn unload(&mut self) { + self.wgpu_texture_view.take(); + } + + fn get(&self) -> &wgpu::TextureView { + self.wgpu_texture_view + .as_ref() + .expect("Texture must be loaded") + } +} + +/// Uniform Buffer +#[derive(Default)] +pub struct UniformBuffer { + wgpu_buffer: Option, +} + +impl UniformBuffer { + /// Loads data into the uniform buffer + pub(crate) fn load<'a>(&mut self, ctx: &Context, data: &'a [u8]) { + if let Some(buffer) = self.wgpu_buffer.as_ref() { + ctx.queue.write_buffer(buffer, 0, data); + } else { + self.wgpu_buffer = Some(ctx.device.create_buffer_init( + &wgpu::util::BufferInitDescriptor { + label: Some("UniformBuffer"), + contents: data, + usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, + }, + )); + } + } + + /// Checks if buffer is empty + pub fn is_empty(&self) -> bool { + self.wgpu_buffer.is_none() + } + + /// Release all resources used by the buffer + pub fn empty(&mut self) { + self.wgpu_buffer.take(); + } + + fn get(&self) -> &wgpu::Buffer { + self.wgpu_buffer + .as_ref() + .expect("Uniform buffer must be loaded") + } +} + +/// Texture Sampler +#[derive(Default)] +pub struct Sampler { + wgpu_sampler: Option, +} + +impl Sampler { + /// Loads the Sampler + pub(crate) fn load(&mut self, ctx: &Context) { + if self.wgpu_sampler.is_some() { + return; + } + self.wgpu_sampler = Some(ctx.device.create_sampler(&wgpu::SamplerDescriptor { + address_mode_u: wgpu::AddressMode::Repeat, + address_mode_v: wgpu::AddressMode::Repeat, + address_mode_w: wgpu::AddressMode::Repeat, + mag_filter: wgpu::FilterMode::Nearest, + min_filter: wgpu::FilterMode::Linear, + mipmap_filter: wgpu::FilterMode::Nearest, + ..Default::default() + })); + } + + /// Checks if the Sampler is empty + pub fn is_empty(&self) -> bool { + self.wgpu_sampler.is_none() + } + + /// Release all resources used by the Sampler + pub fn empty(&mut self) { + self.wgpu_sampler.take(); + } + + fn get(&self) -> &wgpu::Sampler { + self.wgpu_sampler.as_ref().expect("Sampler must be loaded") + } +} + +enum StorageBufferMode { + Read, + ReadWrite, +} + +impl Default for StorageBufferMode { + fn default() -> Self { + Self::Read + } +} + +/// Storage Buffer +#[derive(Default)] +pub struct StorageBuffer { + mode: StorageBufferMode, + wgpu_buffer: Option, +} + +impl StorageBuffer { + /// Create a read only storage buffer + pub fn new_readonly() -> Self { + Self { + mode: StorageBufferMode::Read, + wgpu_buffer: Default::default(), + } + } + + /// Create a read-write storage buffer + pub fn new_readwrite() -> Self { + Self { + mode: StorageBufferMode::ReadWrite, + wgpu_buffer: Default::default(), + } + } + + /// Loads data into the storage buffer + pub(crate) fn load<'a>(&mut self, ctx: &Context, data: &'a [u8]) { + if let Some(buffer) = self.wgpu_buffer.as_ref() { + ctx.queue.write_buffer(buffer, 0, data); + } else { + let usage = match self.mode { + StorageBufferMode::Read => { + wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_DST + } + StorageBufferMode::ReadWrite => { + wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC + } + }; + self.wgpu_buffer = Some(ctx.device.create_buffer_init( + &wgpu::util::BufferInitDescriptor { + label: Some("StorageBuffer"), + contents: data, + usage, + }, + )); + } + } + + /// Checks if buffer is empty + pub fn is_empty(&self) -> bool { + self.wgpu_buffer.is_none() + } + + /// Release all resources used by the buffer + pub fn empty(&mut self) { + self.wgpu_buffer.take(); + } + + fn get(&self) -> &wgpu::Buffer { + self.wgpu_buffer + .as_ref() + .expect("Storage buffer must be loaded") + } +} + +/// Render pipeline backend +pub struct RenderPipelineBackend { + /// WGPU pipeline + wgpu_pipeline: wgpu::RenderPipeline, + depth_buffer_mode: DepthBufferMode, +} + +/// Compute pipeline backend +pub struct ComputePipelineBackend { + /// WGPU pipeline + wgpu_pipeline: wgpu::ComputePipeline, +} + +/// Pipeline backend +pub enum PipelineInstance { + Render(RenderPipelineBackend), + Compute(ComputePipelineBackend), +} + +impl PipelineInstance { + fn render(&self) -> &RenderPipelineBackend { + match self { + Self::Render(pipeline) => pipeline, + Self::Compute(_) => panic!("Compute pipeline used for rendering"), + } + } + fn compute(&self) -> &ComputePipelineBackend { + match self { + Self::Compute(pipeline) => pipeline, + Self::Render(_) => panic!("Render pipeline used for rendering"), + } + } +} + +/// Numbers of Work Groups in all directions +pub struct WorkGroups { + /// Number of Work Groups in X direction + pub x: u32, + /// Number of Work Groups in Y direction + pub y: u32, + /// Number of Work Groups in Z direction + pub z: u32, +} + +/// Compute pipeline backend +pub struct PipelineBackend { + /// WGPU bind group layout + wgpu_bind_group_layouts: Vec, + /// WGPU pipeline + instance: PipelineInstance, +} + +#[inline(always)] +fn visibility(stage: &Stage) -> wgpu::ShaderStages { + match stage { + Stage::All => wgpu::ShaderStages::VERTEX | wgpu::ShaderStages::FRAGMENT, + Stage::Vertex => wgpu::ShaderStages::VERTEX, + Stage::Fragment => wgpu::ShaderStages::FRAGMENT, + Stage::Compute => wgpu::ShaderStages::COMPUTE, + } +} + +impl PipelineBackend { + pub(crate) fn new(ctx: &Context, pipeline: &PipelineLayout) -> Self { + let wgpu_shader_module = pipeline.shader.module.get(); + let wgpu_bind_group_layouts = pipeline + .bindings + .iter() + .map(|bind_group_layout| Self::bind_group_layout(&ctx.device, bind_group_layout)) + .collect::>(); + + // create pipeline layout + let pipeline_layout = ctx + .device + .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: None, + bind_group_layouts: wgpu_bind_group_layouts + .iter() + .collect::>() + .as_slice(), + push_constant_ranges: &[], + }); + + let instance = if let Some(mesh) = pipeline.mesh { + let depth_buffer_mode = pipeline.options.depth_buffer_mode; + + // render pipeline: prepare vertex buffers layout + let mut vertex_array_stride = 0; + let vertex_attributes = mesh + .vertex_buffer_layout() + .iter() + .enumerate() + .map(|(index, attr)| { + let offset = vertex_array_stride; + vertex_array_stride += attr.size(); + wgpu::VertexAttribute { + format: match attr { + AttributeFormat::Float32 => wgpu::VertexFormat::Float32, + AttributeFormat::Float32x2 => wgpu::VertexFormat::Float32x2, + AttributeFormat::Float32x3 => wgpu::VertexFormat::Float32x3, + AttributeFormat::Float32x4 => wgpu::VertexFormat::Float32x4, + AttributeFormat::Uint16x2 => wgpu::VertexFormat::Uint16x2, + AttributeFormat::Uint16x4 => wgpu::VertexFormat::Uint16x4, + AttributeFormat::Uint32 => wgpu::VertexFormat::Uint32, + AttributeFormat::Uint32x2 => wgpu::VertexFormat::Uint32x2, + AttributeFormat::Uint32x3 => wgpu::VertexFormat::Uint32x3, + AttributeFormat::Uint32x4 => wgpu::VertexFormat::Uint32x4, + }, + offset: offset as u64, + shader_location: index as u32, + } + }) + .collect::>(); + + let vertex_buffers = [wgpu::VertexBufferLayout { + array_stride: vertex_array_stride as u64, + step_mode: wgpu::VertexStepMode::Vertex, + attributes: vertex_attributes.as_slice(), + }]; + + // create the pipeline + let wgpu_pipeline = + ctx.device + .create_render_pipeline(&wgpu::RenderPipelineDescriptor { + label: Some(&pipeline.label), + layout: Some(&pipeline_layout), + vertex: wgpu::VertexState { + module: wgpu_shader_module, + entry_point: "vs_main", + buffers: &vertex_buffers, + }, + fragment: Some(wgpu::FragmentState { + module: wgpu_shader_module, + entry_point: "fs_main", + targets: &[if depth_buffer_mode == DepthBufferMode::Disabled { + wgpu::ColorTargetState { + format: ctx.sur_desc.format, + blend: Some(wgpu::BlendState { + color: wgpu::BlendComponent { + src_factor: wgpu::BlendFactor::One, + dst_factor: wgpu::BlendFactor::OneMinusSrcAlpha, + operation: wgpu::BlendOperation::Add, + }, + alpha: wgpu::BlendComponent { + src_factor: wgpu::BlendFactor::OneMinusSrcAlpha, + dst_factor: wgpu::BlendFactor::One, + operation: wgpu::BlendOperation::Add, + }, + }), + write_mask: wgpu::ColorWrites::ALL, + } + } else { + wgpu::ColorTargetState { + format: ctx.sur_desc.format, + blend: Some(wgpu::BlendState { + color: wgpu::BlendComponent::REPLACE, + alpha: wgpu::BlendComponent::REPLACE, + }), + write_mask: wgpu::ColorWrites::ALL, + } + }], + }), + primitive: wgpu::PrimitiveState { + front_face: wgpu::FrontFace::Ccw, + cull_mode: if !pipeline.options.disable_cull_mode { + Some(wgpu::Face::Back) + } else { + None + }, + ..Default::default() + }, + depth_stencil: if depth_buffer_mode != DepthBufferMode::Disabled { + Some(wgpu::DepthStencilState { + format: wgpu::TextureFormat::Depth32Float, + depth_write_enabled: depth_buffer_mode == DepthBufferMode::Write, + depth_compare: wgpu::CompareFunction::Less, + stencil: wgpu::StencilState::default(), + bias: wgpu::DepthBiasState { + constant: 2, // corresponds to bilinear filtering + slope_scale: 2.0, + clamp: 0.0, + }, + }) + } else { + None + }, + multisample: wgpu::MultisampleState::default(), + multiview: None, + }); + + PipelineInstance::Render(RenderPipelineBackend { + wgpu_pipeline, + depth_buffer_mode, + }) + } else { + // compute pipeline + let wgpu_pipeline = + ctx.device + .create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some(&pipeline.label), + layout: Some(&pipeline_layout), + module: wgpu_shader_module, + entry_point: "main", + }); + + PipelineInstance::Compute(ComputePipelineBackend { wgpu_pipeline }) + }; + + Self { + wgpu_bind_group_layouts, + instance, + } + } + + fn bind_group_layout(device: &wgpu::Device, bind_group: &BindGroup) -> wgpu::BindGroupLayout { + let entries = bind_group + .bindings + .iter() + .enumerate() + .map(|(index, binding)| match binding { + Binding::Uniform(_, stage, _) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + Binding::Texture(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Texture { + multisampled: false, + sample_type: wgpu::TextureSampleType::Float { + filterable: texture.format.is_filterable(), + }, + view_dimension: wgpu::TextureViewDimension::D2, + }, + count: None, + }, + Binding::TextureCube(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Texture { + multisampled: false, + sample_type: wgpu::TextureSampleType::Float { + filterable: texture.format.is_filterable(), + }, + view_dimension: wgpu::TextureViewDimension::Cube, + }, + count: None, + }, + Binding::Texture3D(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Texture { + multisampled: false, + sample_type: wgpu::TextureSampleType::Float { + filterable: texture.format.is_filterable(), + }, + view_dimension: wgpu::TextureViewDimension::D3, + }, + count: None, + }, + Binding::StorageTexture(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::StorageTexture { + access: texture.mode.into(), + format: texture.format.into(), + view_dimension: wgpu::TextureViewDimension::D2, + }, + count: None, + }, + Binding::StorageTexture3D(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::StorageTexture { + access: texture.mode.into(), + format: texture.format.into(), + view_dimension: wgpu::TextureViewDimension::D3, + }, + count: None, + }, + Binding::Sampler(_, stage, _) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering), + count: None, + }, + Binding::Storage(_, stage, storage) => { + let read_only = matches!(storage.mode, StorageBufferMode::Read); + wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { read_only }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + } + } + Binding::StorageAsUniform(_, stage, _) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: visibility(stage), + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + }) + .collect::>(); + + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some(bind_group.label), + entries: entries.as_slice(), + }) + } +} + +/// Pipeline Bindings +#[derive(Default)] +pub struct Bindings { + wgpu_bind_groups: Vec, +} + +impl Bindings { + pub(crate) fn load( + &mut self, + ctx: &Context, + pipeline: &PipelineBackend, + bind_groups: &[BindGroup], + ) { + self.wgpu_bind_groups = pipeline + .wgpu_bind_group_layouts + .iter() + .enumerate() + .map(|(group, wgpu_bind_group_layout)| { + ctx.device.create_bind_group(&wgpu::BindGroupDescriptor { + layout: wgpu_bind_group_layout, + entries: bind_groups[group] + .bindings + .iter() + .enumerate() + .map(|(binding, entry)| wgpu::BindGroupEntry { + binding: binding as u32, + resource: match entry { + Binding::Uniform(_, _, uniform) => { + uniform.get().as_entire_binding() + } + Binding::Texture(_, _, texture) + | Binding::TextureCube(_, _, texture) + | Binding::Texture3D(_, _, texture) + | Binding::StorageTexture(_, _, texture) + | Binding::StorageTexture3D(_, _, texture) => { + wgpu::BindingResource::TextureView(texture.get()) + } + Binding::Sampler(_, _, sampler) => { + wgpu::BindingResource::Sampler(sampler.get()) + } + Binding::Storage(_, _, storage) => { + storage.get().as_entire_binding() + } + Binding::StorageAsUniform(_, _, storage) => { + storage.get().as_entire_binding() + } + }, + }) + .collect::>() + .as_slice(), + label: None, + }) + }) + .collect::>(); + } + + /// Returns true if bindings was loaded to GPU + pub fn loaded(&self) -> bool { + !self.wgpu_bind_groups.is_empty() + } + + /// Unloads bindings from GPU + pub fn unload(&mut self) { + self.wgpu_bind_groups.clear(); + } +} + +/// Shader Module +#[derive(Default)] +pub struct ShaderModule { + wgpu_shader_model: Option, +} + +impl ShaderModule { + pub(crate) fn load(&mut self, ctx: &Context, name: &str, code: &str) { + self.wgpu_shader_model = Some(ctx.device.create_shader_module( + &wgpu::ShaderModuleDescriptor { + label: Some(name), + source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(code)), + }, + )); + } + + /// Returns true if shader module was loaded to GPU + pub fn loaded(&self) -> bool { + self.wgpu_shader_model.is_some() + } + + /// Unloads the sahder module from GPU + pub fn unload(&mut self) { + self.wgpu_shader_model.take(); + } + + fn get(&self) -> &wgpu::ShaderModule { + self.wgpu_shader_model + .as_ref() + .expect("Shader model must be loaded") + } +} diff --git a/dotrix_core/src/renderer/bindings.rs b/dotrix_core/src/renderer/bindings.rs index 5e554636..3628db18 100644 --- a/dotrix_core/src/renderer/bindings.rs +++ b/dotrix_core/src/renderer/bindings.rs @@ -30,10 +30,16 @@ pub enum Binding<'a> { Uniform(&'a str, Stage, &'a Buffer), /// Texture binding Texture(&'a str, Stage, &'a Texture), + /// 3D Cube Texture binding + TextureCube(&'a str, Stage, &'a Texture), /// 3D Texture binding Texture3D(&'a str, Stage, &'a Texture), + /// 2D Texture Array binding + TextureArray(&'a str, Stage, &'a Texture), /// Storage texture binding StorageTexture(&'a str, Stage, &'a Texture, Access), + /// Storage texture binding 3D + StorageTexture3D(&'a str, Stage, &'a Texture, Access), /// Texture sampler binding Sampler(&'a str, Stage, &'a Sampler), /// Storage binding @@ -76,25 +82,41 @@ impl<'a> BindGroup<'a> { visibility: stage.into(), ty: wgpu::BindingType::Texture { multisampled: false, - sample_type: wgpu::TextureSampleType::Float { - filterable: texture.is_filterable(), - }, + sample_type: texture.sample_type(), view_dimension: wgpu::TextureViewDimension::D2, }, count: None, }, - Binding::Texture3D(_, stage, texture) => wgpu::BindGroupLayoutEntry { + Binding::TextureArray(_, stage, texture) => wgpu::BindGroupLayoutEntry { binding: index as u32, visibility: stage.into(), ty: wgpu::BindingType::Texture { multisampled: false, - sample_type: wgpu::TextureSampleType::Float { - filterable: texture.is_filterable(), - }, + sample_type: texture.sample_type(), + view_dimension: wgpu::TextureViewDimension::D2Array, + }, + count: None, + }, + Binding::TextureCube(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: stage.into(), + ty: wgpu::BindingType::Texture { + multisampled: false, + sample_type: texture.sample_type(), view_dimension: wgpu::TextureViewDimension::Cube, }, count: None, }, + Binding::Texture3D(_, stage, texture) => wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: stage.into(), + ty: wgpu::BindingType::Texture { + multisampled: false, + sample_type: texture.sample_type(), + view_dimension: wgpu::TextureViewDimension::D3, + }, + count: None, + }, Binding::StorageTexture(_, stage, texture, access) => wgpu::BindGroupLayoutEntry { binding: index as u32, visibility: stage.into(), @@ -105,6 +127,18 @@ impl<'a> BindGroup<'a> { }, count: None, }, + Binding::StorageTexture3D(_, stage, texture, access) => { + wgpu::BindGroupLayoutEntry { + binding: index as u32, + visibility: stage.into(), + ty: wgpu::BindingType::StorageTexture { + access: access.into(), + format: texture.format, + view_dimension: wgpu::TextureViewDimension::D3, + }, + count: None, + } + } Binding::Sampler(_, stage, _) => wgpu::BindGroupLayoutEntry { binding: index as u32, visibility: stage.into(), @@ -169,8 +203,11 @@ impl Bindings { uniform.get().as_entire_binding() } Binding::Texture(_, _, texture) + | Binding::TextureArray(_, _, texture) + | Binding::TextureCube(_, _, texture) | Binding::Texture3D(_, _, texture) - | Binding::StorageTexture(_, _, texture, _) => { + | Binding::StorageTexture(_, _, texture, _) + | Binding::StorageTexture3D(_, _, texture, _) => { wgpu::BindingResource::TextureView(texture.get()) } Binding::Sampler(_, _, sampler) => { diff --git a/dotrix_core/src/renderer/buffer.rs b/dotrix_core/src/renderer/buffer.rs index c7ecf578..1325b722 100644 --- a/dotrix_core/src/renderer/buffer.rs +++ b/dotrix_core/src/renderer/buffer.rs @@ -47,6 +47,16 @@ impl Buffer { Self::new(label).use_as_indirect() } + /// Construct new Map Read buffer + pub fn map_read(label: &str) -> Self { + Self::new(label).use_as_map_read() + } + + /// Construct new Map Write buffer + pub fn map_write(label: &str) -> Self { + Self::new(label).use_as_map_write() + } + /// Allow to use as Vertex Buffer #[must_use] pub fn use_as_vertex(mut self) -> Self { @@ -82,6 +92,20 @@ impl Buffer { self } + /// Allow to use as Map Read Buffer + #[must_use] + pub fn use_as_map_read(mut self) -> Self { + self.usage |= wgpu::BufferUsages::MAP_READ; + self + } + + /// Allow to use as Map Write Buffer + #[must_use] + pub fn use_as_map_write(mut self) -> Self { + self.usage |= wgpu::BufferUsages::MAP_WRITE; + self + } + /// Allow reading from buffer #[must_use] pub fn allow_read(mut self) -> Self { @@ -116,6 +140,18 @@ impl Buffer { } } + /// Create buffer of size without data + /// + /// Typically used for staging buffers + pub fn create(&mut self, ctx: &Context, size: u32, mapped: bool) { + self.wgpu_buffer = Some(ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: Some(self.label.as_str()), + size: size as wgpu::BufferAddress, + usage: self.usage, + mapped_at_creation: mapped, + })); + } + /// Check if buffer isloaded pub fn loaded(&self) -> bool { self.wgpu_buffer.is_some() diff --git a/dotrix_core/src/renderer/context.rs b/dotrix_core/src/renderer/context.rs index 2617576b..a23d639b 100644 --- a/dotrix_core/src/renderer/context.rs +++ b/dotrix_core/src/renderer/context.rs @@ -254,6 +254,48 @@ impl Context { cpass.dispatch(args.work_groups.x, args.work_groups.y, args.work_groups.z); } } + + pub(crate) fn run_copy_texture_to_buffer( + &mut self, + texture: &super::Texture, + buffer: &super::Buffer, + extent: [u32; 3], + bytes_per_pixel: u32, + ) { + let encoder = self.encoder.as_mut().expect("WGPU encoder must be set"); + let unpadded_bytes_per_row: u32 = + std::num::NonZeroU32::new(bytes_per_pixel as u32 * extent[0]) + .unwrap() + .into(); + let align = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT as u32; + let padded_bytes_per_row_padding = (align - unpadded_bytes_per_row % align) % align; + let padded_bytes_per_row = unpadded_bytes_per_row + padded_bytes_per_row_padding; + + encoder.copy_texture_to_buffer( + wgpu::ImageCopyTexture { + texture: texture + .wgpu_texture + .as_ref() + .expect("Texture must be loaded"), + mip_level: 0, + origin: wgpu::Origin3d { x: 0, y: 0, z: 0 }, + aspect: wgpu::TextureAspect::All, + }, + wgpu::ImageCopyBuffer { + buffer: buffer.wgpu_buffer.as_ref().expect("Buffer must be ready"), + layout: wgpu::ImageDataLayout { + offset: 0, + bytes_per_row: Some(std::num::NonZeroU32::new(padded_bytes_per_row).unwrap()), + rows_per_image: Some(std::num::NonZeroU32::new(extent[1]).unwrap()), + }, + }, + wgpu::Extent3d { + width: extent[0], + height: extent[1], + depth_or_array_layers: extent[2], + }, + ); + } } pub(crate) async fn init(window: &winit::window::Window, sample_count: u32) -> Context { diff --git a/dotrix_core/src/renderer/texture.rs b/dotrix_core/src/renderer/texture.rs index 003ca48b..190c7cc7 100644 --- a/dotrix_core/src/renderer/texture.rs +++ b/dotrix_core/src/renderer/texture.rs @@ -1,14 +1,25 @@ -use super::Context; +use super::{Buffer, Context}; use wgpu; +pub enum TextureKind { + D2, + D2Array, + Cube, + D3, +} + /// GPU Texture Implementation pub struct Texture { /// Texture label pub label: String, /// WGPU Texture view pub wgpu_texture_view: Option, + /// WGPU Texture + pub wgpu_texture: Option, /// Texture usage pub usage: wgpu::TextureUsages, + /// Texture kind + pub kind: TextureKind, /// Texture format pub format: wgpu::TextureFormat, } @@ -18,8 +29,10 @@ impl Default for Texture { Self { label: String::from("Noname Texture"), wgpu_texture_view: None, + wgpu_texture: None, usage: wgpu::TextureUsages::empty(), format: wgpu::TextureFormat::Rgba8UnormSrgb, + kind: TextureKind::D2, } } } @@ -34,6 +47,36 @@ impl Texture { } } + /// Constructs a CubeMap GPU Texture + pub fn new_cube(label: &str) -> Self { + Self { + label: String::from(label), + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST, + kind: TextureKind::Cube, + ..Default::default() + } + } + + /// Constructs a 2D Array GPU Texture + pub fn new_array(label: &str) -> Self { + Self { + label: String::from(label), + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST, + kind: TextureKind::D2Array, + ..Default::default() + } + } + + /// Constructs a 3D GPU Texture + pub fn new_3d(label: &str) -> Self { + Self { + label: String::from(label), + usage: wgpu::TextureUsages::TEXTURE_BINDING | wgpu::TextureUsages::COPY_DST, + kind: TextureKind::D3, + ..Default::default() + } + } + /// Constructs GPU Storage Texture pub fn storage(label: &str) -> Self { Self { @@ -101,8 +144,19 @@ impl Texture { self } - /// Loads data into the texture buffer + /// Loads data into the 2d texture buffer + /// This will recreate the backend texture if called twice, as a result + /// any texture will need to be rebound before it updates pub(crate) fn load<'a>(&mut self, ctx: &Context, width: u32, height: u32, layers: &[&'a [u8]]) { + let dimension = match self.kind { + TextureKind::D2 => wgpu::TextureViewDimension::D2, + TextureKind::D2Array => wgpu::TextureViewDimension::D2Array, + TextureKind::Cube => { + assert!(layers.len() == 6); + wgpu::TextureViewDimension::Cube + } + TextureKind::D3 => wgpu::TextureViewDimension::D3, + }; let format = self.format; let usage = self.usage; let depth_or_array_layers = layers.len() as u32; @@ -111,18 +165,21 @@ impl Texture { height, depth_or_array_layers, }; - let layer_size = wgpu::Extent3d { - depth_or_array_layers: 1, - ..size - }; let max_mips = 1; + let tex_dimension: wgpu::TextureDimension = match self.kind { + TextureKind::D2 => wgpu::TextureDimension::D2, + TextureKind::D2Array => wgpu::TextureDimension::D2, + TextureKind::Cube => wgpu::TextureDimension::D2, + TextureKind::D3 => wgpu::TextureDimension::D3, + }; + let texture = ctx.device.create_texture(&wgpu::TextureDescriptor { label: Some(&self.label), size, mip_level_count: max_mips as u32, sample_count: 1, - dimension: wgpu::TextureDimension::D2, + dimension: tex_dimension, format, usage, }); @@ -130,35 +187,78 @@ impl Texture { self.wgpu_texture_view = Some(texture.create_view(&wgpu::TextureViewDescriptor { label: None, format: Some(format), - dimension: Some(if depth_or_array_layers == 6 { - wgpu::TextureViewDimension::Cube - } else { - wgpu::TextureViewDimension::D2 - }), + dimension: Some(dimension), ..wgpu::TextureViewDescriptor::default() })); - for (i, data) in layers.iter().enumerate() { - let bytes_per_row = std::num::NonZeroU32::new(data.len() as u32 / height).unwrap(); - ctx.queue.write_texture( - wgpu::ImageCopyTexture { - texture: &texture, - mip_level: 0, - origin: wgpu::Origin3d { - x: 0, - y: 0, - z: i as u32, + self.wgpu_texture = Some(texture); + + self.update(ctx, width, height, layers) + } + + /// This will write to a texture but not create it + /// This can be used to update a texture's value with out recreating/rebinding it + /// however if the size of the texture is changed it will behave oddly or even panic + pub(crate) fn update<'a>( + &mut self, + ctx: &Context, + width: u32, + height: u32, + layers: &[&'a [u8]], + ) { + let depth_or_array_layers = layers.len() as u32; + let size = wgpu::Extent3d { + width, + height, + depth_or_array_layers, + }; + let layer_size = wgpu::Extent3d { + depth_or_array_layers: 1, + ..size + }; + + if let Some(texture) = self.wgpu_texture.as_ref() { + for (i, data) in layers.iter().enumerate() { + let bytes_per_row = std::num::NonZeroU32::new(data.len() as u32 / height).unwrap(); + ctx.queue.write_texture( + wgpu::ImageCopyTexture { + texture, + mip_level: 0, + origin: wgpu::Origin3d { + x: 0, + y: 0, + z: i as u32, + }, + aspect: wgpu::TextureAspect::All, + }, + data, + wgpu::ImageDataLayout { + offset: 0, + bytes_per_row: Some(bytes_per_row), + rows_per_image: Some(std::num::NonZeroU32::new(height).unwrap()), }, - aspect: wgpu::TextureAspect::All, - }, - data, - wgpu::ImageDataLayout { - offset: 0, - bytes_per_row: Some(bytes_per_row), - rows_per_image: Some(std::num::NonZeroU32::new(height).unwrap()), - }, - layer_size, - ); + layer_size, + ); + } + } + } + + /// This method will update a gpu texture if it exists with new data or + /// load a new texture onto the gpu if it does not. + /// + /// The same cavets of [`update`] apply in that care must be taken to not + /// change the size of the texture between [`load`] and [`update`] + pub(crate) fn update_or_load<'a>( + &mut self, + ctx: &Context, + width: u32, + height: u32, + layers: &[&'a [u8]], + ) { + if self.wgpu_texture.is_none() { + self.load(ctx, width, height, layers); + } else { + self.update(ctx, width, height, layers); } } @@ -169,6 +269,7 @@ impl Texture { /// Release all resources used by the texture pub fn unload(&mut self) { + self.wgpu_texture.take(); self.wgpu_texture_view.take(); } @@ -183,4 +284,90 @@ impl Texture { pub fn is_filterable(&self) -> bool { self.format.describe().guaranteed_format_features.filterable } + + /// Get the texture bytes per pixels + pub fn pixel_bytes(&self) -> u8 { + self.format.describe().block_size + } + + /// Get the number of channels + pub fn num_channels(&self) -> u8 { + self.format.describe().components + } + + /// Get the texture sample type (float/uint etc) + pub fn sample_type(&self) -> wgpu::TextureSampleType { + self.format.describe().sample_type + } + + /// Fetch data from the gpu + /// + /// This is useful textures that are altered on the gpu + /// + /// This operation is slow and should mostly be + /// used for debugging + pub fn fetch_from_gpu( + &self, + dimensions: [u32; 3], + ctx: &mut Context, + ) -> impl std::future::Future, wgpu::BufferAsyncError>> { + let bytes_per_pixel: u32 = self.pixel_bytes() as u32; + let mut staging_buffer = Buffer::map_read("Texture Fetch Staging buffer"); + let unpadded_bytes_per_row: u32 = + std::num::NonZeroU32::new(bytes_per_pixel as u32 * dimensions[0]) + .unwrap() + .into(); + let align = wgpu::COPY_BYTES_PER_ROW_ALIGNMENT as u32; + let padded_bytes_per_row_padding = (align - unpadded_bytes_per_row % align) % align; + let padded_bytes_per_row = unpadded_bytes_per_row + padded_bytes_per_row_padding; + + staging_buffer.create( + ctx, + padded_bytes_per_row * dimensions[0] * dimensions[1], + false, + ); + ctx.run_copy_texture_to_buffer(self, &staging_buffer, dimensions, bytes_per_pixel); + + async move { + // TODO: Urgently work out a better way to await the next frame. + std::thread::sleep(std::time::Duration::from_secs(1)); + + let wgpu_buffer = staging_buffer.wgpu_buffer.expect("Buffer must be loaded"); + let buffer_slice = wgpu_buffer.slice(..); + // Gets the future representing when `staging_buffer` can be read from + let buffer_future = buffer_slice.map_async(wgpu::MapMode::Read); + + match buffer_future.await { + Ok(()) => { + // Gets contents of buffer + let data = buffer_slice.get_mapped_range(); + // This strips the padding on each row + let result: Vec = data + .chunks_exact((padded_bytes_per_row * dimensions[1]) as usize) + .flat_map(|img| { + let rows: Vec> = img + .chunks_exact(padded_bytes_per_row as usize) + .map(|row| row[0..(unpadded_bytes_per_row as usize)].to_vec()) + .collect(); + rows + }) + .flatten() + .collect(); + + // With the current interface, we have to make sure all mapped views are + // dropped before we unmap the buffer. + drop(data); + wgpu_buffer.unmap(); // Unmaps buffer from memory + // If you are familiar with C++ these 2 lines can be thought of similarly to: + // delete myPointer; + // myPointer = NULL; + // It effectively frees the memory + // + Ok(result) + } + + Err(e) => Err(e), + } + } + } } diff --git a/dotrix_math/src/lib.rs b/dotrix_math/src/lib.rs index 0e06311d..522a8702 100644 --- a/dotrix_math/src/lib.rs +++ b/dotrix_math/src/lib.rs @@ -11,6 +11,7 @@ pub use cgmath::num_traits::clamp_min; pub use cgmath::perspective; pub use cgmath::Deg; pub use cgmath::InnerSpace; +pub use cgmath::Matrix; pub use cgmath::MetricSpace; pub use cgmath::Rad; pub use cgmath::Rotation3; diff --git a/dotrix_pbr/src/shaders/light.inc.wgsl b/dotrix_pbr/src/shaders/light.inc.wgsl index c270bc26..ef6edaf0 100644 --- a/dotrix_pbr/src/shaders/light.inc.wgsl +++ b/dotrix_pbr/src/shaders/light.inc.wgsl @@ -51,7 +51,6 @@ var u_light: Light; fn calculate_directional( light: DirectionalLight, - normal: vec3, ) -> LightCalcOutput { let light_direction: vec3 = normalize(-light.direction.xyz); @@ -65,7 +64,6 @@ fn calculate_directional( fn calculate_point( light: PointLight, position: vec3, - normal: vec3, ) -> LightCalcOutput { let light_direction: vec3 = normalize(light.position.xyz - position); @@ -85,7 +83,6 @@ fn calculate_point( fn calculate_simple( light: SimpleLight, position: vec3, - normal: vec3, ) -> LightCalcOutput { let light_direction: vec3 = normalize(light.position.xyz - position.xyz); @@ -99,7 +96,6 @@ fn calculate_simple( fn calculate_spot( light: SpotLight, position: vec3, - normal: vec3, ) -> LightCalcOutput { let light_direction: vec3 = normalize(light.position.xyz - position.xyz); let theta: f32 = dot(light_direction, normalize(-light.direction.xyz)); @@ -113,6 +109,59 @@ fn calculate_spot( return out; } +// This will get the direction direction and intensity of +// the nth light towards a position +// If used in conjectuion with `get_light_count` +// It allows for more consistent iter code by providing +// A standard single data `LightCalcOutput` for any light +// regardless of type +fn calculate_nth_light_ray( + in_camera_index: u32, + position: vec3, +) -> LightCalcOutput { + var camera_index: u32 = in_camera_index; + // directional + let dir_count = min(u32(u_light.count.x), MAX_LIGHTS_COUNT); + if (camera_index < dir_count) { + var light: DirectionalLight = u_light.directional[camera_index]; + return calculate_directional(light); + } + camera_index = camera_index - dir_count; + // point + let point_count = min(u32(u_light.count.y), MAX_LIGHTS_COUNT); + if (camera_index < point_count) { + var light: PointLight = u_light.point[camera_index]; + return calculate_point(light, position); + } + camera_index = camera_index - point_count; + // simple + let simple_count = min(u32(u_light.count.z), MAX_LIGHTS_COUNT); + if (camera_index < simple_count) { + var light: SimpleLight = u_light.simple[camera_index]; + return calculate_simple(light, position); + } + camera_index = camera_index - simple_count; + // spot + let spot_count = min(u32(u_light.count.w), MAX_LIGHTS_COUNT); + if (camera_index < spot_count) { + var light: SpotLight = u_light.spot[camera_index]; + return calculate_spot(light, position); + } + // Trying to access a non existant light + var oob: LightCalcOutput; + oob.light_direction = vec3(0.); + oob.radiance = vec3(0.); + return oob; +} + +fn get_light_count() -> u32 { + return u_light.count.x + u_light.count.y + u_light.count.z + u_light.count.w; +} + +fn get_ambient() -> vec3 { + return u_light.ambient.xyz; +} + fn distribution_ggx(normal: vec3, halfway: vec3, roughness: f32) -> f32 { let a: f32 = roughness*roughness; @@ -202,66 +251,9 @@ fn calculate_lighting( // Directions light var i: u32 = 0u; - var count: u32 = min(u32(u_light.count.x), MAX_LIGHTS_COUNT); + var count: u32 = get_light_count(); for (i = 0u; i< count; i = i + 1u) { - let light_result = calculate_directional( - u_light.directional[i], - normal - ); - light_color = light_color + pbr( - light_result, - camera_direction, - normal, - fresnel_schlick_0, - albedo, - metallic, - roughness - ); - } - // Point light - count = min(u32(u_light.count.y), MAX_LIGHTS_COUNT); - for (i = 0u; i< count; i = i + 1u) { - let light_result = calculate_point( - u_light.point[i], - position, - normal - ); - light_color = light_color + pbr( - light_result, - camera_direction, - normal, - fresnel_schlick_0, - albedo, - metallic, - roughness - ); - } - // Simple light - count = min(u32(u_light.count.z), MAX_LIGHTS_COUNT); - for (i = 0u; i< count; i = i + 1u) { - let light_result = calculate_simple( - u_light.simple[i], - position, - normal - ); - light_color = light_color + pbr( - light_result, - camera_direction, - normal, - fresnel_schlick_0, - albedo, - metallic, - roughness - ); - } - // Spot light - count = min(u32(u_light.count.w), MAX_LIGHTS_COUNT); - for (i = 0u; i< count; i = i + 1u) { - let light_result = calculate_spot( - u_light.spot[i], - position, - normal - ); + let light_result = calculate_nth_light_ray(i, position); light_color = light_color + pbr( light_result, camera_direction, @@ -274,7 +266,7 @@ fn calculate_lighting( } // Ambient - let ambient = u_light.ambient.xyz * albedo * ao; + let ambient = get_ambient() * albedo * ao; light_color = light_color + ambient; // Gamma correct diff --git a/dotrix_sky/src/skybox.rs b/dotrix_sky/src/skybox.rs index 41c10013..02d99591 100644 --- a/dotrix_sky/src/skybox.rs +++ b/dotrix_sky/src/skybox.rs @@ -160,7 +160,7 @@ pub fn render( ), BindGroup::new( "Locals", - vec![Binding::Texture3D( + vec![Binding::TextureCube( "CubeMap", Stage::Fragment, &cubemap.buffer, diff --git a/dotrix_voxel/Cargo.toml b/dotrix_voxel/Cargo.toml new file mode 100644 index 00000000..33fff8e3 --- /dev/null +++ b/dotrix_voxel/Cargo.toml @@ -0,0 +1,39 @@ +[package] +name = "dotrix_voxel" +version = "0.5.0" +authors = ["Elias Kartashov ", "Andrew King "] +edition = "2018" + +description = "Voxels for Dotrix 3D Game Engine (under construction)" +license = "MIT" +resolver = "2" + +[dependencies] +dotrix_core = { version = "0.5", path = "../dotrix_core" } +dotrix_math = { version = "0.4", path = "../dotrix_math" } +dotrix_pbr = { version = "0.2", path = "../dotrix_pbr" } +dotrix_primitives = { version = "0.1", path = "../dotrix_primitives" } +rayon = "1.5" + +[dependencies.num] +version = "^0.3" + +[dependencies.rand] +version = "0.8.4" +features = ["small_rng"] + +[dependencies.wgpu] +version = "0.12" +features = ["trace"] + +[dependencies.bytemuck] +version = "1.4" +features = ["derive"] + +[dependencies.futures] +version = "0.3" +default-features = false +features = ["std", "executor"] + +[dependencies.tera] +version = "1.15.0" diff --git a/dotrix_voxel/src/grid.rs b/dotrix_voxel/src/grid.rs new file mode 100644 index 00000000..49159971 --- /dev/null +++ b/dotrix_voxel/src/grid.rs @@ -0,0 +1,182 @@ +use super::Voxel; +use dotrix_core::{ + renderer::{wgpu, Texture as TextureBuffer}, + Assets, Renderer, +}; + +const DEFAULT_DIM: usize = 3; + +/// A grid of voxels +pub struct Grid { + /// The number of voxels per dimension + /// This is not public to restric resizing + /// which needs a full recreate and rebind + /// of all of it's 3d textures if changed + /// Instead it can only be resized with + /// `[with_dimensions]` which takes self + dimensions: [u32; 3], + /// The voxels + voxels: Vec, + /// 3D Texture buffer + buffer: TextureBuffer, + /// Tracks if changed by incremented a revision number on each change + revision: u32, + /// Revision number of last load + /// Setting this to `None` will force it to reload + last_update: Option, +} + +impl Default for Grid { + fn default() -> Self { + Grid { + dimensions: [DEFAULT_DIM as u32, DEFAULT_DIM as u32, DEFAULT_DIM as u32], + voxels: vec![Default::default(); DEFAULT_DIM * DEFAULT_DIM * DEFAULT_DIM], + buffer: { + let mut buffer = TextureBuffer::new_3d("VoxelGrid"); + buffer.format = wgpu::TextureFormat::Rg8Uint; + buffer + }, + revision: 0, + last_update: None, + } + } +} + +impl Grid { + /// Start building a grid + pub fn build() -> Self { + Default::default() + } + #[must_use] + /// Build the grid with new dimensions, default values will fill the new voxels + pub fn with_dimensions>(mut self, dimensions: T) -> Self { + let dimensions: [u32; 3] = dimensions.into(); + self.dimensions = dimensions; + let count: usize = (dimensions[0] * dimensions[1] * dimensions[2]) as usize; + // Resize number of voxels to match + self.voxels.resize(count, Default::default()); + // Must recreate the binding texture + // Shader must rebind to see the update + self.buffer = { + let mut buffer = TextureBuffer::new_3d("VoxelGrid"); + buffer.format = wgpu::TextureFormat::Rg8Uint; + buffer + }; + Self::flag_changed(self) + } + #[must_use] + /// Build the grid with these values for the voxel + pub fn with_values>(mut self, values: T) -> Self { + self.set_values(values); + self + } + /// Set the values of the voxel, Extra values are ignored + pub fn set_values>(&mut self, values: T) { + let input: &[u8] = values.as_ref(); + let count: usize = (self.dimensions[0] * self.dimensions[1] * self.dimensions[2]) as usize; + + let slice_len = std::cmp::min(input.len(), count); + + // If they provide too much or too little values it is silently ignored + // TODO: Should we panic? + self.voxels[0..slice_len] + .iter_mut() + .enumerate() + .for_each(|(i, v)| (v).value = input[i]); + self.set_changed(); + } + #[must_use] + /// Build the grid with these material values + pub fn with_materials>(mut self, values: T) -> Self { + self.set_materials(values); + self + } + /// Set the material values + pub fn set_materials>(&mut self, values: T) { + let input: &[u8] = values.as_ref(); + let count: usize = (self.dimensions[0] * self.dimensions[1] * self.dimensions[2]) as usize; + + let slice_len = std::cmp::min(input.len(), count); + + // If they provide too much or too little values it is silently ignored + // TODO: Should we panic? + self.voxels[0..slice_len] + .iter_mut() + .enumerate() + .for_each(|(i, v)| (v).material = input[i]); + self.set_changed(); + } + + #[must_use] + /// Build the grid with it flagged as changed + pub fn flag_changed(mut self) -> Self { + self.set_changed(); + self + } + /// Set the grid as changed + pub fn set_changed(&mut self) { + self.revision += 1; + } + + /// Get the current revision + /// This is incremented on each change of the grid + /// that requires a reload + pub fn get_revision(&self) -> u32 { + self.revision + } + + /// The 3DTexture buffer, must first be loaded with [`load`] + pub fn get_buffer(&self) -> &TextureBuffer { + &self.buffer + } + + /// The number of voxels in each dimension + pub fn get_dimensions(&self) -> &[u32; 3] { + &self.dimensions + } + + /// Same as `[get_dimensions]` but as f32 (convenince method) + pub fn get_size(&self) -> [f32; 3] { + [ + self.dimensions[0] as f32, + self.dimensions[1] as f32, + self.dimensions[2] as f32, + ] + } + + /// Load the grid into a 3DTexture + pub fn load(&mut self, renderer: &Renderer, _assets: &Assets) { + if let Some(last_update) = self.last_update { + if last_update == self.revision && self.buffer.loaded() { + return; + } + } + + let data: Vec> = self + .voxels + .chunks(self.dimensions[0] as usize * self.dimensions[1] as usize) + .map(|chunk| { + chunk + .iter() + .flat_map(|voxel| [voxel.value, voxel.material]) + .collect() + }) + .collect(); + + let slices: Vec<&[u8]> = data.iter().map(|chunk| chunk.as_slice()).collect(); + + renderer.update_or_load_texture( + &mut self.buffer, + self.dimensions[0], + self.dimensions[1], + slices.as_slice(), + ); + + self.last_update = Some(self.revision); + } + + /// Unloads the [`Grid`] data from the GPU + pub fn unload(&mut self) { + self.buffer.unload(); + } +} diff --git a/dotrix_voxel/src/lib.rs b/dotrix_voxel/src/lib.rs new file mode 100644 index 00000000..4a4f7dc4 --- /dev/null +++ b/dotrix_voxel/src/lib.rs @@ -0,0 +1,22 @@ +//! Voxel Module +//! +//! Handles general voxel related content, such as conversion to an explicit +//! mesh using marching cubes or direct rendering. +//! + +use dotrix_core::Application; + +mod grid; +mod material_set; +mod sdf; +mod voxel; + +pub use grid::Grid; +pub use material_set::*; +pub use sdf::*; +pub use voxel::Voxel; + +/// Enables Voxel Dotrix Extension +pub fn extension(app: &mut Application) { + sdf::extension(app); +} diff --git a/dotrix_voxel/src/material_set.rs b/dotrix_voxel/src/material_set.rs new file mode 100644 index 00000000..8fdc028c --- /dev/null +++ b/dotrix_voxel/src/material_set.rs @@ -0,0 +1,203 @@ +//! Set of materials used for voxels +//! +//! Each voxel has a material ID, which corresponds to a +//! material in this set +//! +//! All material textures must have the same size +//! + +use dotrix_core::{ + assets::Texture, + renderer::{Buffer, Texture as TextureBuffer}, + Assets, Id, Renderer, +}; +use dotrix_pbr::Material; + +use std::collections::HashMap; + +pub struct MaterialSet { + materials: HashMap, + last_num_textures: u32, + texture_buffer: TextureBuffer, + material_buffer: Buffer, +} + +#[repr(C)] +#[derive(Copy, Clone, Debug)] +struct MaterialData { + // Ids of texture in texture array -1 means not precent + albedo_id: i32, + roughness_id: i32, + metallic_id: i32, + ao_id: i32, + normal_id: i32, + padding_a: [i32; 3], + // These are uses if the texture IDs above are -1 + albedo: [f32; 4], + roughness: f32, + metallic: f32, + ao: f32, + padding_b: [f32; 1], +} +impl Default for MaterialData { + fn default() -> Self { + Self { + albedo_id: -1, + roughness_id: -1, + metallic_id: -1, + ao_id: -1, + normal_id: -1, + albedo: [1., 0.5, 1., 1.], + roughness: 0., + metallic: 0., + ao: 0., + padding_a: Default::default(), + padding_b: Default::default(), + } + } +} +unsafe impl bytemuck::Zeroable for MaterialData {} +unsafe impl bytemuck::Pod for MaterialData {} + +impl Default for MaterialSet { + fn default() -> Self { + Self { + materials: Default::default(), + last_num_textures: 0, + + texture_buffer: TextureBuffer::new_array("MaterialSet") + .use_as_texture() + .rgba_u8norm_srgb(), + material_buffer: Buffer::uniform("MaterialSetIndicies"), + } + } +} + +impl MaterialSet { + /// Set the material for a material ID + pub fn set_material(&mut self, material_id: u8, material: Material) { + self.materials.insert(material_id, material); + } + + /// Clear the material for a material ID + pub fn clear_material(&mut self, material_id: u8) { + self.materials.remove(&material_id); + } + + /// The material buffer. Contains the materials data for the gpu + /// Must be loaed before use with `[load]` + pub fn get_material_buffer(&self) -> &Buffer { + &self.material_buffer + } + + /// The texture array buffer. Contains the textures as referenced by the material buffer for the gpu + /// Must be loaed before use with `[load]` + pub fn get_texture_buffer(&self) -> &TextureBuffer { + &self.texture_buffer + } + + /// Returns true if a full rebind is required + /// returns false if rebind is not required, but may still update textures + /// by replacing current textures. + pub fn load(&mut self, renderer: &Renderer, assets: &Assets) -> bool { + let mut result = false; + + let number_of_materials = 256; + let number_of_textures_per_material = 5; + let mut material_data: Vec = vec![Default::default(); number_of_materials]; + + let mut num_texs_found = 0; + let mut textures: Vec> = vec![]; + let mut texture_data_size = None; + let mut texture_id_idx_map: HashMap, i32> = Default::default(); + for (&material_id, material) in self.materials.iter() { + let i = material_id as usize; + let mut tex_ids = vec![-1; number_of_textures_per_material]; + for (j, tex_id) in [ + material.texture, + material.roughness_texture, + material.metallic_texture, + material.ao_texture, + material.normal_texture, + ] + .iter() + .enumerate() + { + if !tex_id.is_null() { + if let Some(texture) = assets.get(*tex_id) { + let tex_idx = texture_id_idx_map.entry(*tex_id).or_insert_with(|| { + let data = &texture.data; + if let Some(texture_data_size) = texture_data_size { + // TODO: Should we silently ignore/Print a warning/resize/clip? + let (width, height, depth, data_len) = texture_data_size; + assert_eq!(width, texture.width); + assert_eq!(height, texture.height); + assert_eq!(depth, texture.depth); + assert_eq!(data_len, data.len()); + } else { + texture_data_size = + Some((texture.width, texture.height, texture.depth, data.len())) + } + num_texs_found += 1; + let idx = textures.len(); + textures.push(data.clone()); + idx as i32 + }); + tex_ids[j] = *tex_idx; + } + } + material_data[i].albedo_id = tex_ids[0]; + material_data[i].roughness_id = tex_ids[1]; + material_data[i].metallic_id = tex_ids[2]; + material_data[i].ao_id = tex_ids[3]; + material_data[i].normal_id = tex_ids[4]; + material_data[i].albedo = material.albedo.into(); + material_data[i].roughness = material.roughness; + material_data[i].metallic = material.metallic; + material_data[i].ao = material.ao; + } + + if num_texs_found != self.last_num_textures { + // Full reload and bind required + // because number of textures was changed + result = true; + } + + if result { + self.texture_buffer.unload(); + } + + if num_texs_found == 0 { + // Set as dummy texture + renderer.update_or_load_texture( + &mut self.texture_buffer, + 1, + 1, + &[&[0u8, 0u8, 0u8, 0u8]], + ); + } else if let Some(texture_data_size) = texture_data_size { + let (width, height, _, _) = texture_data_size; + + let slices: Vec<&[u8]> = textures.iter().map(|tex| tex.as_slice()).collect(); + + renderer.update_or_load_texture( + &mut self.texture_buffer, + width, + height, + slices.as_slice(), + ); + } else { + unreachable!(); + } + + self.last_num_textures = num_texs_found; + } + + renderer.load_buffer( + &mut self.material_buffer, + bytemuck::cast_slice(material_data.as_slice()), + ); + + result + } +} diff --git a/dotrix_voxel/src/sdf/circle_trace.rs b/dotrix_voxel/src/sdf/circle_trace.rs new file mode 100644 index 00000000..e9c547eb --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace.rs @@ -0,0 +1,232 @@ +use crate::Grid; +use crate::MaterialSet; +use crate::TexSdf; +use dotrix_core::{ + assets::{Mesh, Shader}, + ecs::{Const, Mut, System}, + renderer::{BindGroup, Binding, PipelineLayout, RenderOptions, Sampler, Stage}, + Application, Assets, Globals, Renderer, Transform, World, +}; +use dotrix_math::*; +use dotrix_primitives::Cube; +use tera::{Context, Tera}; + +mod camera; +mod lights; + +use camera::CameraBuffer; +pub use lights::Light; +use lights::LightStorageBuffer; + +const PIPELINE_LABEL: &str = "dotrix_voxel::sdf::circle_trace"; + +#[repr(C)] +#[derive(Default, Copy, Clone)] +struct SdfBufferData { + // This transform scales the 1x1x1 cube so that it totally encloses the + // voxels + pub cube_transform: [[f32; 4]; 4], + // Inverse fo cube_transform + pub inv_cube_transform: [[f32; 4]; 4], + // World transform of the voxel grid + pub world_transform: [[f32; 4]; 4], + // Inverse of world_transform + pub inv_world_transform: [[f32; 4]; 4], + // Converts normals from object space to world space + pub normal_transform: [[f32; 4]; 4], + // Converts normals from world space to object space + pub inv_normal_transform: [[f32; 4]; 4], + // Dimensions of the voxel grid + pub grid_dimensions: [f32; 4], + // World space scale + pub world_scale: [f32; 4], +} + +unsafe impl bytemuck::Zeroable for SdfBufferData {} +unsafe impl bytemuck::Pod for SdfBufferData {} + +pub fn startup(renderer: Const, mut assets: Mut) { + let mut templates = Tera::default(); + templates + .add_raw_templates(vec![ + ( + "circle_trace/render.wgsl", + include_str!("./circle_trace/render.wgsl"), + ), + ( + "circle_trace/map.inc.wgsl", + include_str!("./circle_trace/map.inc.wgsl"), + ), + ( + "circle_trace/accelerated_raytrace.inc.wgsl", + include_str!("./circle_trace/accelerated_raytrace.inc.wgsl"), + ), + ( + "circle_trace/hemisphere_ambient_occulsion.inc.wgsl", + include_str!("./circle_trace/hemisphere_ambient_occulsion.inc.wgsl"), + ), + ( + "circle_trace/lighting.inc.wgsl", + include_str!("./circle_trace/lighting.inc.wgsl"), + ), + ( + "circle_trace/pbr.inc.wgsl", + include_str!("./circle_trace/pbr.inc.wgsl"), + ), + ( + "circle_trace/soft_shadows_closet_approach.inc.wgsl", + include_str!("./circle_trace/soft_shadows_closet_approach.inc.wgsl"), + ), + ( + "circle_trace/triplanar_surface.inc.wgsl", + include_str!("./circle_trace/triplanar_surface.inc.wgsl"), + ), + ]) + .unwrap(); + + let context = Context::new(); + let mut shader = Shader { + name: String::from(PIPELINE_LABEL), + code: templates + .render("circle_trace/render.wgsl", &context) + .unwrap(), + ..Default::default() + }; + shader.load(&renderer); + + assets.store_as(shader, PIPELINE_LABEL); + + let mut mesh = Cube::builder(1.0).with_positions().mesh(); + mesh.load(&renderer); + assets.store_as(mesh, PIPELINE_LABEL); +} + +pub fn render( + mut renderer: Mut, + world: Const, + assets: Const, + globals: Const, +) { + let camera_buffer = globals + .get::() + .expect("ProjView buffer must be loaded"); + + for (grid, sdf, world_transform, material_set) in + world.query::<(&Grid, &mut TexSdf, &Transform, &mut MaterialSet)>() + { + if sdf.pipeline.shader.is_null() { + sdf.pipeline.shader = assets.find::(PIPELINE_LABEL).unwrap_or_default(); + } + if !sdf.pipeline.cycle(&renderer) { + return; + } + let mesh = assets + .get( + assets + .find::(PIPELINE_LABEL) + .expect("Sdf mesh must be initialized with the dotrix_voxel startup system"), + ) + .unwrap(); + + let grid_size = grid.get_size(); + let scale = Mat4::from_nonuniform_scale(grid_size[0], grid_size[1], grid_size[2]); + let world_transform_mat4: Mat4 = world_transform.matrix(); + let mut world_transform_tl: Mat4 = world_transform_mat4; + world_transform_tl.x[3] = 0.; + world_transform_tl.y[3] = 0.; + world_transform_tl.z[3] = 0.; + world_transform_tl.w[0] = 0.; + world_transform_tl.w[1] = 0.; + world_transform_tl.w[2] = 0.; + world_transform_tl.w[3] = 1.; + let normal_transform: Mat4 = world_transform_tl + .invert() + .unwrap_or_else(Mat4::identity) + .transpose(); + let inv_normal_transform: Mat4 = world_transform_tl.transpose(); + let world_scale: [f32; 3] = world_transform.scale.into(); + let uniform = SdfBufferData { + cube_transform: scale.into(), + inv_cube_transform: scale.invert().unwrap_or_else(Mat4::identity).into(), + world_transform: world_transform_mat4.into(), + inv_world_transform: world_transform_mat4 + .invert() + .unwrap_or_else(Mat4::identity) + .into(), + normal_transform: normal_transform.into(), + inv_normal_transform: inv_normal_transform.into(), + grid_dimensions: [grid_size[0], grid_size[1], grid_size[2], 1.], + world_scale: [world_scale[0], world_scale[1], world_scale[2], 1.], + }; + // println!("grid_dimensions: {:?}", uniform.grid_dimensions); + // println!("cube_transform: {:?}", uniform.cube_transform); + // println!("inv_cube_transform: {:?}", uniform.inv_cube_transform); + renderer.load_buffer(&mut sdf.data, bytemuck::cast_slice(&[uniform])); + + let reload_required = material_set.load(&renderer, &assets); + + if reload_required { + sdf.pipeline.bindings.unload(); + } + + if !sdf.pipeline.ready(&renderer) { + let lights_buffer = globals + .get::() + .expect("Light buffer must be loaded"); + + let sampler = globals.get::().expect("Sampler must be loaded"); + + if let Some(shader) = assets.get(sdf.pipeline.shader) { + renderer.bind( + &mut sdf.pipeline, + PipelineLayout::Render { + label: String::from(PIPELINE_LABEL), + mesh, + shader, + bindings: &[ + BindGroup::new( + "Globals", + vec![ + Binding::Uniform("Camera", Stage::All, &camera_buffer.uniform), + Binding::Sampler("Sampler", Stage::Fragment, sampler), + Binding::Storage( + "Lights", + Stage::Fragment, + &lights_buffer.storage, + ), + ], + ), + BindGroup::new( + "Locals", + vec![ + Binding::Uniform("Data", Stage::All, &sdf.data), + Binding::Texture3D("Sdf", Stage::All, &sdf.buffer), + Binding::Uniform( + "Materials", + Stage::All, + material_set.get_material_buffer(), + ), + Binding::TextureArray( + "MaterialTexture", + Stage::All, + material_set.get_texture_buffer(), + ), + ], + ), + ], + options: RenderOptions::default(), + }, + ); + } + } + + renderer.draw(&mut sdf.pipeline, mesh, &Default::default()); + } +} + +pub(super) fn extension(app: &mut Application) { + app.add_system(System::from(startup)); + app.add_system(System::from(render)); + camera::extension(app); + lights::extension(app); +} diff --git a/dotrix_voxel/src/sdf/circle_trace/accelerated_raytrace.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/accelerated_raytrace.inc.wgsl new file mode 100644 index 00000000..8cfba755 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/accelerated_raytrace.inc.wgsl @@ -0,0 +1,71 @@ +// Given: +// A `map` function that gives the distance to the sufrace from any point +// A direction (`rd`) and origin (`ro`) that represents the ray to be marched +// An intial distance already travelled (`t_in`) +// Ray differentails that represent the direction of neighbouring rays (`rdx`, `rdy`) +// Then: +// Compute the point at which this ray intersects the surface +// +// This implementation uses +// Accelerated raymarching +// https://www.researchgate.net/publication/331547302_Accelerating_Sphere_Tracing +// Which attempts to overstep on the ray in order to reduce the number of steps marched +// on the ray +// +struct RaymarchOut { + t: f32; + success: bool; +}; + +// Use pixel based cones to get the size of the pizel +// This is used for an early exit, given the directions and +// the distance traveled it computes the approixmate size of a pixel on screen +// If the distance to the surface is less then the returned pixel size then we +// Stop marching +fn pixel_radius(t: f32, direction: vec3, direction_x: vec3, direction_y: vec3) -> f32 { + let dx: f32 = length(t*(direction_x-direction)); + let dy: f32 = length(t*(direction_y-direction)); + return length(vec2(dx, dy)) * 0.1; // 10% of pixel size is the cut-off +} + +fn raymarch(t_in: f32, ro: vec3, rd: vec3, rdx: vec3, rdy: vec3) -> RaymarchOut { + let o: vec3 = ro; + let d: vec3 = rd; + let dx: vec3 = rdx; + let dy: vec3 = rdy; + + let STEP_SIZE_REDUCTION: f32 = 0.95; + let MAX_DISTANCE: f32 = t_in + length(u_sdf.grid_dimensions.xyz * abs(u_sdf.world_scale.xyz)); + let MAX_ITERATIONS: u32 = 128u; + + var t: f32 = t_in; + var rp: f32 = 0.; // prev + var rc: f32 = map(o + (t)*d);; // current + var rn: f32 = t + MAX_DISTANCE * 2.0; // next (set to effectivly infinity) + + var di: f32 = 0.; + + var out: RaymarchOut; + out.success = false; + + for(var i: u32 = 0u; i < MAX_ITERATIONS && t < MAX_DISTANCE; i = i + 1u) + { + di = rc + STEP_SIZE_REDUCTION * rc * max( (di - rp + rc) / (di + rp - rc), 0.6); + rn = map(o + (t + di)*d); + if(di > rc + rn) { + di = rc; + rn = map(o + (t + di)*d); + } + t = t + di; + out.t = t; + if(rn < pixel_radius(t, d, dx, dy)) { + out.success = true; + return out; + } + + rp = rc; + rc = rn; + } + + return out; +} diff --git a/dotrix_voxel/src/sdf/circle_trace/camera.rs b/dotrix_voxel/src/sdf/circle_trace/camera.rs new file mode 100644 index 00000000..f851c1f7 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/camera.rs @@ -0,0 +1,71 @@ +use dotrix_core::{ + ecs::{Const, Mut, System}, + renderer::Buffer, + Application, Camera, Globals, Renderer, Window, +}; +use dotrix_math::SquareMatrix; + +#[repr(C)] +#[derive(Default, Copy, Clone)] +pub(super) struct Uniform { + proj_view: [[f32; 4]; 4], + static_camera_trans: [[f32; 4]; 4], + pos: [f32; 4], + screen_resolution: [f32; 2], + fov: f32, + padding: [f32; 1], +} + +unsafe impl bytemuck::Zeroable for Uniform {} +unsafe impl bytemuck::Pod for Uniform {} + +pub struct CameraBuffer { + pub uniform: Buffer, +} + +impl Default for CameraBuffer { + fn default() -> Self { + Self { + uniform: Buffer::uniform("SDF Camera Buffer"), + } + } +} + +/// startup system +pub(super) fn startup(mut globals: Mut) { + globals.set(CameraBuffer::default()); +} + +/// startup system +pub(super) fn load( + renderer: Const, + camera: Const, + window: Const, + mut globals: Mut, +) { + let proj_mx = *camera.proj(); + let view_mx = camera.view_matrix(); + let static_camera_mx = camera.view_matrix_static().invert().unwrap(); + let camera_pos = camera.position(); + let inner_size = window.inner_size(); + let uniform = Uniform { + proj_view: (proj_mx * view_mx).into(), + static_camera_trans: static_camera_mx.into(), + pos: [camera_pos[0], camera_pos[1], camera_pos[2], 1.], + screen_resolution: [inner_size[0] as f32, inner_size[1] as f32], + fov: camera.fov, + padding: Default::default(), + }; + + if let Some(uniform_buffer) = globals.get_mut::() { + renderer.load_buffer( + &mut uniform_buffer.uniform, + bytemuck::cast_slice(&[uniform]), + ); + } +} + +pub(super) fn extension(app: &mut Application) { + app.add_system(System::from(startup)); + app.add_system(System::from(load)); +} diff --git a/dotrix_voxel/src/sdf/circle_trace/hemisphere_ambient_occulsion.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/hemisphere_ambient_occulsion.inc.wgsl new file mode 100644 index 00000000..e9153e39 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/hemisphere_ambient_occulsion.inc.wgsl @@ -0,0 +1,105 @@ +// This is a modified AO implementation. +// It attempt to probe the space around a point and determine +// How full that space is and therefore how occulded it is +// +// http://www.aduprat.com/portfolio/?page=articles/hemisphericalSDFAO +// +// It takes the standard concept of marching along the normal +// but expands it to march along multiple rays in a hemisphere +// like arrangment around the normal +// +// Uniform points on a hemisphere +// http://holger.dammertz.org/stuff/notes_HammersleyOnHemisphere.html +// +// The Hammersley function is used for uniform hemisphere points due to its +// Simplicilty on the gpu + +struct AoResult { + ao: f32; +}; + +struct AoInput { + origin: vec3; + direction: vec3; + samples: u32; + steps: u32; + ao_step_size: f32; +}; + +let PI: f32 = 3.14159265358979; +fn radicalInverse_VdC(in_bits: u32) -> f32 { + var bits: u32 = in_bits; + bits = (bits << 16u) | (bits >> 16u); + bits = ((bits & 0x55555555u) << 1u) | ((bits & 0xAAAAAAAAu) >> 1u); + bits = ((bits & 0x33333333u) << 2u) | ((bits & 0xCCCCCCCCu) >> 2u); + bits = ((bits & 0x0F0F0F0Fu) << 4u) | ((bits & 0xF0F0F0F0u) >> 4u); + bits = ((bits & 0x00FF00FFu) << 8u) | ((bits & 0xFF00FF00u) >> 8u); + return f32(bits) * 2.3283064365386963e-10; // / 0x100000000 +} +fn hammersley2d(i: u32, N: u32) -> vec2 { + return vec2(f32(i)/f32(N), radicalInverse_VdC(i)); +} +fn hemisphereSample_uniform(u: f32, v: f32) -> vec3 { + let phi: f32 = v * 2.0 * PI; + let cosTheta: f32 = 1.0 - u; + let sinTheta: f32 = sqrt(1.0 - cosTheta * cosTheta); + return vec3(cos(phi) * sinTheta, sin(phi) * sinTheta, cosTheta); +} + +fn ambient_occlusion(input: AoInput) -> AoResult +{ + let nb_ite: u32 = input.samples; + let nb_ite_inv: f32 = 1./f32(nb_ite); + let rad: f32 = 1. - (1. * nb_ite_inv); //Hemispherical factor (self occlusion correction) + + var ao: f32 = 0.0; + + // Tangent space tranformation + let a: vec3 = vec3(0., 0., 1.); + let b: vec3 = input.direction; + let v: vec3 = cross(a,b); + let s: f32 = length(v); + let I: mat3x3 = mat3x3(vec3(1.,0.,0.), vec3(0.,1.,0.), vec3(0.,0.,1.)); + var R: mat3x3; + if (abs(s) < 0.01) { + R = I; + } else { + let c: f32 = dot(a, b); + let sx: mat3x3 = mat3x3(vec3(0.,v.z,-v.y), vec3(-v.z,0.,v.x), vec3(v.y,-v.x,0.)); + + // R = I + sx + sx * sx * (1./(1. + c)); mat + mat addition is broken upstream /~https://github.com/gfx-rs/naga/issues/1527 + // Workaround start + let ISx: mat3x3 = mat3x3(I.x + sx.x, I.y + sx.y, I.z + sx.z); + let sxsx: mat3x3 = sx * sx * (1./(1. + c)); + + R = mat3x3(ISx.x + sxsx.x, ISx.y + sxsx.y, ISx.z + sxsx.z); + // Workaround end + } + + + + for( var i: u32 = 0u; i < nb_ite; i = i + 1u ) + { + let hammersley: vec2 = hammersley2d(i, nb_ite); + let rd = hemisphereSample_uniform(hammersley.x, hammersley.y); + + // In tangent space + let direction: vec3 = R * rd; + + // Stepping on the ray + var sum: f32 = 0.; + var max_sum: f32 = 0.; + for (var j: u32 = 0u; j < input.steps; j = j + 1u) + { + let p: vec3 = input.origin + direction * f32(j + 1u) * input.ao_step_size; + sum = sum + 1. / pow(2., f32(j)) * max(map(p), 0.); + max_sum = max_sum + 1. / pow(2., f32(j)) * f32(j + 1u) * input.ao_step_size; + } + + ao = ao + (sum / max_sum) / f32(nb_ite); + } + + var ray_out: AoResult; + ray_out.ao = clamp(ao, 0., 1.); + return ray_out; +} diff --git a/dotrix_voxel/src/sdf/circle_trace/lighting.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/lighting.inc.wgsl new file mode 100644 index 00000000..075c83c3 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/lighting.inc.wgsl @@ -0,0 +1,161 @@ +struct LightCalcOutput { + light_direction: vec3; + radiance: vec3; +}; + +struct DirectionalLight { + direction: vec4; + color: vec4; +}; + +struct PointLight { + position: vec4; + color: vec4; + attenuation: vec4; + // attenuation + // a_constant: f32; + // a_linear: f32; + // a_quadratic: f32; + // unused: f32; +}; + +struct SimpleLight { + position: vec4; + color: vec4; +}; + +struct SpotLight { + position: vec4; + direction: vec4; + color: vec4; + cut_off: f32; + outer_cut_off: f32; +}; + +struct GenericLight { + position: vec4; + direction: vec4; + color: vec4; + parameters: vec4; + kind: u32; // 1 = DirectionalLight, 2 = PointLight, 3 = SimpleLight, 4 = SpotLight, 0 = None +}; + +struct Lights { + generic_lights: array; +}; + +[[group(0), binding(2)]] +var s_lights: Lights; + +fn calculate_directional( + light: DirectionalLight, +) -> LightCalcOutput { + let light_direction: vec3 = normalize(-light.direction.xyz); + + var out: LightCalcOutput; + out.light_direction = light_direction; + out.radiance = light.color.rgb; + return out; +} + + +fn calculate_point( + light: PointLight, + position: vec3, +) -> LightCalcOutput { + let light_direction: vec3 = normalize(light.position.xyz - position); + + let light_distance: f32 = length(light.position.xyz - position.xyz); + let attenuation: f32 = 1.0 / ( + light.attenuation.x + + light.attenuation.y * light_distance + + light.attenuation.z * (light_distance * light_distance) + ); + + var out: LightCalcOutput; + out.light_direction = light_direction; + out.radiance = light.color.rgb * attenuation; + return out; +} + +fn calculate_simple( + light: SimpleLight, + position: vec3, +) -> LightCalcOutput { + let light_direction: vec3 = normalize(light.position.xyz - position.xyz); + + var out: LightCalcOutput; + out.light_direction = light_direction; + out.radiance = light.color.rgb; + return out; +} + + +fn calculate_spot( + light: SpotLight, + position: vec3, +) -> LightCalcOutput { + let light_direction: vec3 = normalize(light.position.xyz - position.xyz); + let theta: f32 = dot(light_direction, normalize(-light.direction.xyz)); + + let epsilon: f32 = light.cut_off - light.outer_cut_off; + let intensity: f32 = clamp((theta - light.outer_cut_off) / epsilon, 0.0, 1.0); + + var out: LightCalcOutput; + out.light_direction = light_direction; + out.radiance = light.color.rgb * intensity; + return out; +} + +fn calculate_light_ray_for( + camera_index: u32, + position: vec3, +) -> LightCalcOutput { + var generic_light: GenericLight = s_lights.generic_lights[camera_index]; + switch (generic_light.kind) { + case 1: { + var light: DirectionalLight; + light.direction = generic_light.direction; + light.color = generic_light.color; + return calculate_directional(light); + } + case 2: { + var light: PointLight; + light.position = generic_light.position; + light.color = generic_light.color; + light.attenuation = generic_light.parameters; + return calculate_point(light, position); + } + case 3: { + var light: SimpleLight; + light.position = generic_light.position; + light.color = generic_light.color; + return calculate_simple(light, position); + } + case 4: { + var light: SpotLight; + light.direction = generic_light.direction; + light.position = generic_light.position; + light.color = generic_light.color; + light.cut_off = generic_light.parameters.x; + light.outer_cut_off = generic_light.parameters.y; + return calculate_spot(light, position); + } + default: { + var out: LightCalcOutput; + out.light_direction = vec3(0.); + out.radiance = vec3(0.); + return out; + } + } +} + +// Ambient is stored in the last light +fn get_light_count() -> u32 { + return arrayLength(&s_lights.generic_lights) - 1u; +} + +fn get_ambient() -> vec3 { + let idx: u32 = u32(arrayLength(&s_lights.generic_lights)) - 1u; + return s_lights.generic_lights[idx].color.xyz; +} diff --git a/dotrix_voxel/src/sdf/circle_trace/lights.rs b/dotrix_voxel/src/sdf/circle_trace/lights.rs new file mode 100644 index 00000000..54e46105 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/lights.rs @@ -0,0 +1,12 @@ +use dotrix_core::{ecs::System, Application}; + +mod data; +mod storage; + +pub use data::Light; +pub(super) use storage::LightStorageBuffer; + +pub(super) fn extension(app: &mut Application) { + app.add_system(System::from(storage::load)); + app.add_system(System::from(storage::startup)); +} diff --git a/dotrix_voxel/src/sdf/circle_trace/lights/data.rs b/dotrix_voxel/src/sdf/circle_trace/lights/data.rs new file mode 100644 index 00000000..6882ff5c --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/lights/data.rs @@ -0,0 +1,99 @@ +/// Light component of different types and settings +use super::storage::GenericLight; +pub use dotrix_pbr::Light; + +/// Directional light uniform data +pub(super) struct DirectionalLight { + /// Light direction + pub(super) direction: [f32; 4], + /// Light color + pub(super) color: [f32; 4], +} + +impl From for GenericLight { + fn from(src: DirectionalLight) -> Self { + Self { + position: Default::default(), + direction: src.direction, + color: src.color, + parameters: [0., 0., 0., 0.], + kind: 1, + padding: Default::default(), + } + } +} + +/// Point light uniform data +pub(super) struct PointLight { + /// Light source position + pub(super) position: [f32; 4], + /// Light color + pub(super) color: [f32; 4], + /// Constant light attenuation + pub(super) a_constant: f32, + /// Linear light attenuation + pub(super) a_linear: f32, + /// Quadratic light attenuation + pub(super) a_quadratic: f32, +} + +impl From for GenericLight { + fn from(src: PointLight) -> Self { + Self { + position: src.position, + direction: Default::default(), + color: src.color, + parameters: [src.a_constant, src.a_linear, src.a_quadratic, 0.], + kind: 2, + padding: Default::default(), + } + } +} + +/// Simple light uniform data +pub(super) struct SimpleLight { + /// Light source position + pub(super) position: [f32; 4], + /// Light color + pub(super) color: [f32; 4], +} + +impl From for GenericLight { + fn from(src: SimpleLight) -> Self { + Self { + position: src.position, + direction: Default::default(), + color: src.color, + parameters: Default::default(), + kind: 3, + padding: Default::default(), + } + } +} + +/// Spot Light uniform data +pub(super) struct SpotLight { + /// Light source position + pub(super) position: [f32; 4], + /// Light source direction + pub(super) direction: [f32; 4], + /// Light source color + pub(super) color: [f32; 4], + /// Light source cut off + pub(super) cut_off: f32, + /// Light source outer cut off + pub(super) outer_cut_off: f32, +} + +impl From for GenericLight { + fn from(src: SpotLight) -> Self { + Self { + position: src.position, + direction: src.direction, + color: src.color, + parameters: [src.cut_off, src.outer_cut_off, 0., 0.], + kind: 4, + padding: Default::default(), + } + } +} diff --git a/dotrix_voxel/src/sdf/circle_trace/lights/storage.rs b/dotrix_voxel/src/sdf/circle_trace/lights/storage.rs new file mode 100644 index 00000000..19dfc578 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/lights/storage.rs @@ -0,0 +1,134 @@ +// Originally from /~https://github.com/lowenware/dotrix/blob/4582922bbdd8bd039857271f8c8eb4cfd42fcb00/dotrix_pbr/src/light.rs +// +use super::data::*; +use dotrix_core::{ + ecs::{Const, Mut}, + renderer::Buffer, + Globals, Renderer, World, +}; + +#[repr(C)] +#[derive(Default, Copy, Clone)] +pub(super) struct GenericLight { + pub(super) position: [f32; 4], + pub(super) direction: [f32; 4], + pub(super) color: [f32; 4], + pub(super) parameters: [f32; 4], + pub(super) kind: u32, // 1 = DirectionalLight, 2 = PointLight, 3 = SimpleLight, 4 = SpotLight, 0 = None + pub(super) padding: [f32; 3], +} + +unsafe impl bytemuck::Zeroable for GenericLight {} +unsafe impl bytemuck::Pod for GenericLight {} + +pub struct LightStorageBuffer { + pub storage: Buffer, +} + +impl Default for LightStorageBuffer { + fn default() -> Self { + Self { + storage: Buffer::storage("Light Storage"), + } + } +} + +pub fn startup(mut globals: Mut) { + globals.set(LightStorageBuffer::default()); +} + +/// Lights binding system +pub fn load(world: Const, renderer: Const, mut globals: Mut) { + if let Some(lights) = globals.get_mut::() { + let mut generic_lights: Vec = world + .query::<(&Light,)>() + .flat_map(|(light,)| match light { + Light::Directional { + color, + direction, + intensity, + enabled, + } if *enabled => Some( + DirectionalLight { + direction: [direction.x, direction.y, direction.z, 1.0], + color: (*color * (*intensity)).into(), + } + .into(), + ), + Light::Point { + color, + position, + intensity, + enabled, + constant, + linear, + quadratic, + } if *enabled => Some( + PointLight { + position: [position.x, position.y, position.z, 1.0], + color: (*color * (*intensity)).into(), + a_constant: *constant, + a_linear: *linear, + a_quadratic: *quadratic, + } + .into(), + ), + Light::Simple { + color, + position, + intensity, + enabled, + } if *enabled => Some( + SimpleLight { + position: [position.x, position.y, position.z, 1.0], + color: (*color * (*intensity)).into(), + } + .into(), + ), + Light::Spot { + color, + position, + direction, + intensity, + enabled, + cut_off, + outer_cut_off, + } if *enabled => Some( + SpotLight { + position: [position.x, position.y, position.z, 1.0], + direction: [direction.x, direction.y, direction.z, 1.0], + color: (*color * (*intensity)).into(), + cut_off: *cut_off, + outer_cut_off: *outer_cut_off, + } + .into(), + ), + _ => None, + }) + .collect(); + + let ambient: GenericLight = world + .query::<(&Light,)>() + .flat_map(|(light,)| match light { + Light::Ambient { color, intensity } => Some(GenericLight { + color: (*color * (*intensity)).into(), + ..Default::default() + }), + _ => None, + }) + .fold(GenericLight::default(), |mut ambient, light| { + ambient.color[0] += light.color[0]; + ambient.color[1] += light.color[1]; + ambient.color[2] += light.color[2]; + ambient.color[3] += light.color[3]; + ambient + }); + + generic_lights.push(ambient); + + renderer.load_buffer( + &mut lights.storage, + bytemuck::cast_slice(generic_lights.as_slice()), + ); + } +} diff --git a/dotrix_voxel/src/sdf/circle_trace/map.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/map.inc.wgsl new file mode 100644 index 00000000..232e9cbe --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/map.inc.wgsl @@ -0,0 +1,157 @@ +/// This function takes a position in world space, consults the +/// sdf and returns the distance to the surface +/// +/// This varient works on SDFs stored in 3D textures + + +// SDF for a box, b is half box size +fn sdBox(p: vec3,b: vec3) -> f32 +{ + let q: vec3 = abs(p) - b; + return length(max(q,vec3(0.0))) + min(max(q.x,max(q.y,q.z)),0.0); +} + +// Given a location in the texture of the kind x = (-0.5,0.5), y = (-0.5,0.5), z = (-0.5,0.5) +// i.e. a coordinate inside a box of unit length center at the origin +// Return the interpolated value +fn tri_linear_interpolation(cube_p: vec4) -> f32 { + let seed_dim: vec3 = textureDimensions(sdf_texture) - vec3(1); + let seed_dim_f32: vec3 = vec3(f32(seed_dim.x), f32(seed_dim.y), f32(seed_dim.z)); + let pixel_coords: vec3 = ((cube_p.xyz + vec3(0.5)) /2.) * seed_dim_f32 * 2.; + + var i: i32 = i32(pixel_coords.x); + var j: i32 = i32(pixel_coords.y); + var k: i32 = i32(pixel_coords.z); + + + i = clamp(i, 0, seed_dim.x); + j = clamp(j, 0, seed_dim.y); + k = clamp(k, 0, seed_dim.z); + let x: f32 = clamp(pixel_coords.x - f32(i), 0., 1.); + let y: f32 = clamp(pixel_coords.y - f32(j), 0., 1.); + let z: f32 = clamp(pixel_coords.z - f32(k), 0., 1.); + + let f000: f32 = textureLoad(sdf_texture, + vec3( + i, + j, + k, + ) + ,0).r; + let f001: f32 = textureLoad(sdf_texture, + vec3( + i, + j, + k + 1, + ) + ,0).r; + let f010: f32 = textureLoad(sdf_texture, + vec3( + i, + j + 1, + k, + ) + ,0).r; + let f011: f32 = textureLoad(sdf_texture, + vec3( + i, + j + 1, + k + 1, + ) + ,0).r; + let f100: f32 = textureLoad(sdf_texture, + vec3( + i + 1, + j, + k, + ) + ,0).r; + let f101: f32 = textureLoad(sdf_texture, + vec3( + i + 1, + j, + k + 1, + ) + ,0).r; + let f110: f32 = textureLoad(sdf_texture, + vec3( + i + 1, + j + 1, + k, + ) + ,0).r; + let f111: f32 = textureLoad(sdf_texture, + vec3( + i + 1, + j + 1, + k + 1, + ) + ,0).r; + + return ( + f000*(1.-x)*(1.-y)*(1.-z) + +f001*(1.-x)*(1.-y)*z + +f010*(1.-x)*y *(1.-z) + +f011*(1.-x)*y *z + +f100*x *(1.-y)*(1.-z) + +f101*x *(1.-y)*z + +f110*x *y *(1.-z) + +f111*x *y *z + ); + +} + +// Get distance to surface from a point in world space +fn map(p: vec3) -> f32 +{ + let local_p: vec4 = (u_sdf.inv_world_transform * vec4(p, 1.)); + let cube_p: vec4 = (u_sdf.inv_cube_transform * local_p); + + let internal_dist = tri_linear_interpolation(cube_p); + + // Enclosing box used for clipping + let enclosing_box: f32 = sdBox(local_p.xyz, (u_sdf.grid_dimensions.xyz * 1.00)/vec3(2.)); + + // + // Distance are built on the assumption that voxel size is one + // we must correct that + // If scale is non_uniform we can only provide a bound on the distance + let scale: vec3 = u_sdf.world_scale.xyz; + let min_scale: f32 = min(abs(scale.x), min(abs(scale.y), abs(scale.z))); + + // Return intersection of voxel sdf and enclosing (clipping) box + let dist = max(internal_dist, enclosing_box) * min_scale; + return dist; +} +// Get the material id at a point +// +// Using nearest neighbour +fn map_material(p: vec3) -> u32 +{ + let local_p: vec4 = (u_sdf.inv_world_transform * vec4(p, 1.)); + let cube_p: vec4 = (u_sdf.inv_cube_transform * local_p); + + let seed_dim: vec3 = textureDimensions(sdf_texture) - vec3(1); + let seed_dim_f32: vec3 = vec3(f32(seed_dim.x), f32(seed_dim.y), f32(seed_dim.z)); + let pixel_coords: vec3 = ((cube_p.xyz + vec3(0.5)) /2.) * seed_dim_f32 * 2.; + + let nearest_pos: vec3 = round(pixel_coords); + let nearest_coord: vec3 = vec3(i32(nearest_pos.x), i32(nearest_pos.y), i32(nearest_pos.z)); + + return u32(textureLoad(sdf_texture, nearest_coord, 0).g); +} + +// Surface gradient (is the normal) +fn map_normal (p: vec3) -> vec3 +{ + let eps: vec3 = abs(u_sdf.world_scale.xyz) * 0.05; + + return normalize + ( vec3 + ( + map(p + vec3(eps.x, 0., 0.) ) - map(p - vec3(eps.x, 0., 0.)), + map(p + vec3(0., eps.y, 0.) ) - map(p - vec3(0., eps.y, 0.)), + map(p + vec3(0., 0., eps.z) ) - map(p - vec3(0., 0., eps.z)) + ) + ); +} diff --git a/dotrix_voxel/src/sdf/circle_trace/pbr.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/pbr.inc.wgsl new file mode 100644 index 00000000..950e1536 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/pbr.inc.wgsl @@ -0,0 +1,115 @@ +fn distribution_ggx(normal: vec3, halfway: vec3, roughness: f32) -> f32 +{ + let a: f32 = roughness*roughness; + let a2: f32 = a*a; + let n_dot_h: f32 = max(dot(normal, halfway), 0.0); + let n_dot_h_2: f32 = n_dot_h*n_dot_h; + + let num: f32 = a2; + var denom: f32 = (n_dot_h_2 * (a2 - 1.0) + 1.0); + denom = PI * denom * denom; + + return num / denom; +} + +fn geometry_schlick_ggx(n_dot_v: f32, roughness: f32) -> f32 +{ + let r: f32 = (roughness + 1.0); + let k: f32 = (r*r) / 8.0; + + let num: f32 = n_dot_v; + let denom: f32 = n_dot_v * (1.0 - k) + k; + + return num / denom; +} +fn geometry_smith(normal: vec3, camera_direction: vec3, light_direction: vec3, roughness: f32) -> f32 +{ + let n_dot_v: f32 = max(dot(normal, camera_direction), 0.0); + let n_dot_l: f32 = max(dot(normal, light_direction), 0.0); + let ggx2: f32 = geometry_schlick_ggx(n_dot_v, roughness); + let ggx1: f32 = geometry_schlick_ggx(n_dot_l, roughness); + + return ggx1 * ggx2; +} + +// Calulates the amount of light that refects (specular) and that which scatters (diffuse) +fn calculate_fresnel_schlick(cos_theta: f32, fresnel_schlick_0: vec3) -> vec3 +{ + return fresnel_schlick_0 + (vec3(1.0) - fresnel_schlick_0) * pow(clamp(1.0 - cos_theta, 0.0, 1.0), 5.0); +} + +fn pbr( + light_out: LightCalcOutput, + camera_direction: vec3, // Camera direction + normal: vec3, // normal + fresnel_schlick_0: vec3, // surface reflection at zero incidence + albedo: vec3, // scatter color in linear space + metallic: f32, // Metallic (reflectance) + roughness: f32 // Roughness (random scatter) +) -> vec3 { + let light_direction: vec3 = light_out.light_direction; + let halfway: vec3 = normalize(camera_direction + light_direction); + + // cook-torrance brdf + let normal_distribution_function: f32 = distribution_ggx(normal, halfway, roughness); + let geometry_function: f32 = geometry_smith(normal, camera_direction, light_direction, roughness); + let fresnel_schlick: vec3 = calculate_fresnel_schlick(max(dot(halfway, camera_direction), 0.0), fresnel_schlick_0); + + let reflection_specular_fraction: vec3 = fresnel_schlick; + var refraction_diffuse_fraction: vec3 = vec3(1.0); // - reflection_specular_fraction; // refraction/diffuse fraction + refraction_diffuse_fraction = refraction_diffuse_fraction * (1.0 - metallic); + + let numerator: vec3 = normal_distribution_function * geometry_function * fresnel_schlick; + let denominator: f32 = 4.0 * max(dot(normal, camera_direction), 0.0) * max(dot(normal, light_direction), 0.0) + 0.0001; + let specular: vec3 = numerator / denominator; + + // get the outgoing radiance + let n_dot_l: f32 = max(dot(normal, light_direction), 0.0); + return (refraction_diffuse_fraction * albedo / PI + specular) * light_out.radiance * n_dot_l; + // return light_out.radiance * (1.0 - metallic); +} + +fn calculate_lighting( + position: vec3, + normal_in: vec3, + albedo: vec3, + roughness: f32, + metallic: f32, + ao: f32, +) -> vec4 { + let camera_position: vec3 = u_camera.pos.xyz; + var light_color: vec3 = vec3(0.); + + let normal: vec3 = normalize(normal_in ); + let camera_direction: vec3 = normalize(camera_position - position); + + var fresnel_schlick_0: vec3 = vec3(0.04); + fresnel_schlick_0 = mix(fresnel_schlick_0, albedo, vec3(metallic)); + + var i: u32 = 0u; + var count: u32 = get_light_count(); + for (i = 0u; i< count; i = i + 1u) { + let light_result = calculate_light_ray_for(i, position); + light_color = light_color + + // light_result.radiance; + pbr( + light_result, + camera_direction, + normal, + fresnel_schlick_0, + albedo, + metallic, + roughness + ); + } + + // Ambient + let ambient = get_ambient() * albedo * ao; + light_color = light_color + ambient; + + // Gamma correct + light_color = light_color / (light_color + vec3(1.0)); + light_color = pow(light_color, vec3(1.0/2.2)); + + return vec4(light_color, 1.0); +} diff --git a/dotrix_voxel/src/sdf/circle_trace/render.wgsl b/dotrix_voxel/src/sdf/circle_trace/render.wgsl new file mode 100644 index 00000000..7e619baf --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/render.wgsl @@ -0,0 +1,226 @@ +struct Camera { + proj_view: mat4x4; + static_camera_trans: mat4x4; + pos: vec4; + screen_resolution: vec2; + fov: f32; +}; +[[group(0), binding(0)]] +var u_camera: Camera; + +[[group(0), binding(1)]] +var r_sampler: sampler; + +struct SdfData { + // This transform scales the 1x1x1 cube so that it totally encloses the + // voxels + cube_transform: mat4x4; + // Inverse cube_transform + inv_cube_transform: mat4x4; + // World transform of the voxel grid + world_transform: mat4x4; + // Inverse World transform of the voxel grid + inv_world_transform: mat4x4; + // Matrix to convert objectspace normal to world space + normal_transform: mat4x4; + // Matrix to convert world space normal to object space + inv_normal_transform: mat4x4; + // Dimensions of the voxel + grid_dimensions: vec4; + // Scale in world space + world_scale: vec4; +}; +[[group(1), binding(0)]] +var u_sdf: SdfData; + +[[group(1), binding(1)]] +var sdf_texture: texture_3d; + +struct Material { + albedo_id: i32; + roughness_id: i32; + metallic_id: i32; + ao_id: i32; + normal_id: i32; + albedo: vec4; + roughness: f32; + metallic: f32; + ao: f32; +}; +struct Materials { + materials: [[stride(64)]] array; +}; +[[group(1), binding(2)]] +var u_materials: Materials; + +[[group(1), binding(3)]] +var material_textures: texture_2d_array; + + +struct VertexOutput { + [[builtin(position)]] position: vec4; + [[location(0)]] world_position: vec3; + [[location(1)]] clip_coords: vec4; +}; + +[[stage(vertex)]] +fn vs_main( + [[location(0)]] position: vec3, +) -> VertexOutput { + let pos_4: vec4 = vec4(position, 1.); + let local_coords: vec4 = u_sdf.cube_transform * pos_4; + let world_coords: vec4 = u_sdf.world_transform * local_coords; + let clip_coords: vec4 = u_camera.proj_view * world_coords; + + var out: VertexOutput; + out.position = clip_coords; + out.world_position = world_coords.xyz; + out.clip_coords = clip_coords; + return out; +} + +// Given pixel coordinates get the ray direction +fn get_ray_direction(pixel: vec2, resolution: vec2) -> vec3 { + let pixel_f32: vec2 = vec2(f32(pixel.x), f32(pixel.y)); + let p: vec2 = (2.0 * pixel_f32 - resolution.xy)/(resolution.y); + let z: f32 = -1. / tan(u_camera.fov * 0.5); + let view_coordinate: vec4 = vec4(p.x, p.y, z, 1.); + let world_coordinate: vec4 = u_camera.static_camera_trans * view_coordinate; + + return normalize(world_coordinate.xyz); +} + +{% include "circle_trace/map.inc.wgsl" %} + +{% include "circle_trace/accelerated_raytrace.inc.wgsl" %} + +{% include "circle_trace/hemisphere_ambient_occulsion.inc.wgsl" %} + +{% include "circle_trace/soft_shadows_closet_approach.inc.wgsl" %} + +{% include "circle_trace/lighting.inc.wgsl" %} + +{% include "circle_trace/pbr.inc.wgsl" %} + +{% include "circle_trace/triplanar_surface.inc.wgsl" %} + + +struct FragmentOutput { + [[location(0)]] color: vec4; + [[builtin(frag_depth)]] depth: f32; +}; + +[[stage(fragment)]] +fn fs_main(in: VertexOutput) -> FragmentOutput { + let debug: bool = false; + let resolution: vec2 = u_camera.screen_resolution.xy; + let pixel_coords: vec2 = vec2(u32(in.position.x), u32(resolution.y - in.position.y)); + + let ro: vec3 = u_camera.pos.xyz; + + // This can also be achieved by using world coords but we + // do it in pixels coords to get the pixel differentials + let rd: vec3 = get_ray_direction(pixel_coords.xy, resolution); + let rdx: vec3 = get_ray_direction(pixel_coords.xy + vec2(1u, 0u), resolution); + let rdy: vec3 = get_ray_direction(pixel_coords.xy + vec2(0u, 1u), resolution); + // let rd: vec3 = normalize(in.world_position - ro); // Use world_coords instead + + // Current distance from camera to grid + let t: f32 = length(in.world_position - ro); + + // March that ray + let r_out: RaymarchOut = raymarch(t, ro, rd, rdx, rdy); + let t_out: f32 = r_out.t; + if (!debug && !r_out.success) { + discard; + } + + // Final position of the ray + let pos: vec3 = ro + rd*t_out; + + // Normal of the surface + let nor: vec3 = map_normal(pos); + + // AO + var ray_in: AoInput; + ray_in.origin = ro; + ray_in.direction = nor; + ray_in.samples = 32u; + ray_in.steps = 8u; + ray_in.ao_step_size = 0.01; + + let ao: f32 = 1. - clamp(0., .1, ambient_occlusion(ray_in).ao); + + // Shadows + var total_radiance: vec3 = vec3(0.); + total_radiance = total_radiance + get_ambient(); + + let light_count: u32 = get_light_count(); + for (var i: u32 = 0u; i = light_out.light_direction; + + let intensity: f32 = dot(light_out.light_direction, nor); + // If perpendicular don't bother (numerically unstable) + if (abs(intensity) > 0.1 ) { + var ray_in: SoftShadowInput; + ray_in.origin = pos; + ray_in.direction = light_out.light_direction; + ray_in.max_iterations = 128u; + ray_in.min_distance = 0.01; + ray_in.max_distance = 100.; + ray_in.k = 8.; + + let ray_out: SoftShadowResult = softshadow(ray_in); + + total_radiance = total_radiance + intensity * ray_out.radiance; + } + } + total_radiance = clamp(vec3(0.), vec3(1.), total_radiance); + + // TODO: Work out how to bind textures effectivly + // // Ray differntials + // let dp_dxy: DpDxy = calcDpDxy( ro, rd, rdx, rdy, t, nor ); + // + // Material ID + let material_id: u32 = u32(map_material(pos)); + var material_data: Material = u_materials.materials[material_id]; + // + + // Partial derivates + let dpdx = t_out*(rdx*dot(rd,nor)/dot(rdx,nor) - rd); + let dpdy = t_out*(rdy*dot(rd,nor)/dot(rdy,nor) - rd); + // // Surface material + let sur: Surface = get_surface(pos, nor, material_id, dpdx, dpdy); + // + // // Lighting and PBR + let shaded: vec4 = calculate_lighting( + pos, + sur.normal, + sur.albedo.rgb, + sur.roughness, + sur.metallic, + sur.ao, + ); + + var out: FragmentOutput; + if (r_out.success) { + out.color = vec4(total_radiance, 1.) * shaded; + // out.color = vec4(total_radiance, 1.) * vec4(sur.albedo, 1. ); + // out.color = vec4(total_radiance, 1.) * vec4(vec3(sur.roughness), 1. ); + // out.color = vec4(total_radiance, 1.) * vec4(vec3(sur.metallic), 1. ); + // out.color = vec4(total_radiance, 1.) * vec4(vec3(sur.ao), 1. ); + // out.color = vec4(sur.albedo, 1. ); + // out.color = vec4(vec3(sur.metallic), 1. ); + // out.color = vec4(abs(sur.normal), 1.); + // out.color = material_data.albedo; + // out.color = vec4(vec3(material_data.metallic), 1.); + // out.color = shaded; + } else { + out.color = vec4(0.5,0.1,0.1,1.0); + } + + let pos_clip: vec4 = u_camera.proj_view * vec4(in.world_position, 1.); + out.depth = pos_clip.z / pos_clip.w; + return out; +} diff --git a/dotrix_voxel/src/sdf/circle_trace/soft_shadows_closet_approach.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/soft_shadows_closet_approach.inc.wgsl new file mode 100644 index 00000000..328a1d3c --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/soft_shadows_closet_approach.inc.wgsl @@ -0,0 +1,71 @@ +// Soft shadows as suggested by +// Sebastian Aaltonen et al. Soft Shadows at his GDC presentation +// +// Marches a ray from a point towards a single light +// If that ray hits the surface then radience will be none +// +// If that ray goes near but does not hit a surface then we are in soft (partial) +// shadow +// +// Calcualation of near surface it simple because of the SDF map +// +// Sebastian Aaltonen et al. provides an improved implementation of nearmiss +// Soft shadows with fewer artifacts that more accuratly works out the closest +// approach point of a ray +// +// +struct SoftShadowResult { + radiance: f32; +}; + +struct SoftShadowInput { + origin: vec3; + direction: vec3; + max_iterations: u32; + min_distance: f32; + max_distance: f32; + k: f32; +}; + +fn softshadow (input: SoftShadowInput) -> SoftShadowResult +{ + let o: vec3 = input.origin; + let d: vec3 = input.direction; + + var di: f32 = 0.; + var t: f32 = input.min_distance; + + let STEP_SIZE_REDUCTION: f32 = 0.95; + var rp: f32 = 0.; // prev + var rc: f32 = 0.; // current large such that y=0.0 at first + var rn: f32 = map(o + (t)*d); // next + + var radiance: f32 = 1.; + + for(var i: u32 = 0u; i < input.max_iterations && t < input.max_distance; i = i + 1u) + { + let y: f32 = rn*rn/(2.0*rc); + let approx_distance: f32 = sqrt(rn*rn-y*y); + radiance = min(radiance, input.k * approx_distance/max(0.0,t-y)); + + di = rc + STEP_SIZE_REDUCTION * rc * max( (di - rp + rc) / (di + rp - rc), 0.6); + rn = map(o + (t + di)*d); + if(di > rc + rn) + { + di = rc; + rn = map(o + (t + di)*d); + } + // if(rn < 0.001) { + // var out: SoftShadowResult; + // out.radiance = 0.; + // return out; + // } + t = t + di; + + rp = rc; + rc = rn; + } + var out: SoftShadowResult; + out.radiance = radiance; + return out; +} diff --git a/dotrix_voxel/src/sdf/circle_trace/triplanar_surface.inc.wgsl b/dotrix_voxel/src/sdf/circle_trace/triplanar_surface.inc.wgsl new file mode 100644 index 00000000..f5291460 --- /dev/null +++ b/dotrix_voxel/src/sdf/circle_trace/triplanar_surface.inc.wgsl @@ -0,0 +1,177 @@ +//! This get the material properties at a point on the surface +//! +//! This uses triplanar rendering which means we blend a +//! up front and right texture according to the normal +//! +//! Triplanar math is from https://www.volume-gfx.com/volume-rendering/triplanar-texturing/ +//! +//! Normal math is from https://bgolus.medium.com/normal-mapping-for-a-triplanar-shader-10bf39dca05a +//! + +struct Surface { + normal: vec3; + albedo: vec3; + roughness: f32; + metallic: f32; + ao: f32; +}; + +fn average(input: vec4) -> f32 { + return (input.x + input.y + input.z + input.w) / 4.; +} + +// Get the surface using triplanar texturing +// This is the final materials values at a point +fn get_surface(world_pos: vec3, world_normal: vec3, material_id: u32, ddx: vec3, ddy: vec3) -> Surface { + // Convert normal to local space + let local_normal: vec3 = normalize((u_sdf.inv_normal_transform * vec4(world_normal, 0.)).xyz); + // Convert position to local space without scale + // This is so that materials don't stretch over the voxel and a large voxel will have + // more repeating material + // TODO: Make this configurable + let local_pos: vec3 = (u_sdf.inv_world_transform * vec4(world_pos, 1.)).xyz * u_sdf.world_scale.xyz; + + // Tri planar parameters + let delta: f32 = 1e-3; // Controls plateua where it is considered to be one plane + let m: f32 = 1.; // Controls triplanar transision speed + + // Tile over 1x1x1 space + let x: f32 = fract(local_pos.x); + let y: f32 = fract(local_pos.y); + let z: f32 = fract(local_pos.z); + + + // Triplanar weights + let b: vec3 = normalize(vec3( + pow(max(abs(local_normal.x) - delta, 0.), m), + pow(max(abs(local_normal.y) - delta, 0.), m), + pow(max(abs(local_normal.z) - delta, 0.), m) + )); + + var material_right: Material = u_materials.materials[material_id]; + var material_top: Material = u_materials.materials[material_id]; + var material_front: Material = u_materials.materials[material_id]; + + var out: Surface; + + // Albedo + var albedo_right: vec4; + var albedo_top: vec4; + var albedo_front: vec4; + if (material_right.albedo_id < 0) { + albedo_right = material_right.albedo; + } else { + albedo_right = textureSampleGrad(material_textures, r_sampler, vec2(y,z), material_right.albedo_id, ddx.yz, ddy.yz); + } + if (material_top.albedo_id < 0) { + albedo_top = material_top.albedo; + } else { + albedo_top = textureSampleGrad(material_textures, r_sampler, vec2(x,z), material_top.albedo_id, ddx.xz, ddy.xz); + } + if (material_front.albedo_id < 0) { + albedo_front = material_front.albedo; + } else { + albedo_front = textureSampleGrad(material_textures, r_sampler, vec2(x,y), material_front.albedo_id, ddx.xy, ddy.xy); + } + out.albedo = albedo_right.xyz * b.x + albedo_top.xyz * b.y + albedo_front.xyz * b.z; + + // Roughness + var roughness_right: f32; + var roughness_top: f32; + var roughness_front: f32; + if (material_right.roughness_id < 0) { + roughness_right = material_right.roughness; + } else { + roughness_right = average(textureSampleGrad(material_textures, r_sampler, vec2(y,z), material_right.roughness_id, ddx.yz, ddy.yz)); + } + if (material_top.roughness_id < 0) { + roughness_top = material_top.roughness; + } else { + roughness_top = average(textureSampleGrad(material_textures, r_sampler, vec2(x,z), material_top.roughness_id, ddx.xz, ddy.xz)); + } + if (material_front.roughness_id < 0) { + roughness_front = material_front.roughness; + } else { + roughness_front = average(textureSampleGrad(material_textures, r_sampler, vec2(x,y), material_front.roughness_id, ddx.xy, ddy.xy)); + } + out.roughness = roughness_right * b.x + roughness_top * b.y + roughness_front * b.z; + + // Metallic + var metallic_right: f32; + var metallic_top: f32; + var metallic_front: f32; + if (material_right.metallic_id < 0) { + metallic_right = material_right.metallic; + } else { + metallic_right = average(textureSampleGrad(material_textures, r_sampler, vec2(y,z), material_right.metallic_id, ddx.yz, ddy.yz)); + } + if (material_top.metallic_id < 0) { + metallic_top = material_top.metallic; + } else { + metallic_top = average(textureSampleGrad(material_textures, r_sampler, vec2(x,z), material_top.metallic_id, ddx.xz, ddy.xz)); + } + if (material_front.metallic_id < 0) { + metallic_front = material_front.metallic; + } else { + metallic_front = average(textureSampleGrad(material_textures, r_sampler, vec2(x,y), material_front.metallic_id, ddx.xy, ddy.xy)); + } + out.metallic = metallic_right * b.x + metallic_top * b.y + metallic_front * b.z; + + // Ao + var ao_right: f32; + var ao_top: f32; + var ao_front: f32; + if (material_right.ao_id < 0) { + ao_right = material_right.ao; + } else { + ao_right = average(textureSampleGrad(material_textures, r_sampler, vec2(y,z), material_right.ao_id, ddx.yz, ddy.yz)); + } + if (material_top.ao_id < 0) { + ao_top = material_top.ao; + } else { + ao_top = average(textureSampleGrad(material_textures, r_sampler, vec2(x,z), material_top.ao_id, ddx.xz, ddy.xz)); + } + if (material_front.ao_id < 0) { + ao_front = material_front.ao; + } else { + ao_front = average(textureSampleGrad(material_textures, r_sampler, vec2(x,y), material_front.ao_id, ddx.xy, ddy.xy)); + } + out.ao = ao_right * b.x + ao_top * b.y + ao_front * b.z; + + // Normal Map + // + // Using Whiteout blend normal map in local space + var normal_right: vec3; + var normal_top: vec3; + var normal_front: vec3; + if (material_right.normal_id < 0) { + normal_right = vec3(0.,0.,1.); + } else { + normal_right = (textureSampleGrad(material_textures, r_sampler, vec2(y,z), material_right.normal_id, ddx.yz, ddy.yz).xyz - 0.5) * 2.; + } + + if (material_top.normal_id < 0) { + normal_right = vec3(0.,0.,1.); + } else { + normal_top = (textureSampleGrad(material_textures, r_sampler, vec2(x,z), material_top.normal_id, ddx.xz, ddy.xz).xyz - 0.5) * 2.; + } + if (material_front.normal_id < 0) { + normal_right = vec3(0.,0.,1.); + } else { + normal_front = (textureSampleGrad(material_textures, r_sampler, vec2(x,y), material_front.normal_id, ddx.xy, ddy.xy).xyz - 0.5) * 2.; + } + + // Apply whiteout blend + normal_right = vec3(normal_right.x+local_normal.z, normal_right.y+local_normal.y, abs(normal_right.z) * local_normal.x); + normal_top = vec3(normal_top.x+local_normal.x, normal_top.y+local_normal.z, abs(normal_top.z) * local_normal.y); + normal_front = vec3(normal_front.x+local_normal.x, normal_front.y+local_normal.y, abs(normal_front.z) * local_normal.z); + + // Swizzel into one normal + let new_local_normal = normal_right.zyx * b.x + normal_top.xzy * b.y + normal_front.xyz * b.z; + // let new_local_normal = normal_right.xyz * b.x + normal_top.xyz * b.y + normal_front.xyz * b.z; + + // Convert back to world space + out.normal = normalize((u_sdf.normal_transform * vec4(new_local_normal, 0.)).xyz); + + return out; +} diff --git a/dotrix_voxel/src/sdf/jump_flood.rs b/dotrix_voxel/src/sdf/jump_flood.rs new file mode 100644 index 00000000..2603f578 --- /dev/null +++ b/dotrix_voxel/src/sdf/jump_flood.rs @@ -0,0 +1,411 @@ +/// Use the jump flood algorithm to convert +/// a voxel into a distance field +/// +/// Paper: +/// Jump Flooding in GPU with Applications to Voronoi Diagram and Distance Transform +/// Guodong Rong et al. +/// +/// In 2006 ACM Symposium on Interactive 3D +/// Graphics and Games, 14-17 March, Redwood City, +/// CA, USA, pp. 109-116, pp. 228. +/// +use crate::{Grid, TexSdf}; +use dotrix_core::{ + assets::Shader, + ecs::{Const, Mut}, + renderer::{ + wgpu, Access, BindGroup, Binding, Buffer, Compute, ComputeArgs, ComputeOptions, + PipelineLayout, Stage, Texture as TextureBuffer, WorkGroups, + }, + Assets, Renderer, World, +}; + +const VOXEL_TO_JUMP_FLOOD_PIPELINE: &str = "dotrix_voxel::sdf::jump_flood_voxel_seed"; +const JUMP_FLOOD_PIPELINE: &str = "dotrix_voxel::sdf::jump_flood"; +const JUMP_FLOOD_TO_DF_PIPELINE: &str = "dotrix_voxel::sdf::jump_flood_sdf"; +const VOXELS_PER_WORKGROUP: [usize; 3] = [8, 8, 4]; +const SCALE_FACTOR: u32 = 4; + +/// Jump flood varients +/// detailed in section 3.3.1 of the original paper +pub enum JumpFlood { + /// Standard (fastest, most errors) + Jfa, + /// 1 Additional round + Jfa1, + /// 2 Additional rounds + Jfa2, + /// log2(n) additional rounds (slowest, least errors) + JfaSquare, +} + +/// Component for generating a SDF +/// which tells the renderer how far +/// a point is from the surface. +/// Computed with the jump flooding +/// algorithm, which is an approximate +/// algorithm with O(log2(n)) complexity +pub struct VoxelJumpFlood { + jump_flood_varient: JumpFlood, + ping_buffer: TextureBuffer, + pong_buffer: TextureBuffer, + init_pipeline: Option, + jumpflood_pipelines: Vec, + jumpflood_data: Vec, + sdf_pipeline: Option, + last_revision: Option, +} + +impl Default for VoxelJumpFlood { + fn default() -> Self { + Self::new() + } +} + +impl VoxelJumpFlood { + /// Use default parameters for a new jump flood + pub fn new() -> Self { + Self { + jump_flood_varient: JumpFlood::JfaSquare, + ping_buffer: { + let mut buffer = TextureBuffer::new_3d("PingBuffer") + .use_as_storage() + .allow_write(); + buffer.format = wgpu::TextureFormat::Rgba32Float; + buffer + }, + pong_buffer: { + let mut buffer = TextureBuffer::new_3d("PongBuffer") + .use_as_storage() + .allow_write(); + buffer.format = wgpu::TextureFormat::Rgba32Float; + buffer + }, + init_pipeline: None, + jumpflood_pipelines: vec![], + jumpflood_data: vec![], + sdf_pipeline: None, + last_revision: None, + } + } + + #[must_use] + pub fn with_algorithm(mut self, jfa: JumpFlood) -> Self { + self.jump_flood_varient = jfa; + self + } + + /// Reset the algorithm for recomputation, call this is the size of the grid has changed + /// This does not need to be called if only the values of the grid have changes + pub fn reset(&mut self) { + self.init_pipeline = None; + self.jumpflood_pipelines = vec![]; + self.jumpflood_data = vec![]; + self.sdf_pipeline = None; + self.last_revision = None; + self.ping_buffer = { + let mut buffer = TextureBuffer::new_3d("PingBuffer") + .use_as_storage() + .allow_write(); + buffer.format = wgpu::TextureFormat::Rgba32Float; + buffer + }; + self.pong_buffer = { + let mut buffer = TextureBuffer::new_3d("PongBuffer") + .use_as_storage() + .allow_write(); + buffer.format = wgpu::TextureFormat::Rgba32Float; + buffer + }; + } + + /// Load the voxel data for computation + pub fn load(&mut self, renderer: &Renderer, grid: &Grid) { + let pixel_size = 4 * 4; + let dimensions = grid.get_dimensions(); + let dim: [u32; 3] = [ + dimensions[0] * SCALE_FACTOR, + dimensions[1] * SCALE_FACTOR, + dimensions[2] * SCALE_FACTOR, + ]; + let data: Vec> = + vec![0u8; pixel_size * dim[0] as usize * dim[1] as usize * dim[2] as usize] + .chunks(dim[0] as usize * dim[1] as usize * pixel_size) + .map(|chunk| chunk.to_vec()) + .collect(); + + let slices: Vec<&[u8]> = data.iter().map(|chunk| chunk.as_slice()).collect(); + + renderer.update_or_load_texture(&mut self.ping_buffer, dim[0], dim[1], slices.as_slice()); + + renderer.update_or_load_texture(&mut self.pong_buffer, dim[0], dim[1], slices.as_slice()); + } +} + +/// Uniform structure for ancillary data of the jump flood calculation +#[repr(C)] +#[derive(Default, Clone, Copy, Debug)] +struct Data { + k: u32, + padding: [f32; 3], +} +unsafe impl bytemuck::Zeroable for Data {} +unsafe impl bytemuck::Pod for Data {} + +pub(super) fn startup(renderer: Const, mut assets: Mut) { + let mut shader = Shader { + name: String::from(JUMP_FLOOD_PIPELINE), + code: String::from(include_str!("./jump_flood/jump_flood.wgsl")), + ..Default::default() + }; + shader.load(&renderer); + + assets.store_as(shader, JUMP_FLOOD_PIPELINE); + + let mut shader = Shader { + name: String::from(VOXEL_TO_JUMP_FLOOD_PIPELINE), + code: String::from(include_str!("./jump_flood/jump_flood_voxel_seed.wgsl")), + ..Default::default() + }; + shader.load(&renderer); + + assets.store_as(shader, VOXEL_TO_JUMP_FLOOD_PIPELINE); + + let mut shader = Shader { + name: String::from(JUMP_FLOOD_TO_DF_PIPELINE), + code: String::from(include_str!("./jump_flood/jump_flood_sdf.wgsl")), + ..Default::default() + }; + shader.load(&renderer); + + assets.store_as(shader, JUMP_FLOOD_TO_DF_PIPELINE); +} + +// Compute the SDF from the grid +pub(super) fn compute(world: Const, assets: Const, mut renderer: Mut) { + for (grid, jump_flood, sdf) in world.query::<(&mut Grid, &mut VoxelJumpFlood, &mut TexSdf)>() { + let grid_dim = grid.get_dimensions(); + let dimensions: [u32; 3] = [ + grid_dim[0] * SCALE_FACTOR, + grid_dim[1] * SCALE_FACTOR, + grid_dim[2] * SCALE_FACTOR, + ]; + let workgroup_size_x = + (dimensions[0] as f32 / VOXELS_PER_WORKGROUP[0] as f32).ceil() as u32; + let workgroup_size_y = + (dimensions[1] as f32 / VOXELS_PER_WORKGROUP[1] as f32).ceil() as u32; + let workgroup_size_z = + (dimensions[2] as f32 / VOXELS_PER_WORKGROUP[2] as f32).ceil() as u32; + + // Set up pipelines once + if jump_flood.init_pipeline.is_none() { + grid.load(&renderer, &assets); + jump_flood.load(&renderer, grid); + + let mut voxel_to_jump_flood: Compute = Default::default(); + + if voxel_to_jump_flood.pipeline.shader.is_null() { + voxel_to_jump_flood.pipeline.shader = assets + .find::(VOXEL_TO_JUMP_FLOOD_PIPELINE) + .unwrap_or_default(); + } + + if let Some(shader) = assets.get(voxel_to_jump_flood.pipeline.shader) { + if !shader.loaded() { + continue; + } + + renderer.bind( + &mut voxel_to_jump_flood.pipeline, + PipelineLayout::Compute { + label: "Voxel_2_JumpFlood".into(), + shader, + bindings: &[BindGroup::new( + "Globals", + vec![ + Binding::Texture3D( + "VoxelTexture", + Stage::Compute, + grid.get_buffer(), + ), + Binding::StorageTexture3D( + "InitSeeds", + Stage::Compute, + &jump_flood.ping_buffer, + Access::WriteOnly, + ), + ], + )], + options: ComputeOptions { cs_main: "main" }, + }, + ); + + jump_flood.init_pipeline = Some(voxel_to_jump_flood); + } + } + + if jump_flood.init_pipeline.is_some() && jump_flood.jumpflood_pipelines.is_empty() { + let n = *dimensions.iter().max().unwrap(); + + let mut ping_buffer = &jump_flood.ping_buffer; + let mut pong_buffer = &jump_flood.pong_buffer; + + let n_log2 = (n as f32).log2().ceil() as u32; + let n_ceil = 2u32.pow(n_log2); + + let limit = match jump_flood.jump_flood_varient { + JumpFlood::Jfa => n_log2 as usize, + JumpFlood::Jfa1 => n_log2 as usize + 1, + JumpFlood::Jfa2 => n_log2 as usize + 2, + JumpFlood::JfaSquare => n_log2 as usize * 2, + }; + + for i in 0..limit { + let k = if i < n_log2 as usize { + n_ceil / 2u32.pow(i as u32 + 1) + } else { + 2u32.pow((i as u32) - n_log2) + }; + + let mut buffer = Buffer::uniform("Jump Flood Params"); + let data = Data { + k, + padding: Default::default(), + }; + renderer.load_buffer(&mut buffer, bytemuck::cast_slice(&[data])); + + let mut jump_flood_compute: Compute = Default::default(); + + if jump_flood_compute.pipeline.shader.is_null() { + jump_flood_compute.pipeline.shader = assets + .find::(JUMP_FLOOD_PIPELINE) + .unwrap_or_default(); + } + + if let Some(shader) = assets.get(jump_flood_compute.pipeline.shader) { + if !shader.loaded() { + continue; + } + + renderer.bind( + &mut jump_flood_compute.pipeline, + PipelineLayout::Compute { + label: "JumpFlood".into(), + shader, + bindings: &[BindGroup::new( + "Globals", + vec![ + Binding::Uniform("Params", Stage::Compute, &buffer), + Binding::Texture3D("VoxelTexture", Stage::Compute, ping_buffer), + Binding::StorageTexture3D( + "InitSeeds", + Stage::Compute, + pong_buffer, + Access::WriteOnly, + ), + ], + )], + options: ComputeOptions { cs_main: "main" }, + }, + ); + + jump_flood.jumpflood_pipelines.push(jump_flood_compute); + jump_flood.jumpflood_data.push(buffer); + (ping_buffer, pong_buffer) = (pong_buffer, ping_buffer); + } + } + + // SDF conversion + if jump_flood.sdf_pipeline.is_none() { + sdf.load(&renderer, &dimensions); + + let mut jump_flood_sdf: Compute = Default::default(); + + if jump_flood_sdf.pipeline.shader.is_null() { + jump_flood_sdf.pipeline.shader = assets + .find::(JUMP_FLOOD_TO_DF_PIPELINE) + .unwrap_or_default(); + } + + if let Some(shader) = assets.get(jump_flood_sdf.pipeline.shader) { + if !shader.loaded() { + continue; + } + + renderer.bind( + &mut jump_flood_sdf.pipeline, + PipelineLayout::Compute { + label: "JumpFlood_2_SDF".into(), + shader, + bindings: &[BindGroup::new( + "Globals", + vec![ + Binding::Texture3D("Voxel", Stage::Compute, grid.get_buffer()), + Binding::Texture3D("JumpFlood", Stage::Compute, pong_buffer), + Binding::StorageTexture3D( + "SDF", + Stage::Compute, + &sdf.buffer, + Access::WriteOnly, + ), + ], + )], + options: ComputeOptions { cs_main: "main" }, + }, + ); + + jump_flood.sdf_pipeline = Some(jump_flood_sdf); + } + } + } + + // Call compute when ever the grid changes + if jump_flood.last_revision.is_none() + || jump_flood.last_revision.unwrap() != grid.get_revision() + { + grid.load(&renderer, &assets); + jump_flood.load(&renderer, grid); + if let (Some(voxel_to_jump_flood), Some(jump_flood_sdf)) = ( + jump_flood.init_pipeline.as_mut(), + jump_flood.sdf_pipeline.as_mut(), + ) { + jump_flood.last_revision = Some(grid.get_revision()); + + renderer.compute( + &mut voxel_to_jump_flood.pipeline, + &ComputeArgs { + work_groups: WorkGroups { + x: workgroup_size_x, + y: workgroup_size_y, + z: workgroup_size_z, + }, + }, + ); + + for jump_flood_compute in jump_flood.jumpflood_pipelines.iter_mut() { + renderer.compute( + &mut jump_flood_compute.pipeline, + &ComputeArgs { + work_groups: WorkGroups { + x: workgroup_size_x, + y: workgroup_size_y, + z: workgroup_size_z, + }, + }, + ); + } + + renderer.compute( + &mut jump_flood_sdf.pipeline, + &ComputeArgs { + work_groups: WorkGroups { + x: workgroup_size_x, + y: workgroup_size_y, + z: workgroup_size_z, + }, + }, + ); + } + } + } +} diff --git a/dotrix_voxel/src/sdf/jump_flood/jump_flood.wgsl b/dotrix_voxel/src/sdf/jump_flood/jump_flood.wgsl new file mode 100644 index 00000000..045577a7 --- /dev/null +++ b/dotrix_voxel/src/sdf/jump_flood/jump_flood.wgsl @@ -0,0 +1,154 @@ +// This compute applies the jump flood algorithm +// +// The algorithm is a fast (approximate) method +// for voronoi diagrams and distance transforms +// +// It is O(log2(n)) +// +// This algorihm should be called as a ping pong buffer +// Each call should decrease k until k==1 +// while swapping the input/output texture buffers +// +// Texture buffers are of kind: +// r,g,b,a where r,g,b are the xyz values of the nearest seed +// and a is used as a flag for invalid seed when a<0 +// +// Paper: +// Jump Flooding in GPU with Applications to Voronoi Diagram and Distance Transform +// Guodong Rong et al. +// +// In 2006 ACM Symposium on Interactive 3D +// Graphics and Games, 14-17 March, Redwood City, +// CA, USA, pp. 109-116, pp. 228. + +struct Data { + // The current iterations step size must be >=1 + k: i32; +}; + +[[group(0), binding(0)]] +var data: Data; + +// The previous run seed values stored in each pixel +[[group(0), binding(1)]] +var init_seeds: texture_3d; + +// The next run's seed values +[[group(0), binding(2)]] +var out_seeds: texture_storage_3d; + +fn value_at(coord: vec3) -> vec3 { + return textureLoad(init_seeds, coord, 0).rgb; +} + +// Write location of current nearest seed for this pixel +// Written into the RGB channels +fn set_value_at(coord: vec3, value: vec3) { + textureStore(out_seeds, coord, vec4(value, 0.)); +} + +/// Checks if it is has an invalid seed location +fn is_invalid_at(coord: vec3) -> bool { + return textureLoad(init_seeds, coord, 0).a < 0.; +} + +fn is_out_of_bounds(coord: vec3) -> bool { + let tex_dim: vec3 = textureDimensions(init_seeds); + return ( + coord[0] < 0 + || coord[1] < 0 + || coord[2] < 0 + || coord[0] >= tex_dim[0] + || coord[1] >= tex_dim[1] + || coord[2] >= tex_dim[2] + ); +} + +// For a given voxel get its origin in local seed space +fn origin(coord: vec3) -> vec3 { + return vec3(f32(coord[0]),f32(coord[1]),f32(coord[2])); +} + +// For a given pixel tries to read the seed value, +// then compares to a reference seed distance +// and identifies if it is a better seed distance +fn is_seed_better(origin_coord: vec3, delta: vec3, best_seed: ptr>) { + let coord: vec3 = origin_coord + delta; + if (is_out_of_bounds(coord)) { + return; + } + if (is_invalid_at(coord)) { + return; + } + let new_seed: vec3 = value_at(coord); + let origin_pos: vec3 = origin(origin_coord); + if (is_invalid_at(origin_coord) || distance(new_seed, origin_pos) < distance(*best_seed, origin_pos)) { + *best_seed = new_seed; + } +} + + +[[stage(compute), workgroup_size(8,8,4)]] +fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3) { + // Jump Flood Algorithm: + // + // n = number of pixels in largest dimension + // Loop over ceil(log2(n)) times over the image i=[1, ceil(log2(n))] + // For n = 8, k = n/2, n/4, n/8 + // For n = 16, k= n/2, n/4, n/8, n/16 + // For n = 17, k= n/2, n/4, n/8, n/16, n/32 + // k = n/(2^(i)) + // + // Look in all seeds at location of origin±k + // If seed found in neighbouring cell is better than current + // then use that one + // + // This compute does only a single value of k + // it must be enqueued multiple times to complete the jump flood + // with a ping pong style buffer + // + var k: i32 = data.k; + if (k<1) { + return; + } + + let origin_coord: vec3 = vec3( + i32(global_invocation_id[0]), + i32(global_invocation_id[1]), + i32(global_invocation_id[2]), + ); + + var best_seed: vec3 = value_at(origin_coord); + + is_seed_better(origin_coord, vec3(-k,-k,-k), &best_seed); + is_seed_better(origin_coord, vec3( 0,-k,-k), &best_seed); + is_seed_better(origin_coord, vec3( k,-k,-k), &best_seed); + is_seed_better(origin_coord, vec3(-k, 0,-k), &best_seed); + is_seed_better(origin_coord, vec3( 0, 0,-k), &best_seed); + is_seed_better(origin_coord, vec3( k, 0,-k), &best_seed); + is_seed_better(origin_coord, vec3(-k, k,-k), &best_seed); + is_seed_better(origin_coord, vec3( 0, k,-k), &best_seed); + is_seed_better(origin_coord, vec3( k, k,-k), &best_seed); + + is_seed_better(origin_coord, vec3(-k,-k, 0), &best_seed); + is_seed_better(origin_coord, vec3( 0,-k, 0), &best_seed); + is_seed_better(origin_coord, vec3( k,-k, 0), &best_seed); + is_seed_better(origin_coord, vec3(-k, 0, 0), &best_seed); + // is_seed_better(origin_coord, vec3( 0, 0, 0), &best_seed); + is_seed_better(origin_coord, vec3( k, 0, 0), &best_seed); + is_seed_better(origin_coord, vec3(-k, k, 0), &best_seed); + is_seed_better(origin_coord, vec3( 0, k, 0), &best_seed); + is_seed_better(origin_coord, vec3( k, k, 0), &best_seed); + + is_seed_better(origin_coord, vec3(-k,-k, k), &best_seed); + is_seed_better(origin_coord, vec3( 0,-k, k), &best_seed); + is_seed_better(origin_coord, vec3( k,-k, k), &best_seed); + is_seed_better(origin_coord, vec3(-k, 0, k), &best_seed); + is_seed_better(origin_coord, vec3( 0, 0, k), &best_seed); + is_seed_better(origin_coord, vec3( k, 0, k), &best_seed); + is_seed_better(origin_coord, vec3(-k, k, k), &best_seed); + is_seed_better(origin_coord, vec3( 0, k, k), &best_seed); + is_seed_better(origin_coord, vec3( k, k, k), &best_seed); + + set_value_at(origin_coord, best_seed); +} diff --git a/dotrix_voxel/src/sdf/jump_flood/jump_flood_sdf.wgsl b/dotrix_voxel/src/sdf/jump_flood/jump_flood_sdf.wgsl new file mode 100644 index 00000000..5bf3786a --- /dev/null +++ b/dotrix_voxel/src/sdf/jump_flood/jump_flood_sdf.wgsl @@ -0,0 +1,158 @@ +// Takes a jump flood result and computes the SDF +// +// Jump flood has seeds in jump flood space +// SDF should have distance in voxel space +// Conversion is done where appropiate +// + +let ISO_SURFACE: f32 = 0.5; + +// The density for the voxel should be stored in the r channel +// The material should be in the g channel +[[group(0), binding(0)]] +var voxels: texture_3d; + +// The rgb channels will be set to contain the nearest seed location +[[group(0), binding(1)]] +var jump_flood: texture_3d; + + +// The r channel will contain the DF +// The g channel will copy the material ID from the voxel +[[group(0), binding(2)]] +var sdf: texture_storage_3d; + + +// For a given jump_flood coord get its coord in local voxel space +fn seed_coord_to_voxel_pos(coord: vec3) -> vec3 { + let seed_dim: vec3 = textureDimensions(jump_flood) - vec3(1); + let voxel_dim: vec3 = textureDimensions(voxels) - vec3(1); + let seed_dim_f32: vec3 = vec3(f32(seed_dim.x), f32(seed_dim.y), f32(seed_dim.z)); + let voxel_dim_f32: vec3 = vec3(f32(voxel_dim.x), f32(voxel_dim.y), f32(voxel_dim.z)); + return vec3(f32(coord[0]),f32(coord[1]),f32(coord[2])) * voxel_dim_f32 / seed_dim_f32; +} + +// Get the seed position from the jump flood in local voxel space +fn seed_position(coord: vec3) -> vec3 { + let seed_dim: vec3 = textureDimensions(jump_flood) - vec3(1); + let voxel_dim: vec3 = textureDimensions(voxels) - vec3(1); + let seed_dim_f32: vec3 = vec3(f32(seed_dim.x), f32(seed_dim.y), f32(seed_dim.z)); + let voxel_dim_f32: vec3 = vec3(f32(voxel_dim.x), f32(voxel_dim.y), f32(voxel_dim.z)); + return textureLoad(jump_flood, coord, 0).rgb * voxel_dim_f32 / seed_dim_f32; +} + +// Get density from voxel using voxel coord +// - Used to work out if a point is inside or not +fn voxel_value(coord: vec3) -> f32 { + let voxel_dim: vec3 = textureDimensions(voxels) - vec3(1); + let i: i32 = clamp(i32(coord.x), 0, voxel_dim.x); + let j: i32 = clamp(i32(coord.y), 0, voxel_dim.y); + let k: i32 = clamp(i32(coord.z), 0, voxel_dim.z); + let x: f32 = clamp(coord.x - f32(i), 0., 1.); + let y: f32 = clamp(coord.y - f32(j), 0., 1.); + let z: f32 = clamp(coord.z - f32(k), 0., 1.); + + let f000: f32 = f32(textureLoad(voxels, + vec3( + i, + j, + k, + ) + ,0).r); + let f001: f32 = f32(textureLoad(voxels, + vec3( + i, + j, + k + 1, + ) + ,0).r); + let f010: f32 = f32(textureLoad(voxels, + vec3( + i, + j + 1, + k, + ) + ,0).r); + let f011: f32 = f32(textureLoad(voxels, + vec3( + i, + j + 1, + k + 1, + ) + ,0).r); + let f100: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j, + k, + ) + ,0).r); + let f101: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j, + k + 1, + ) + ,0).r); + let f110: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j + 1, + k, + ) + ,0).r); + let f111: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j + 1, + k + 1, + ) + ,0).r); + + return + ( + f000*(1.-x)*(1.-y)*(1.-z) + +f001*(1.-x)*(1.-y)*z + +f010*(1.-x)*y *(1.-z) + +f011*(1.-x)*y *z + +f100*x *(1.-y)*(1.-z) + +f101*x *(1.-y)*z + +f110*x *y *(1.-z) + +f111*x *y *z + ) - ISO_SURFACE; +} + +// Get the material from the voxel +// using nearest neighbour +fn material(coord: vec3) -> u32 { + let nearest_pos: vec3 = round(coord); + let nearest_coord: vec3 = vec3(i32(nearest_pos.x), i32(nearest_pos.y), i32(nearest_pos.z)); + return textureLoad(voxels, nearest_coord, 0).g; +} + +// Save the distance field value into the SDF texture +fn save_sdf(coord: vec3, dist: f32, material_id: f32) { + textureStore(sdf, coord, vec4(dist, material_id, 0., 0.)); +} + +// Check if a position in inside or outside the surface +fn is_outside(pos_voxel: vec3) -> bool { + // return true; + return voxel_value(pos_voxel) > 1e-4; +} + +[[stage(compute), workgroup_size(8,8,4)]] +fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3) { + let coord_seed: vec3 = vec3( + i32(global_invocation_id[0]), + i32(global_invocation_id[1]), + i32(global_invocation_id[2]), + ); + + let pos_voxel: vec3 = seed_coord_to_voxel_pos(coord_seed); + let seed_pos_voxel: vec3 = seed_position(coord_seed); + let dist: f32 = select(1.,-1., is_outside(pos_voxel)) * distance(pos_voxel, seed_pos_voxel); + let material_id: f32 = f32(material(pos_voxel)); + + save_sdf(coord_seed, dist, material_id); +} diff --git a/dotrix_voxel/src/sdf/jump_flood/jump_flood_voxel_seed.wgsl b/dotrix_voxel/src/sdf/jump_flood/jump_flood_voxel_seed.wgsl new file mode 100644 index 00000000..c128904a --- /dev/null +++ b/dotrix_voxel/src/sdf/jump_flood/jump_flood_voxel_seed.wgsl @@ -0,0 +1,194 @@ +// This compute takes a voxel +// and write the init seed location for the +// jump flood algorithm +// +// + +let ISO_SURFACE: f32 = 0.5; + +// The density for the voxel should be stored in the r channel +[[group(0), binding(0)]] +var voxels: texture_3d; +// The rgb channels will be set to contain the nearest seed location +// of voxels that cross the root +[[group(0), binding(1)]] +var init_seeds: texture_storage_3d; + +// Write location of current nearest seed for this pixel +// Written into the RGB channels +fn set_seed_at(value: vec3, coord: vec3) { + textureStore(init_seeds, coord, vec4(value, 0.)); +} + +/// Marks a cell as being invalid with no known data yet +fn set_seed_invalid_at(coord: vec3) { + textureStore(init_seeds, coord, vec4(0.,0.,0., -1.)); +} + +// Get density from voxel +fn voxel_value(seed_coord: vec3) -> f32 { + let voxel_dim: vec3 = textureDimensions(voxels) - vec3(1); + let voxel_dim_f32: vec3 = vec3(f32(voxel_dim.x), f32(voxel_dim.y), f32(voxel_dim.z)); + let seed_dim: vec3 = textureDimensions(init_seeds) - vec3(1); + let seed_dim_f32: vec3 = vec3(f32(seed_dim.x), f32(seed_dim.y), f32(seed_dim.z)); + + let coord: vec3 = seed_coord / seed_dim_f32 * voxel_dim_f32; + let i: i32 = clamp(i32(coord.x), 0, voxel_dim.x); + let j: i32 = clamp(i32(coord.y), 0, voxel_dim.y); + let k: i32 = clamp(i32(coord.z), 0, voxel_dim.z); + let x: f32 = clamp(coord.x - f32(i), 0., 1.); + let y: f32 = clamp(coord.y - f32(j), 0., 1.); + let z: f32 = clamp(coord.z - f32(k), 0., 1.); + + let f000: f32 = f32(textureLoad(voxels, + vec3( + i, + j, + k, + ) + ,0).r); + let f001: f32 = f32(textureLoad(voxels, + vec3( + i, + j, + k + 1, + ) + ,0).r); + let f010: f32 = f32(textureLoad(voxels, + vec3( + i, + j + 1, + k, + ) + ,0).r); + let f011: f32 = f32(textureLoad(voxels, + vec3( + i, + j + 1, + k + 1, + ) + ,0).r); + let f100: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j, + k, + ) + ,0).r); + let f101: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j, + k + 1, + ) + ,0).r); + let f110: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j + 1, + k, + ) + ,0).r); + let f111: f32 = f32(textureLoad(voxels, + vec3( + i + 1, + j + 1, + k + 1, + ) + ,0).r); + + return + ( + f000*(1.-x)*(1.-y)*(1.-z) + +f001*(1.-x)*(1.-y)*z + +f010*(1.-x)*y *(1.-z) + +f011*(1.-x)*y *z + +f100*x *(1.-y)*(1.-z) + +f101*x *(1.-y)*z + +f110*x *y *(1.-z) + +f111*x *y *z + ) - ISO_SURFACE; +} + +// For a given pixel get its pos in local seed space +fn seed_coord(seed_pixel: vec3) -> vec3 { + return vec3(f32(seed_pixel[0]),f32(seed_pixel[1]),f32(seed_pixel[2])); +} + +fn approximate_root(seed_pixel_a: vec3, delta: vec3, current_best: ptr) { + let seed_pixel_b: vec3 = seed_pixel_a + delta; + let seed_pos_a: vec3 = seed_coord(seed_pixel_a); + let seed_pos_b: vec3 = seed_coord(seed_pixel_b); + let voxel_value_a = voxel_value(seed_pos_a); + let voxel_value_b = voxel_value(seed_pos_b); + let weight: f32 = voxel_value_a/(voxel_value_a-voxel_value_b); + + if (weight<0. || weight > 1.) { + return; + } + + let root_pos: vec3 = mix(seed_pos_a, seed_pos_b, weight); + let dist: f32 = distance(seed_pos_a, root_pos); + if (dist < *current_best) { + set_seed_at(root_pos, seed_pixel_a); + *current_best = dist; + } + return; +} + +[[stage(compute), workgroup_size(8,8,4)]] +fn main([[builtin(global_invocation_id)]] global_invocation_id: vec3) { + // Seeding the jump flood + // x---x---x + // | \ | / | + // x---o---x + // | / | \ | + // x---x---x + // + // - From o travel to x in all directions + // - If travelling from o-x crosses the ISO_SURFACE + // then use linear interpolation to find the + // ISO_SURFACE point + // - Calculate the distance o to the ISO_SURFACE + // - If distance is better then current_best distance + // set the seed as ISO_SURFACE and update current_best + // distance + let seed_pixel: vec3 = vec3( + i32(global_invocation_id[0]), + i32(global_invocation_id[1]), + i32(global_invocation_id[2]), + ); + + var best_value: f32 = 4000.; + set_seed_invalid_at(seed_pixel); + approximate_root(seed_pixel, vec3(-1,-1,-1), &best_value); + approximate_root(seed_pixel, vec3( 0,-1,-1), &best_value); + approximate_root(seed_pixel, vec3( 1,-1,-1), &best_value); + approximate_root(seed_pixel, vec3(-1, 0,-1), &best_value); + approximate_root(seed_pixel, vec3( 0, 0,-1), &best_value); + approximate_root(seed_pixel, vec3( 1, 0,-1), &best_value); + approximate_root(seed_pixel, vec3(-1, 1,-1), &best_value); + approximate_root(seed_pixel, vec3( 0, 1,-1), &best_value); + approximate_root(seed_pixel, vec3( 1, 1,-1), &best_value); + + approximate_root(seed_pixel, vec3(-1,-1, 0), &best_value); + approximate_root(seed_pixel, vec3( 0,-1, 0), &best_value); + approximate_root(seed_pixel, vec3( 1,-1, 0), &best_value); + approximate_root(seed_pixel, vec3(-1, 0, 0), &best_value); + // approximate_root(seed_pixel, vec3( 0, 0, 0), &best_value); + approximate_root(seed_pixel, vec3( 1, 0, 0), &best_value); + approximate_root(seed_pixel, vec3(-1, 1, 0), &best_value); + approximate_root(seed_pixel, vec3( 0, 1, 0), &best_value); + approximate_root(seed_pixel, vec3( 1, 1, 0), &best_value); + + approximate_root(seed_pixel, vec3(-1,-1, 1), &best_value); + approximate_root(seed_pixel, vec3( 0,-1, 1), &best_value); + approximate_root(seed_pixel, vec3( 1,-1, 1), &best_value); + approximate_root(seed_pixel, vec3(-1, 0, 1), &best_value); + approximate_root(seed_pixel, vec3( 0, 0, 1), &best_value); + approximate_root(seed_pixel, vec3( 1, 0, 1), &best_value); + approximate_root(seed_pixel, vec3(-1, 1, 1), &best_value); + approximate_root(seed_pixel, vec3( 0, 1, 1), &best_value); + approximate_root(seed_pixel, vec3( 1, 1, 1), &best_value); + +} diff --git a/dotrix_voxel/src/sdf/mod.rs b/dotrix_voxel/src/sdf/mod.rs new file mode 100644 index 00000000..f3762bb1 --- /dev/null +++ b/dotrix_voxel/src/sdf/mod.rs @@ -0,0 +1,18 @@ +use dotrix_core::ecs::System; +use dotrix_core::Application; + +mod circle_trace; +mod jump_flood; +mod tex_sdf; + +pub use circle_trace::*; +pub use jump_flood::*; +pub use tex_sdf::*; + +/// Enables Voxel SDF Dotrix Extension +pub fn extension(app: &mut Application) { + app.add_system(System::from(jump_flood::startup)); + app.add_system(System::from(jump_flood::compute)); + + circle_trace::extension(app); +} diff --git a/dotrix_voxel/src/sdf/tex_sdf.rs b/dotrix_voxel/src/sdf/tex_sdf.rs new file mode 100644 index 00000000..b0527b7e --- /dev/null +++ b/dotrix_voxel/src/sdf/tex_sdf.rs @@ -0,0 +1,53 @@ +use dotrix_core::renderer::{Buffer, Pipeline, Renderer, Texture as TextureBuffer}; + +/// Object to hold the 3D texture containing an Sdf +pub struct TexSdf { + /// Texture buffer containing a 3d texture + /// with r channel of the distance anf g channel of the material ID + pub buffer: TextureBuffer, + /// Pipeline for renderering this SDF + pub pipeline: Pipeline, + /// Uniform that holds render related data + pub data: Buffer, +} + +impl Default for TexSdf { + fn default() -> Self { + Self { + buffer: { + let mut buffer = TextureBuffer::new_3d("TexSDF") + .use_as_storage() + .allow_write(); + buffer.format = wgpu::TextureFormat::Rg32Float; + buffer + }, + pipeline: Default::default(), + data: Buffer::uniform("TexSdf Data"), + } + } +} + +impl TexSdf { + pub fn load(&mut self, renderer: &Renderer, dimensions: &[u32; 3]) { + let pixel_size = 4 * 2; + let data: Vec> = vec![ + 0u8; + pixel_size + * dimensions[0] as usize + * dimensions[1] as usize + * dimensions[2] as usize + ] + .chunks(dimensions[0] as usize * dimensions[1] as usize * pixel_size) + .map(|chunk| chunk.to_vec()) + .collect(); + + let slices: Vec<&[u8]> = data.iter().map(|chunk| chunk.as_slice()).collect(); + + renderer.update_or_load_texture( + &mut self.buffer, + dimensions[0], + dimensions[1], + slices.as_slice(), + ); + } +} diff --git a/dotrix_voxel/src/voxel.rs b/dotrix_voxel/src/voxel.rs new file mode 100644 index 00000000..4974c6af --- /dev/null +++ b/dotrix_voxel/src/voxel.rs @@ -0,0 +1,8 @@ +/// An individual voxel +#[derive(Default, Clone, Copy, Debug)] +pub struct Voxel { + /// Voxel density value + pub value: u8, + /// Voxel material id + pub material: u8, +} diff --git a/examples/voxel_sdf/main.rs b/examples/voxel_sdf/main.rs new file mode 100644 index 00000000..46b54600 --- /dev/null +++ b/examples/voxel_sdf/main.rs @@ -0,0 +1,208 @@ +use dotrix::egui::{DragValue, Egui, TopBottomPanel}; +use dotrix::overlay::Overlay; +use dotrix::{ + assets::Texture, + camera, + ecs::{Const, Mut}, + egui, overlay, Assets, Camera, Dotrix, System, Transform, World, +}; +use dotrix_pbr::Material; +use dotrix_voxel::{Grid, Light, MaterialSet, TexSdf, VoxelJumpFlood}; +use rand::Rng; + +fn main() { + Dotrix::application("Dotrix: Voxel SDF") + .with(System::from(startup)) + .with(System::from(camera::control)) + .with(System::from(self::ui)) + .with(overlay::extension) + .with(egui::extension) + .with(dotrix_voxel::extension) + .run(); +} + +fn startup(mut camera: Mut, mut world: Mut, mut assets: Mut) { + camera.target = [0., 0., 0.].into(); + camera.distance = 30.0; + camera.tilt = 0.0; + + let mut grid = Grid::default(); + randomize_grid(&mut grid); + + assets.import("assets/textures/mossy_bricks/Bricks076C_1K_AmbientOcclusion.jpg"); + assets.import("assets/textures/mossy_bricks/Bricks076C_1K_Color.jpg"); + assets.import("assets/textures/mossy_bricks/Bricks076C_1K_NormalDX.jpg"); + assets.import("assets/textures/mossy_bricks/Bricks076C_1K_Roughness.jpg"); + + let ao = assets.register::("Bricks076C_1K_AmbientOcclusion"); + let albedo = assets.register::("Bricks076C_1K_Color"); + let normal = assets.register::("Bricks076C_1K_NormalDX"); + let roughness = assets.register::("Bricks076C_1K_Roughness"); + let mut material_set = MaterialSet::default(); + material_set.set_material( + 0, + Material { + texture: albedo, + albedo: [0.5, 0.5, 0.5].into(), + roughness_texture: roughness, + ao_texture: ao, + normal_texture: normal, + metallic: 0.0, + ..Default::default() + }, + ); + + assets.import("assets/textures/PaintedPlaster010/PaintedPlaster010_1K_AmbientOcclusion.png"); + assets.import("assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Color.png"); + assets.import("assets/textures/PaintedPlaster010/PaintedPlaster010_1K_NormalDX.png"); + assets.import("assets/textures/PaintedPlaster010/PaintedPlaster010_1K_Roughness.png"); + + let ao = assets.register::("PaintedPlaster010_1K_AmbientOcclusion"); + let albedo = assets.register::("PaintedPlaster010_1K_Color"); + let normal = assets.register::("PaintedPlaster010_1K_NormalDX"); + let roughness = assets.register::("PaintedPlaster010_1K_Roughness"); + material_set.set_material( + 1, + Material { + texture: albedo, + albedo: [0.5, 0.5, 0.5].into(), + roughness_texture: roughness, + ao_texture: ao, + normal_texture: normal, + metallic: 0.0, + ..Default::default() + }, + ); + + assets.import("assets/textures/Bricks075B/Bricks075B_1K_AmbientOcclusion.jpg"); + assets.import("assets/textures/Bricks075B/Bricks075B_1K_Color.jpg"); + assets.import("assets/textures/Bricks075B/Bricks075B_1K_NormalDX.jpg"); + assets.import("assets/textures/Bricks075B/Bricks075B_1K_Roughness.jpg"); + + let ao = assets.register::("Bricks075B_1K_AmbientOcclusion"); + let albedo = assets.register::("Bricks075B_1K_Color"); + let normal = assets.register::("Bricks075B_1K_NormalDX"); + let roughness = assets.register::("Bricks075B_1K_Roughness"); + material_set.set_material( + 2, + Material { + texture: albedo, + albedo: [0.5, 0.5, 0.5].into(), + roughness_texture: roughness, + ao_texture: ao, + normal_texture: normal, + metallic: 0.0, + ..Default::default() + }, + ); + + assets.import("assets/textures/PavingStones113/PavingStones113_1K_AmbientOcclusion.jpg"); + assets.import("assets/textures/PavingStones113/PavingStones113_1K_Color.jpg"); + assets.import("assets/textures/PavingStones113/PavingStones113_1K_NormalDX.jpg"); + assets.import("assets/textures/PavingStones113/PavingStones113_1K_Roughness.jpg"); + + let ao = assets.register::("PavingStones113_1K_AmbientOcclusion"); + let albedo = assets.register::("PavingStones113_1K_Color"); + let normal = assets.register::("PavingStones113_1K_NormalDX"); + let roughness = assets.register::("PavingStones113_1K_Roughness"); + material_set.set_material( + 3, + Material { + texture: albedo, + albedo: [0.5, 0.5, 0.5].into(), + roughness_texture: roughness, + ao_texture: ao, + normal_texture: normal, + metallic: 0.0, + ..Default::default() + }, + ); + + world.spawn(vec![( + grid, + material_set, + // Instruct it to use the JumpFlood algorithm to convert the Voxel to an SDF + VoxelJumpFlood::default(), + // Render as a 3D texture based SDF + TexSdf::default(), + // Transform the voxel where you like + Transform::builder() + // .with_translate([2.,2.,2.].into()) + .with_scale([1., 3., 1.].into()) + .build(), + )]); + + world.spawn(Some((Light::Ambient { + color: [0., 0., 0.].into(), + intensity: 0., + },))); + world.spawn(Some((Light::Directional { + color: [1., 1., 1.].into(), + direction: [100., -100., -100.].into(), + intensity: 1., + enabled: true, + },))); +} + +pub fn ui(overlay: Mut, world: Const) { + let egui = overlay + .get::() + .expect("Renderer does not contain an Overlay instance"); + TopBottomPanel::bottom("my_panel").show(&egui.ctx, |ui| { + for (grid, transform) in world.query::<(&mut Grid, &mut Transform)>() { + if ui.button("Randomize").clicked() { + randomize_grid(grid); + } + ui.add( + DragValue::new(&mut transform.scale[0]) + .speed(0.1) + .prefix("X:"), + ); + ui.add( + DragValue::new(&mut transform.scale[1]) + .speed(0.1) + .prefix("Y:"), + ); + ui.add( + DragValue::new(&mut transform.scale[2]) + .speed(0.1) + .prefix("Z:"), + ); + } + }); +} + +fn randomize_grid(grid: &mut Grid) { + let dims = grid.get_dimensions(); + let total_size: usize = dims.iter().fold(1usize, |acc, &item| acc * (item as usize)); + let values: Vec = vec![0u8; total_size] + .iter() + .map(|_v| { + let chance: u8 = rand::thread_rng().gen(); + if chance > 128 { + 1 + } else { + 0 + } + }) + .collect(); + + grid.set_values(values); + + let material_values: Vec = vec![0u8; total_size] + .iter() + .map(|_v| { + let chance: u8 = rand::thread_rng().gen(); + if chance > 192 { + 3 + } else if chance > 128 { + 2 + } else if chance > 64 { + 1 + } else { + 0 + } + }) + .collect(); + grid.set_materials(material_values); +} diff --git a/rustfmt.toml b/rustfmt.toml new file mode 100644 index 00000000..32a9786f --- /dev/null +++ b/rustfmt.toml @@ -0,0 +1 @@ +edition = "2018" diff --git a/src/lib.rs b/src/lib.rs index 180947b7..9f6d6953 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -50,6 +50,9 @@ pub use dotrix_sky as sky; #[cfg(feature = "terrain")] pub use dotrix_terrain as terrain; +#[cfg(feature = "voxel")] +pub use dotrix_voxel as voxel; + pub mod prelude { pub use crate::Dotrix; pub use dotrix_core::ecs::{Const, Context, Mut, System};