working version

10 files changed, 2526 insertions, 0 deletions

diff --git a/.gitignore b/.gitignore
new file mode 100644
index 0000000..ea8c4bf
--- /dev/null
+++ b/.gitignore
@@ -0,0 +1 @@
+/target
diff --git a/Cargo.lock b/Cargo.lock
new file mode 100644
index 0000000..b7d4bcf
--- /dev/null
+++ b/Cargo.lock
@@ -0,0 +1,1325 @@
+# This file is automatically @generated by Cargo.
+# It is not intended for manual editing.
+version = 4
+
+[[package]]
+name = "aho-corasick"
+version = "1.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ddd31a130427c27518df266943a5308ed92d4b226cc639f5a8f1002816174301"
+dependencies = [
+ "memchr",
+]
+
+[[package]]
+name = "allocator-api2"
+version = "0.2.21"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "683d7910e743518b0e34f1186f92494becacb047c7b6bf616c96772180fef923"
+
+[[package]]
+name = "android_system_properties"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "anstream"
+version = "0.6.21"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "43d5b281e737544384e969a5ccad3f1cdd24b48086a0fc1b2a5262a26b8f4f4a"
+dependencies = [
+ "anstyle",
+ "anstyle-parse",
+ "anstyle-query",
+ "anstyle-wincon",
+ "colorchoice",
+ "is_terminal_polyfill",
+ "utf8parse",
+]
+
+[[package]]
+name = "anstyle"
+version = "1.0.13"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5192cca8006f1fd4f7237516f40fa183bb07f8fbdfedaa0036de5ea9b0b45e78"
+
+[[package]]
+name = "anstyle-parse"
+version = "0.2.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4e7644824f0aa2c7b9384579234ef10eb7efb6a0deb83f9630a49594dd9c15c2"
+dependencies = [
+ "utf8parse",
+]
+
+[[package]]
+name = "anstyle-query"
+version = "1.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "40c48f72fd53cd289104fc64099abca73db4166ad86ea0b4341abe65af83dadc"
+dependencies = [
+ "windows-sys",
+]
+
+[[package]]
+name = "anstyle-wincon"
+version = "3.0.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "291e6a250ff86cd4a820112fb8898808a366d8f9f58ce16d1f538353ad55747d"
+dependencies = [
+ "anstyle",
+ "once_cell_polyfill",
+ "windows-sys",
+]
+
+[[package]]
+name = "arrayvec"
+version = "0.7.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7c02d123df017efcdfbd739ef81735b36c5ba83ec3c59c80a9d7ecc718f92e50"
+
+[[package]]
+name = "ash"
+version = "0.38.0+1.3.281"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f"
+dependencies = [
+ "libloading",
+]
+
+[[package]]
+name = "autocfg"
+version = "1.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c08606f8c3cbf4ce6ec8e28fb0014a2c086708fe954eaa885384a6165172e7e8"
+
+[[package]]
+name = "bit-set"
+version = "0.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "08807e080ed7f9d5433fa9b275196cfc35414f66a0c79d864dc51a0d825231a3"
+dependencies = [
+ "bit-vec",
+]
+
+[[package]]
+name = "bit-vec"
+version = "0.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5e764a1d40d510daf35e07be9eb06e75770908c27d411ee6c92109c9840eaaf7"
+
+[[package]]
+name = "bitcoin-rgpu"
+version = "0.1.0"
+dependencies = [
+ "bytemuck",
+ "env_logger",
+ "hex",
+ "pollster",
+ "wgpu",
+]
+
+[[package]]
+name = "bitflags"
+version = "2.11.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "843867be96c8daad0d758b57df9392b6d8d271134fce549de6ce169ff98a92af"
+
+[[package]]
+name = "block"
+version = "0.1.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a"
+
+[[package]]
+name = "bumpalo"
+version = "3.20.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5d20789868f4b01b2f2caec9f5c4e0213b41e3e5702a50157d699ae31ced2fcb"
+
+[[package]]
+name = "bytemuck"
+version = "1.25.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c8efb64bd706a16a1bdde310ae86b351e4d21550d98d056f22f8a7f7a2183fec"
+dependencies = [
+ "bytemuck_derive",
+]
+
+[[package]]
+name = "bytemuck_derive"
+version = "1.10.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f9abbd1bc6865053c427f7198e6af43bfdedc55ab791faed4fbd361d789575ff"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "cfg-if"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801"
+
+[[package]]
+name = "cfg_aliases"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "613afe47fcd5fac7ccf1db93babcb082c5994d996f20b8b159f2ad1658eb5724"
+
+[[package]]
+name = "codespan-reporting"
+version = "0.12.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fe6d2e5af09e8c8ad56c969f2157a3d4238cebc7c55f0a517728c38f7b200f81"
+dependencies = [
+ "serde",
+ "termcolor",
+ "unicode-width",
+]
+
+[[package]]
+name = "colorchoice"
+version = "1.0.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b05b61dc5112cbb17e4b6cd61790d9845d13888356391624cbe7e41efeac1e75"
+
+[[package]]
+name = "core-foundation"
+version = "0.10.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2a6cd9ae233e7f62ba4e9353e81a88df7fc8a5987b8d445b4d90c879bd156f6"
+dependencies = [
+ "core-foundation-sys",
+ "libc",
+]
+
+[[package]]
+name = "core-foundation-sys"
+version = "0.8.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "773648b94d0e5d620f64f280777445740e61fe701025087ec8b57f45c791888b"
+
+[[package]]
+name = "core-graphics-types"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3d44a101f213f6c4cdc1853d4b78aef6db6bdfa3468798cc1d9912f4735013eb"
+dependencies = [
+ "bitflags",
+ "core-foundation",
+ "libc",
+]
+
+[[package]]
+name = "crunchy"
+version = "0.2.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "460fbee9c2c2f33933d720630a6a0bac33ba7053db5344fac858d4b8952d77d5"
+
+[[package]]
+name = "document-features"
+version = "0.2.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d4b8a88685455ed29a21542a33abd9cb6510b6b129abadabdcef0f4c55bc8f61"
+dependencies = [
+ "litrs",
+]
+
+[[package]]
+name = "env_filter"
+version = "1.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7a1c3cc8e57274ec99de65301228b537f1e4eedc1b8e0f9411c6caac8ae7308f"
+dependencies = [
+ "log",
+ "regex",
+]
+
+[[package]]
+name = "env_logger"
+version = "0.11.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b2daee4ea451f429a58296525ddf28b45a3b64f1acf6587e2067437bb11e218d"
+dependencies = [
+ "anstream",
+ "anstyle",
+ "env_filter",
+ "jiff",
+ "log",
+]
+
+[[package]]
+name = "equivalent"
+version = "1.0.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "877a4ace8713b0bcf2a4e7eec82529c029f1d0619886d18145fea96c3ffe5c0f"
+
+[[package]]
+name = "foldhash"
+version = "0.1.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d9c4f5dac5e15c24eb999c26181a6ca40b39fe946cbe4c263c7209467bc83af2"
+
+[[package]]
+name = "foldhash"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "77ce24cb58228fbb8aa041425bb1050850ac19177686ea6e0f41a70416f56fdb"
+
+[[package]]
+name = "foreign-types"
+version = "0.5.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d737d9aa519fb7b749cbc3b962edcf310a8dd1f4b67c91c4f83975dbdd17d965"
+dependencies = [
+ "foreign-types-macros",
+ "foreign-types-shared",
+]
+
+[[package]]
+name = "foreign-types-macros"
+version = "0.2.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1a5c6c585bc94aaf2c7b51dd4c2ba22680844aba4c687be581871a6f518c5742"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "foreign-types-shared"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "aa9a19cbb55df58761df49b23516a86d432839add4af60fc256da840f66ed35b"
+
+[[package]]
+name = "futures-core"
+version = "0.3.32"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7e3450815272ef58cec6d564423f6e755e25379b217b0bc688e295ba24df6b1d"
+
+[[package]]
+name = "futures-task"
+version = "0.3.32"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "037711b3d59c33004d3856fbdc83b99d4ff37a24768fa1be9ce3538a1cde4393"
+
+[[package]]
+name = "futures-util"
+version = "0.3.32"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "389ca41296e6190b48053de0321d02a77f32f8a5d2461dd38762c0593805c6d6"
+dependencies = [
+ "futures-core",
+ "futures-task",
+ "pin-project-lite",
+ "slab",
+]
+
+[[package]]
+name = "gl_generator"
+version = "0.14.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1a95dfc23a2b4a9a2f5ab41d194f8bfda3cabec42af4e39f08c339eb2a0c124d"
+dependencies = [
+ "khronos_api",
+ "log",
+ "xml-rs",
+]
+
+[[package]]
+name = "glow"
+version = "0.16.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c5e5ea60d70410161c8bf5da3fdfeaa1c72ed2c15f8bbb9d19fe3a4fad085f08"
+dependencies = [
+ "js-sys",
+ "slotmap",
+ "wasm-bindgen",
+ "web-sys",
+]
+
+[[package]]
+name = "glutin_wgl_sys"
+version = "0.6.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2c4ee00b289aba7a9e5306d57c2d05499b2e5dc427f84ac708bd2c090212cf3e"
+dependencies = [
+ "gl_generator",
+]
+
+[[package]]
+name = "gpu-allocator"
+version = "0.28.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "51255ea7cfaadb6c5f1528d43e92a82acb2b96c43365989a28b2d44ee38f8795"
+dependencies = [
+ "ash",
+ "hashbrown 0.16.1",
+ "log",
+ "presser",
+ "thiserror",
+ "windows",
+]
+
+[[package]]
+name = "gpu-descriptor"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b89c83349105e3732062a895becfc71a8f921bb71ecbbdd8ff99263e3b53a0ca"
+dependencies = [
+ "bitflags",
+ "gpu-descriptor-types",
+ "hashbrown 0.15.5",
+]
+
+[[package]]
+name = "gpu-descriptor-types"
+version = "0.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "fdf242682df893b86f33a73828fb09ca4b2d3bb6cc95249707fc684d27484b91"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "half"
+version = "2.7.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6ea2d84b969582b4b1864a92dc5d27cd2b77b622a8d79306834f1be5ba20d84b"
+dependencies = [
+ "cfg-if",
+ "crunchy",
+ "num-traits",
+ "zerocopy",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.15.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9229cfe53dfd69f0609a49f65461bd93001ea1ef889cd5529dd176593f5338a1"
+dependencies = [
+ "foldhash 0.1.5",
+]
+
+[[package]]
+name = "hashbrown"
+version = "0.16.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "841d1cc9bed7f9236f321df977030373f4a4163ae1a7dbfe1a51a2c1a51d9100"
+dependencies = [
+ "allocator-api2",
+ "equivalent",
+ "foldhash 0.2.0",
+]
+
+[[package]]
+name = "hex"
+version = "0.4.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f24254aa9a54b5c858eaee2f5bccdb46aaf0e486a595ed5fd8f86ba55232a70"
+
+[[package]]
+name = "hexf-parse"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df"
+
+[[package]]
+name = "indexmap"
+version = "2.13.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7714e70437a7dc3ac8eb7e6f8df75fd8eb422675fc7678aff7364301092b1017"
+dependencies = [
+ "equivalent",
+ "hashbrown 0.16.1",
+]
+
+[[package]]
+name = "is_terminal_polyfill"
+version = "1.70.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a6cb138bb79a146c1bd460005623e142ef0181e3d0219cb493e02f7d08a35695"
+
+[[package]]
+name = "jiff"
+version = "0.2.20"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c867c356cc096b33f4981825ab281ecba3db0acefe60329f044c1789d94c6543"
+dependencies = [
+ "jiff-static",
+ "log",
+ "portable-atomic",
+ "portable-atomic-util",
+ "serde_core",
+]
+
+[[package]]
+name = "jiff-static"
+version = "0.2.20"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f7946b4325269738f270bb55b3c19ab5c5040525f83fd625259422a9d25d9be5"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "jni-sys"
+version = "0.3.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130"
+
+[[package]]
+name = "js-sys"
+version = "0.3.87"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "93f0862381daaec758576dcc22eb7bbf4d7efd67328553f3b45a412a51a3fb21"
+dependencies = [
+ "once_cell",
+ "wasm-bindgen",
+]
+
+[[package]]
+name = "khronos-egl"
+version = "6.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76"
+dependencies = [
+ "libc",
+ "libloading",
+ "pkg-config",
+]
+
+[[package]]
+name = "khronos_api"
+version = "3.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc"
+
+[[package]]
+name = "libc"
+version = "0.2.182"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6800badb6cb2082ffd7b6a67e6125bb39f18782f793520caee8cb8846be06112"
+
+[[package]]
+name = "libloading"
+version = "0.8.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55"
+dependencies = [
+ "cfg-if",
+ "windows-link",
+]
+
+[[package]]
+name = "libm"
+version = "0.2.16"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b6d2cec3eae94f9f509c767b45932f1ada8350c4bdb85af2fcab4a3c14807981"
+
+[[package]]
+name = "litrs"
+version = "1.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092"
+
+[[package]]
+name = "lock_api"
+version = "0.4.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "224399e74b87b5f3557511d98dff8b14089b3dadafcab6bb93eab67d3aace965"
+dependencies = [
+ "scopeguard",
+]
+
+[[package]]
+name = "log"
+version = "0.4.29"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "5e5032e24019045c762d3c0f28f5b6b8bbf38563a65908389bf7978758920897"
+
+[[package]]
+name = "malloc_buf"
+version = "0.0.6"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb"
+dependencies = [
+ "libc",
+]
+
+[[package]]
+name = "memchr"
+version = "2.8.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f8ca58f447f06ed17d5fc4043ce1b10dd205e060fb3ce5b979b8ed8e59ff3f79"
+
+[[package]]
+name = "metal"
+version = "0.33.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c7047791b5bc903b8cd963014b355f71dc9864a9a0b727057676c1dcae5cbc15"
+dependencies = [
+ "bitflags",
+ "block",
+ "core-graphics-types",
+ "foreign-types",
+ "log",
+ "objc",
+ "paste",
+]
+
+[[package]]
+name = "naga"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "618f667225063219ddfc61251087db8a9aec3c3f0950c916b614e403486f1135"
+dependencies = [
+ "arrayvec",
+ "bit-set",
+ "bitflags",
+ "cfg-if",
+ "cfg_aliases",
+ "codespan-reporting",
+ "half",
+ "hashbrown 0.16.1",
+ "hexf-parse",
+ "indexmap",
+ "libm",
+ "log",
+ "num-traits",
+ "once_cell",
+ "rustc-hash",
+ "spirv",
+ "thiserror",
+ "unicode-ident",
+]
+
+[[package]]
+name = "ndk-sys"
+version = "0.6.0+11769913"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ee6cda3051665f1fb8d9e08fc35c96d5a244fb1be711a03b71118828afc9a873"
+dependencies = [
+ "jni-sys",
+]
+
+[[package]]
+name = "num-traits"
+version = "0.2.19"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841"
+dependencies = [
+ "autocfg",
+ "libm",
+]
+
+[[package]]
+name = "objc"
+version = "0.2.7"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1"
+dependencies = [
+ "malloc_buf",
+]
+
+[[package]]
+name = "once_cell"
+version = "1.21.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "42f5e15c9953c5e4ccceeb2e7382a716482c34515315f7b03532b8b4e8393d2d"
+
+[[package]]
+name = "once_cell_polyfill"
+version = "1.70.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "384b8ab6d37215f3c5301a95a4accb5d64aa607f1fcb26a11b5303878451b4fe"
+
+[[package]]
+name = "ordered-float"
+version = "5.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7f4779c6901a562440c3786d08192c6fbda7c1c2060edd10006b05ee35d10f2d"
+dependencies = [
+ "num-traits",
+]
+
+[[package]]
+name = "parking_lot"
+version = "0.12.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "93857453250e3077bd71ff98b6a65ea6621a19bb0f559a85248955ac12c45a1a"
+dependencies = [
+ "lock_api",
+ "parking_lot_core",
+]
+
+[[package]]
+name = "parking_lot_core"
+version = "0.9.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2621685985a2ebf1c516881c026032ac7deafcda1a2c9b7850dc81e3dfcb64c1"
+dependencies = [
+ "cfg-if",
+ "libc",
+ "redox_syscall",
+ "smallvec",
+ "windows-link",
+]
+
+[[package]]
+name = "paste"
+version = "1.0.15"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a"
+
+[[package]]
+name = "pin-project-lite"
+version = "0.2.16"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3b3cff922bd51709b605d9ead9aa71031d81447142d828eb4a6eba76fe619f9b"
+
+[[package]]
+name = "pkg-config"
+version = "0.3.32"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7edddbd0b52d732b21ad9a5fab5c704c14cd949e5e9a1ec5929a24fded1b904c"
+
+[[package]]
+name = "pollster"
+version = "0.4.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "2f3a9f18d041e6d0e102a0a46750538147e5e8992d3b4873aaafee2520b00ce3"
+
+[[package]]
+name = "portable-atomic"
+version = "1.13.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c33a9471896f1c69cecef8d20cbe2f7accd12527ce60845ff44c153bb2a21b49"
+
+[[package]]
+name = "portable-atomic-util"
+version = "0.2.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7a9db96d7fa8782dd8c15ce32ffe8680bbd1e978a43bf51a34d39483540495f5"
+dependencies = [
+ "portable-atomic",
+]
+
+[[package]]
+name = "presser"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa"
+
+[[package]]
+name = "proc-macro2"
+version = "1.0.106"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8fd00f0bb2e90d81d1044c2b32617f68fcb9fa3bb7640c23e9c748e53fb30934"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "profiling"
+version = "1.0.17"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3eb8486b569e12e2c32ad3e204dbaba5e4b5b216e9367044f25f1dba42341773"
+
+[[package]]
+name = "quote"
+version = "1.0.44"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "21b2ebcf727b7760c461f091f9f0f539b77b8e87f2fd88131e7f1b433b3cece4"
+dependencies = [
+ "proc-macro2",
+]
+
+[[package]]
+name = "range-alloc"
+version = "0.1.4"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c3d6831663a5098ea164f89cff59c6284e95f4e3c76ce9848d4529f5ccca9bde"
+
+[[package]]
+name = "raw-window-handle"
+version = "0.6.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539"
+
+[[package]]
+name = "redox_syscall"
+version = "0.5.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ed2bf2547551a7053d6fdfafda3f938979645c44812fbfcda098faae3f1a362d"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "regex"
+version = "1.12.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e10754a14b9137dd7b1e3e5b0493cc9171fdd105e0ab477f51b72e7f3ac0e276"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-automata",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-automata"
+version = "0.4.14"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e1dd4122fc1595e8162618945476892eefca7b88c52820e74af6262213cae8f"
+dependencies = [
+ "aho-corasick",
+ "memchr",
+ "regex-syntax",
+]
+
+[[package]]
+name = "regex-syntax"
+version = "0.8.9"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a96887878f22d7bad8a3b6dc5b7440e0ada9a245242924394987b21cf2210a4c"
+
+[[package]]
+name = "renderdoc-sys"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832"
+
+[[package]]
+name = "rustc-hash"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2"
+
+[[package]]
+name = "rustversion"
+version = "1.0.22"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b39cdef0fa800fc44525c84ccb54a029961a8215f9619753635a9c0d2538d46d"
+
+[[package]]
+name = "scopeguard"
+version = "1.2.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "94143f37725109f92c262ed2cf5e59bce7498c01bcc1502d7b9afe439a4e9f49"
+
+[[package]]
+name = "serde"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "9a8e94ea7f378bd32cbbd37198a4a91436180c5bb472411e48b5ec2e2124ae9e"
+dependencies = [
+ "serde_core",
+ "serde_derive",
+]
+
+[[package]]
+name = "serde_core"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "41d385c7d4ca58e59fc732af25c3983b67ac852c1a25000afe1175de458b67ad"
+dependencies = [
+ "serde_derive",
+]
+
+[[package]]
+name = "serde_derive"
+version = "1.0.228"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "d540f220d3187173da220f885ab66608367b6574e925011a9353e4badda91d79"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "slab"
+version = "0.4.12"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0c790de23124f9ab44544d7ac05d60440adc586479ce501c1d6d7da3cd8c9cf5"
+
+[[package]]
+name = "slotmap"
+version = "1.1.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "bdd58c3c93c3d278ca835519292445cb4b0d4dc59ccfdf7ceadaab3f8aeb4038"
+dependencies = [
+ "version_check",
+]
+
+[[package]]
+name = "smallvec"
+version = "1.15.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "67b1b7a3b5fe4f1376887184045fcf45c69e92af734b7aaddc05fb777b6fbd03"
+
+[[package]]
+name = "spirv"
+version = "0.3.0+sdk-1.3.268.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "eda41003dc44290527a59b13432d4a0379379fa074b70174882adfbdfd917844"
+dependencies = [
+ "bitflags",
+]
+
+[[package]]
+name = "static_assertions"
+version = "1.1.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f"
+
+[[package]]
+name = "syn"
+version = "2.0.117"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e665b8803e7b1d2a727f4023456bbbbe74da67099c585258af0ad9c5013b9b99"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "unicode-ident",
+]
+
+[[package]]
+name = "termcolor"
+version = "1.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "06794f8f6c5c898b3275aebefa6b8a1cb24cd2c6c79397ab15774837a0bc5755"
+dependencies = [
+ "winapi-util",
+]
+
+[[package]]
+name = "thiserror"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4288b5bcbc7920c07a1149a35cf9590a2aa808e0bc1eafaade0b80947865fbc4"
+dependencies = [
+ "thiserror-impl",
+]
+
+[[package]]
+name = "thiserror-impl"
+version = "2.0.18"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ebc4ee7f67670e9b64d05fa4253e753e016c6c95ff35b89b7941d6b856dec1d5"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "unicode-ident"
+version = "1.0.24"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75"
+
+[[package]]
+name = "unicode-width"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b4ac048d71ede7ee76d585517add45da530660ef4390e49b098733c6e897f254"
+
+[[package]]
+name = "utf8parse"
+version = "0.2.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "06abde3611657adf66d383f00b093d7faecc7fa57071cce2578660c9f1010821"
+
+[[package]]
+name = "version_check"
+version = "0.9.5"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0b928f33d975fc6ad9f86c8f283853ad26bdd5b10b7f1542aa2fa15e2289105a"
+
+[[package]]
+name = "wasm-bindgen"
+version = "0.2.110"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "1de241cdc66a9d91bd84f097039eb140cdc6eec47e0cdbaf9d932a1dd6c35866"
+dependencies = [
+ "cfg-if",
+ "once_cell",
+ "rustversion",
+ "wasm-bindgen-macro",
+ "wasm-bindgen-shared",
+]
+
+[[package]]
+name = "wasm-bindgen-futures"
+version = "0.4.60"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "a42e96ea38f49b191e08a1bab66c7ffdba24b06f9995b39a9dd60222e5b6f1da"
+dependencies = [
+ "cfg-if",
+ "futures-util",
+ "js-sys",
+ "once_cell",
+ "wasm-bindgen",
+ "web-sys",
+]
+
+[[package]]
+name = "wasm-bindgen-macro"
+version = "0.2.110"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e12fdf6649048f2e3de6d7d5ff3ced779cdedee0e0baffd7dff5cdfa3abc8a52"
+dependencies = [
+ "quote",
+ "wasm-bindgen-macro-support",
+]
+
+[[package]]
+name = "wasm-bindgen-macro-support"
+version = "0.2.110"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "0e63d1795c565ac3462334c1e396fd46dbf481c40f51f5072c310717bc4fb309"
+dependencies = [
+ "bumpalo",
+ "proc-macro2",
+ "quote",
+ "syn",
+ "wasm-bindgen-shared",
+]
+
+[[package]]
+name = "wasm-bindgen-shared"
+version = "0.2.110"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e9f9cdac23a5ce71f6bf9f8824898a501e511892791ea2a0c6b8568c68b9cb53"
+dependencies = [
+ "unicode-ident",
+]
+
+[[package]]
+name = "web-sys"
+version = "0.3.87"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f2c7c5718134e770ee62af3b6b4a84518ec10101aad610c024b64d6ff29bb1ff"
+dependencies = [
+ "js-sys",
+ "wasm-bindgen",
+]
+
+[[package]]
+name = "wgpu"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f9cb534d5ffd109c7d1135f34cdae29e60eab94855a625dcfe1705f8bc7ad79f"
+dependencies = [
+ "arrayvec",
+ "bitflags",
+ "bytemuck",
+ "cfg-if",
+ "cfg_aliases",
+ "document-features",
+ "hashbrown 0.16.1",
+ "js-sys",
+ "log",
+ "naga",
+ "parking_lot",
+ "portable-atomic",
+ "profiling",
+ "raw-window-handle",
+ "smallvec",
+ "static_assertions",
+ "wasm-bindgen",
+ "wasm-bindgen-futures",
+ "web-sys",
+ "wgpu-core",
+ "wgpu-hal",
+ "wgpu-types",
+]
+
+[[package]]
+name = "wgpu-core"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "8bb4c8b5db5f00e56f1f08869d870a0dff7c8bc7ebc01091fec140b0cf0211a9"
+dependencies = [
+ "arrayvec",
+ "bit-set",
+ "bit-vec",
+ "bitflags",
+ "bytemuck",
+ "cfg_aliases",
+ "document-features",
+ "hashbrown 0.16.1",
+ "indexmap",
+ "log",
+ "naga",
+ "once_cell",
+ "parking_lot",
+ "portable-atomic",
+ "profiling",
+ "raw-window-handle",
+ "rustc-hash",
+ "smallvec",
+ "thiserror",
+ "wgpu-core-deps-apple",
+ "wgpu-core-deps-emscripten",
+ "wgpu-core-deps-windows-linux-android",
+ "wgpu-hal",
+ "wgpu-types",
+]
+
+[[package]]
+name = "wgpu-core-deps-apple"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "87b7b696b918f337c486bf93142454080a32a37832ba8a31e4f48221890047da"
+dependencies = [
+ "wgpu-hal",
+]
+
+[[package]]
+name = "wgpu-core-deps-emscripten"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "34b251c331f84feac147de3c4aa3aa45112622a95dd7ee1b74384fa0458dbd79"
+dependencies = [
+ "wgpu-hal",
+]
+
+[[package]]
+name = "wgpu-core-deps-windows-linux-android"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "68ca976e72b2c9964eb243e281f6ce7f14a514e409920920dcda12ae40febaae"
+dependencies = [
+ "wgpu-hal",
+]
+
+[[package]]
+name = "wgpu-hal"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "293080d77fdd14d6b08a67c5487dfddbf874534bb7921526db56a7b75d7e3bef"
+dependencies = [
+ "android_system_properties",
+ "arrayvec",
+ "ash",
+ "bit-set",
+ "bitflags",
+ "block",
+ "bytemuck",
+ "cfg-if",
+ "cfg_aliases",
+ "core-graphics-types",
+ "glow",
+ "glutin_wgl_sys",
+ "gpu-allocator",
+ "gpu-descriptor",
+ "hashbrown 0.16.1",
+ "js-sys",
+ "khronos-egl",
+ "libc",
+ "libloading",
+ "log",
+ "metal",
+ "naga",
+ "ndk-sys",
+ "objc",
+ "once_cell",
+ "ordered-float",
+ "parking_lot",
+ "portable-atomic",
+ "portable-atomic-util",
+ "profiling",
+ "range-alloc",
+ "raw-window-handle",
+ "renderdoc-sys",
+ "smallvec",
+ "thiserror",
+ "wasm-bindgen",
+ "web-sys",
+ "wgpu-types",
+ "windows",
+ "windows-core",
+]
+
+[[package]]
+name = "wgpu-types"
+version = "28.0.0"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e18308757e594ed2cd27dddbb16a139c42a683819d32a2e0b1b0167552f5840c"
+dependencies = [
+ "bitflags",
+ "bytemuck",
+ "js-sys",
+ "log",
+ "web-sys",
+]
+
+[[package]]
+name = "winapi-util"
+version = "0.1.11"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22"
+dependencies = [
+ "windows-sys",
+]
+
+[[package]]
+name = "windows"
+version = "0.62.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "527fadee13e0c05939a6a05d5bd6eec6cd2e3dbd648b9f8e447c6518133d8580"
+dependencies = [
+ "windows-collections",
+ "windows-core",
+ "windows-future",
+ "windows-numerics",
+]
+
+[[package]]
+name = "windows-collections"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "23b2d95af1a8a14a3c7367e1ed4fc9c20e0a26e79551b1454d72583c97cc6610"
+dependencies = [
+ "windows-core",
+]
+
+[[package]]
+name = "windows-core"
+version = "0.62.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "b8e83a14d34d0623b51dce9581199302a221863196a1dde71a7663a4c2be9deb"
+dependencies = [
+ "windows-implement",
+ "windows-interface",
+ "windows-link",
+ "windows-result",
+ "windows-strings",
+]
+
+[[package]]
+name = "windows-future"
+version = "0.3.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "e1d6f90251fe18a279739e78025bd6ddc52a7e22f921070ccdc67dde84c605cb"
+dependencies = [
+ "windows-core",
+ "windows-link",
+ "windows-threading",
+]
+
+[[package]]
+name = "windows-implement"
+version = "0.60.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "053e2e040ab57b9dc951b72c264860db7eb3b0200ba345b4e4c3b14f67855ddf"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "windows-interface"
+version = "0.59.3"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3f316c4a2570ba26bbec722032c4099d8c8bc095efccdc15688708623367e358"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
+
+[[package]]
+name = "windows-link"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5"
+
+[[package]]
+name = "windows-numerics"
+version = "0.3.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "6e2e40844ac143cdb44aead537bbf727de9b044e107a0f1220392177d15b0f26"
+dependencies = [
+ "windows-core",
+ "windows-link",
+]
+
+[[package]]
+name = "windows-result"
+version = "0.4.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7781fa89eaf60850ac3d2da7af8e5242a5ea78d1a11c49bf2910bb5a73853eb5"
+dependencies = [
+ "windows-link",
+]
+
+[[package]]
+name = "windows-strings"
+version = "0.5.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "7837d08f69c77cf6b07689544538e017c1bfcf57e34b4c0ff58e6c2cd3b37091"
+dependencies = [
+ "windows-link",
+]
+
+[[package]]
+name = "windows-sys"
+version = "0.61.2"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "ae137229bcbd6cdf0f7b80a31df61766145077ddf49416a728b02cb3921ff3fc"
+dependencies = [
+ "windows-link",
+]
+
+[[package]]
+name = "windows-threading"
+version = "0.2.1"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3949bd5b99cafdf1c7ca86b43ca564028dfe27d66958f2470940f73d86d75b37"
+dependencies = [
+ "windows-link",
+]
+
+[[package]]
+name = "xml-rs"
+version = "0.8.28"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "3ae8337f8a065cfc972643663ea4279e04e7256de865aa66fe25cec5fb912d3f"
+
+[[package]]
+name = "zerocopy"
+version = "0.8.39"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "db6d35d663eadb6c932438e763b262fe1a70987f9ae936e60158176d710cae4a"
+dependencies = [
+ "zerocopy-derive",
+]
+
+[[package]]
+name = "zerocopy-derive"
+version = "0.8.39"
+source = "registry+https://github.com/rust-lang/crates.io-index"
+checksum = "4122cd3169e94605190e77839c9a40d40ed048d305bfdc146e7df40ab0f3e517"
+dependencies = [
+ "proc-macro2",
+ "quote",
+ "syn",
+]
diff --git a/Cargo.toml b/Cargo.toml
new file mode 100644
index 0000000..95bf379
--- /dev/null
+++ b/Cargo.toml
@@ -0,0 +1,11 @@
+[package]
+name = "bitcoin-rgpu"
+version = "0.1.0"
+edition = "2024"
+
+[dependencies]
+bytemuck = "1.22.0"
+env_logger = "0.11"
+hex = "0.4.3"
+pollster = "0.4"
+wgpu = "28.0.0"
diff --git a/flake.lock b/flake.lock
new file mode 100644
index 0000000..4a7ecf7
--- /dev/null
+++ b/flake.lock
@@ -0,0 +1,137 @@
+{
+  "nodes": {
+    "crane": {
+      "locked": {
+        "lastModified": 1771438068,
+        "narHash": "sha256-nGBbXvEZVe/egCPVPFcu89RFtd8Rf6J+4RFoVCFec0A=",
+        "owner": "ipetkov",
+        "repo": "crane",
+        "rev": "b5090e53e9d68c523a4bb9ad42b4737ee6747597",
+        "type": "github"
+      },
+      "original": {
+        "owner": "ipetkov",
+        "repo": "crane",
+        "type": "github"
+      }
+    },
+    "fenix": {
+      "inputs": {
+        "nixpkgs": [
+          "nixpkgs"
+        ],
+        "rust-analyzer-src": "rust-analyzer-src"
+      },
+      "locked": {
+        "lastModified": 1771657407,
+        "narHash": "sha256-ZkKTcOgfy8lIkdTNpGfc01G21d5oT7xMqGT4cazFl60=",
+        "owner": "nix-community",
+        "repo": "fenix",
+        "rev": "6d86ae5186ba275b7874d02908505fc658448658",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nix-community",
+        "repo": "fenix",
+        "type": "github"
+      }
+    },
+    "flake-utils": {
+      "inputs": {
+        "systems": "systems"
+      },
+      "locked": {
+        "lastModified": 1731533236,
+        "narHash": "sha256-l0KFg5HjrsfsO/JpG+r7fRrqm12kzFHyUHqHCVpMMbI=",
+        "owner": "numtide",
+        "repo": "flake-utils",
+        "rev": "11707dc2f618dd54ca8739b309ec4fc024de578b",
+        "type": "github"
+      },
+      "original": {
+        "owner": "numtide",
+        "repo": "flake-utils",
+        "type": "github"
+      }
+    },
+    "nixpkgs": {
+      "locked": {
+        "lastModified": 1771369470,
+        "narHash": "sha256-0NBlEBKkN3lufyvFegY4TYv5mCNHbi5OmBDrzihbBMQ=",
+        "owner": "nixos",
+        "repo": "nixpkgs",
+        "rev": "0182a361324364ae3f436a63005877674cf45efb",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nixos",
+        "ref": "nixos-unstable",
+        "repo": "nixpkgs",
+        "type": "github"
+      }
+    },
+    "root": {
+      "inputs": {
+        "crane": "crane",
+        "fenix": "fenix",
+        "flake-utils": "flake-utils",
+        "nixpkgs": "nixpkgs",
+        "treefmt-nix": "treefmt-nix"
+      }
+    },
+    "rust-analyzer-src": {
+      "flake": false,
+      "locked": {
+        "lastModified": 1771585531,
+        "narHash": "sha256-xMDlGXSLxr5F2lt7ydy3ve4rE7xFVvBx1u8OyYOXyDI=",
+        "owner": "rust-lang",
+        "repo": "rust-analyzer",
+        "rev": "46a214b5d35f2e2a02c0b278d9287b20760821eb",
+        "type": "github"
+      },
+      "original": {
+        "owner": "rust-lang",
+        "ref": "nightly",
+        "repo": "rust-analyzer",
+        "type": "github"
+      }
+    },
+    "systems": {
+      "locked": {
+        "lastModified": 1681028828,
+        "narHash": "sha256-Vy1rq5AaRuLzOxct8nz4T6wlgyUR7zLU309k9mBC768=",
+        "owner": "nix-systems",
+        "repo": "default",
+        "rev": "da67096a3b9bf56a91d16901293e51ba5b49a27e",
+        "type": "github"
+      },
+      "original": {
+        "owner": "nix-systems",
+        "repo": "default",
+        "type": "github"
+      }
+    },
+    "treefmt-nix": {
+      "inputs": {
+        "nixpkgs": [
+          "nixpkgs"
+        ]
+      },
+      "locked": {
+        "lastModified": 1770228511,
+        "narHash": "sha256-wQ6NJSuFqAEmIg2VMnLdCnUc0b7vslUohqqGGD+Fyxk=",
+        "owner": "numtide",
+        "repo": "treefmt-nix",
+        "rev": "337a4fe074be1042a35086f15481d763b8ddc0e7",
+        "type": "github"
+      },
+      "original": {
+        "owner": "numtide",
+        "repo": "treefmt-nix",
+        "type": "github"
+      }
+    }
+  },
+  "root": "root",
+  "version": 7
+}
diff --git a/flake.nix b/flake.nix
new file mode 100644
index 0000000..a413f57
--- /dev/null
+++ b/flake.nix
@@ -0,0 +1,115 @@
+{
+  inputs = {
+    nixpkgs.url = "github:nixos/nixpkgs/nixos-unstable";
+    flake-utils.url = "github:numtide/flake-utils";
+    crane.url = "github:ipetkov/crane";
+    treefmt-nix = {
+      url = "github:numtide/treefmt-nix";
+      inputs.nixpkgs.follows = "nixpkgs";
+    };
+
+    fenix = {
+      url = "github:nix-community/fenix";
+      inputs.nixpkgs.follows = "nixpkgs";
+    };
+  };
+  outputs =
+    {
+      self,
+      nixpkgs,
+      flake-utils,
+      crane,
+      fenix,
+      treefmt-nix,
+      ...
+    }:
+    flake-utils.lib.eachDefaultSystem (
+      system:
+      let
+        pkgs = nixpkgs.legacyPackages.${system};
+        lib = pkgs.lib;
+
+        rustToolchain = fenix.packages.${system}.fromToolchainFile {
+          file = ./rust-toolchain.toml;
+          sha256 = "sha256-lEeYHNJjEVlYS+kLomkfpFWNEJQFhgSjuISZwMdWmb8=" ;
+        };
+
+        craneLib = (crane.mkLib pkgs).overrideToolchain rustToolchain;
+        src = (craneLib.cleanCargoSource ./.);
+
+        commonArgs = {
+          inherit src;
+          strictDeps = true;
+          cargoToml = ./Cargo.toml;
+    
+          nativeBuildInputs = [
+            pkgs.pkg-config
+          ];
+        };
+
+        cargoArtifacts = craneLib.buildDepsOnly commonArgs;
+
+        simple-pow = craneLib.buildPackage (
+          commonArgs
+          // {
+            inherit cargoArtifacts;
+          }
+        );
+
+        tree-fmt-eval = treefmt-nix.lib.evalModule pkgs (
+          { ... }:
+          {
+            # Used to find the project root
+            projectRootFile = "flake.nix";
+            programs.nixfmt.enable = true;
+            programs.deadnix.enable = true;
+            programs.rustfmt = {
+              enable = true;
+              package = rustToolchain;
+            };
+          }
+        );
+
+      in
+      {
+        checks = {
+          inherit simple-pow;
+          str-clippy = craneLib.cargoClippy (
+            commonArgs
+            // {
+              inherit cargoArtifacts;
+              cargoClippyExtraArgs = "--all-targets -- --deny warnings";
+            }
+          );
+          formatting = tree-fmt-eval.config.build.check self;
+        };
+
+        formatter = tree-fmt-eval.config.build.wrapper;
+
+        apps.default = flake-utils.lib.mkApp {
+          drv = simple-pow;
+        };
+
+        devShells.default = craneLib.devShell {
+          checks = self.checks.${system};
+
+          packages = with pkgs; [
+            cargo-watch
+            python3
+            wgsl-analyzer
+          ];
+
+          env = {
+            "LD_LIBRARY_PATH" = lib.makeLibraryPath (with pkgs; [ stdenv.cc.cc.lib
+   vulkan-loader
+          xorg.libX11
+          xorg.libXcursor
+          xorg.libXi
+          xorg.libXrandr
+          libxkbcommon
+            ]); 
+          };
+        };
+      }
+    );
+}
diff --git a/palette.txt b/palette.txt
new file mode 100644
index 0000000..8d188c4
--- /dev/null
+++ b/palette.txt
@@ -0,0 +1 @@
+!"#$%&'()*+,-./0123456789:;<=>?@ABCDEFGHIJKLMNOPQRSTUVWXYZ[\]^_`abcdefghijklmnopqrstuvwxyz{|}~
diff --git a/rust-toolchain.toml b/rust-toolchain.toml
new file mode 100644
index 0000000..39171ab
--- /dev/null
+++ b/rust-toolchain.toml
@@ -0,0 +1,13 @@
+[toolchain]
+channel = "nightly"
+targets = [
+    #linux 
+    "x86_64-unknown-linux-gnu",
+    "aarch64-unknown-linux-gnu",
+    "x86_64-unknown-linux-musl",
+    "aarch64-unknown-linux-musl",
+
+    # WebAssembly
+    "wasm32-unknown-unknown"
+]
+components = ["rust-src", "clippy", "rustfmt","rust-analyzer", "llvm-tools"] 
diff --git a/src/main.rs b/src/main.rs
new file mode 100644
index 0000000..54d9610
--- /dev/null
+++ b/src/main.rs
@@ -0,0 +1,489 @@
+use std::env::{self, Args};
+use std::mem;
+use std::num::NonZeroU64;
+use std::ops::RangeBounds;
+use std::time::Instant;
+
+use wgpu::{BufferUsages, SubmissionIndex, include_wgsl};
+use wgpu::util::DeviceExt;
+use wgpu::wgc::command;
+
+
+
+const PALETTE: &[u8] = include_bytes!("../palette.txt"); 
+const PALETTE_SIZE: usize =  PALETTE.len();
+
+const BATCH_SIZE : u64 = (1 << 16) - 64; 
+const MAX_ITER : u64 = 1 << 32;
+// const BATCH_SIZE : u64 = 8; 
+// const MAX_ITER : u64 = 2; 
+
+const INPUT_SIZE:u64 =  512 / 8 * 8;
+const INSIZE_SIZE:u64= 32 / 8; 
+const OUTPUT_SIZE:u64= 256 / 8; 
+const INPUT_BUF_SIZE :u64= INPUT_SIZE * BATCH_SIZE; 
+const SIZES_BUF_SIZE :u64= INSIZE_SIZE * BATCH_SIZE;
+const OUTPUT_BUF_SIZE:u64= OUTPUT_SIZE * BATCH_SIZE;  
+
+fn get_arg(args: &mut Args, pname: &String , name: &'static str) -> String {
+    args.next().unwrap_or_else(|| panic!("Argument missing: {}\nUsage: {} value prev bits", name, pname)).clone()
+}
+
+
+struct NonceIter {
+    digits: Vec<usize>, 
+}
+
+impl NonceIter {
+    fn new_with_capacity(capacity: usize) -> Self {
+        Self {
+            digits: Vec::with_capacity(capacity)    
+        }
+    }
+} 
+
+impl Iterator for NonceIter {
+    type Item = NonceElement;
+    fn next(&mut self) -> Option<Self::Item> {
+        let mut carry = 1; 
+        let mut cursor = 0; 
+        while carry > 0 { 
+           if let Some(d) = self.digits.get_mut(cursor)  {
+                if *d + carry >= PALETTE_SIZE {
+                   *d = (*d + carry) % PALETTE_SIZE; 
+                } else {
+                   *d += carry; 
+                   carry = 0; 
+                }
+           } else {
+                self.digits.push(carry);
+                carry = 0; 
+           } 
+           cursor += 1;
+        }
+        Some(NonceElement::from_digits(&self.digits))
+    }
+
+}
+
+#[derive(Debug)]
+struct NonceElement {
+    pub bytes: Vec<u8>
+}
+
+impl NonceElement {
+    fn from_digits(digits: &Vec<usize>) -> Self {
+        let mut bytes = Vec::with_capacity(digits.len());
+        for d in digits {
+           bytes.push(PALETTE[*d]) 
+        } 
+        Self {
+            bytes
+        }
+    }
+    fn to_hash_input(&self, value: &String, prev: &String) -> String {
+        let mut out = String::with_capacity(self.bytes.len() + value.len() + prev.len()); 
+        out += value; 
+        out += prev; 
+        out += str::from_utf8(self.bytes.as_slice()).expect("Invalid bytes in pallette");
+        out 
+    }
+}
+
+
+/// Swap chain. We Have Swap Chains At Home Edition. 
+/// When given an index, returns the first tuple entry if the index is even, and the second if it's
+/// odd. 
+fn swaptuple_get<T>(tup :&(T,T), i: u64) -> &T {
+    if i.is_multiple_of(2){ &tup.0 } else { &tup.1 }  
+} 
+/// Swap chain. We Have Swap Chains At Home Edition. Mutable Edition.  
+/// When given an index, returns the first tuple entry if the index is even, and the second if it's
+/// odd, but now mutable. 
+fn swaptuple_get_mut<T>(tup :& mut (T,T), i: u64) -> & mut T {
+    if i.is_multiple_of(2) { &mut tup.0 } else { &mut tup.1 }  
+} 
+
+fn main() {
+    let mut args = env::args(); 
+
+    let pname: String = args.next().expect("Program name should always be included.");
+    let value: String = get_arg(&mut args, &pname, "value");
+    let prev: String = get_arg(&mut args, &pname, "prev");
+    let bits: u16 = get_arg(&mut args, &pname, "bits").parse().unwrap();
+
+    // We first initialize an wgpu `Instance`, which contains any "global" state wgpu needs.
+    //
+    // This is what loads the vulkan/dx12/metal/opengl libraries.
+    env_logger::init();
+
+    let instance = wgpu::Instance::new(&wgpu::InstanceDescriptor::from_env_or_default());
+    let areq = wgpu::RequestAdapterOptions {power_preference: wgpu::PowerPreference::HighPerformance, ..Default::default()};
+    let adapter =
+        pollster::block_on(instance.request_adapter(&areq))
+            .expect("Failed to create adapter");
+
+    // Print out some basic information about the adapter.
+    println!("Running on Adapter: {:#?}", adapter.get_info());
+
+    // Check to see if the adapter supports compute shaders. While WebGPU guarantees support for
+    // compute shaders, wgpu supports a wider range of devices through the use of "downlevel" devices.
+    let downlevel_capabilities = adapter.get_downlevel_capabilities();
+    if !downlevel_capabilities
+        .flags
+        .contains(wgpu::DownlevelFlags::COMPUTE_SHADERS)
+    {
+        panic!("Adapter does not support compute shaders");
+    }
+
+    // We then create a `Device` and a `Queue` from the `Adapter`.
+    //
+    // The `Device` is used to create and manage GPU resources.
+    // The `Queue` is a queue used to submit work for the GPU to process.
+    let (device, queue) = pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
+        label: None,
+        required_features: wgpu::Features::empty(),
+        required_limits: wgpu::Limits::downlevel_defaults(),
+        experimental_features: wgpu::ExperimentalFeatures::disabled(),
+        memory_hints: wgpu::MemoryHints::MemoryUsage,
+        trace: wgpu::Trace::Off,
+    }))
+    .expect("Failed to create device");
+
+    // Create a shader module from our shader code. This will parse and validate the shader.
+    //
+    // `include_wgsl` is a macro provided by wgpu like `include_str` which constructs a ShaderModuleDescriptor.
+    // If you want to load shaders differently, you can construct the ShaderModuleDescriptor manually.
+    let module = device.create_shader_module(wgpu::include_wgsl!("shader_own.wgsl"));
+    
+    let mut iter = NonceIter::new_with_capacity((512 - prev.len() - value.len()));
+
+    // Create a buffer with the data we want to process on the GPU.
+    //
+    // `create_buffer_init` is a utility provided by `wgpu::util::DeviceExt` which simplifies creating
+    // a buffer with some initial data.
+    //
+    // We use the `bytemuck` crate to cast the slice of f32 to a &[u8] to be uploaded to the GPU.
+    let input_data_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("Input"),
+        size:  INPUT_BUF_SIZE,
+        usage: wgpu::BufferUsages::union(BufferUsages::STORAGE, BufferUsages::COPY_DST),
+        mapped_at_creation: false,
+    });
+    
+    let data_size_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("Data Size"),
+        size:  SIZES_BUF_SIZE,
+        usage: wgpu::BufferUsages::union(BufferUsages::STORAGE, BufferUsages::COPY_DST),
+        mapped_at_creation: false,
+    });
+
+    // Now we create a buffer to store the output data.
+    let output_data_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("Output"),
+        size: OUTPUT_BUF_SIZE,
+        usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
+        mapped_at_creation: false,
+    });
+
+    // Finally we create a buffer which can be read by the CPU. This buffer is how we will read
+    // the data. We need to use a separate buffer because we need to have a usage of `MAP_READ`,
+    // and that usage can only be used with `COPY_DST`.
+    let download_buffer_0 = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("Download_0"),
+        size: OUTPUT_BUF_SIZE,
+        usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+        mapped_at_creation: false,
+    });
+
+    // We use a second buffer to flip stuff
+    let download_buffer_1 = device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some("Download_1"),
+        size: OUTPUT_BUF_SIZE,
+        usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+        mapped_at_creation: false,
+    });
+
+    let dl_bufs = (download_buffer_0, download_buffer_1); 
+
+    // A bind group layout describes the types of resources that a bind group can contain. Think
+    // of this like a C-style header declaration, ensuring both the pipeline and bind group agree
+    // on the types of resources.
+    let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+        label: None,
+        entries: &[
+            // Input buffer
+            wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: true },
+                    // This is the size of a single element in the buffer.
+                    min_binding_size: Some(NonZeroU64::new(512).unwrap()),
+                    has_dynamic_offset: false,
+                },
+                count: None,
+            },
+            wgpu::BindGroupLayoutEntry {
+                binding: 1, 
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: true },
+                    // This is the size of a single element in the buffer.
+                    min_binding_size: Some(NonZeroU64::new(4).unwrap()),
+                    has_dynamic_offset: false,
+                },
+                count: None,
+            },
+        
+            // Output buffer
+            wgpu::BindGroupLayoutEntry {
+                binding: 2,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: false },
+                    // This is the size of a single element in the buffer.
+                    min_binding_size: Some(NonZeroU64::new(32).unwrap()),
+                    has_dynamic_offset: false,
+                },
+                count: None,
+            },
+        ],
+    });
+
+    // The bind group contains the actual resources to bind to the pipeline.
+    //
+    // Even when the buffers are individually dropped, wgpu will keep the bind group and buffers
+    // alive until the bind group itself is dropped.
+    let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+        label: None,
+        layout: &bind_group_layout,
+        entries: &[
+            wgpu::BindGroupEntry {
+                binding: 0,
+                resource: input_data_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 1,
+                resource: data_size_buffer.as_entire_binding(),
+            },
+            wgpu::BindGroupEntry {
+                binding: 2,
+                resource: output_data_buffer.as_entire_binding(),
+            },
+        ],
+    });
+
+    // The pipeline layout describes the bind groups that a pipeline expects
+    let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+        label: None,
+        bind_group_layouts: &[&bind_group_layout],
+        immediate_size: 0,
+    });
+
+    // The pipeline is the ready-to-go program state for the GPU. It contains the shader modules,
+    // the interfaces (bind group layouts) and the shader entry point.
+    let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+        label: None,
+        layout: Some(&pipeline_layout),
+        module: &module,
+        entry_point: Some("main"),
+        compilation_options: wgpu::PipelineCompilationOptions::default(),
+        cache: None,
+    });
+
+        
+    let mut nibbles: [u16; OUTPUT_SIZE as usize/ 8] = [0;4];
+    let mut b = bits as u16; 
+    for n in 0..4 {
+        if b > 64 {
+            nibbles[n] = 64;
+            b -= 64;
+        } else {
+            nibbles[n] = b; 
+            break;
+        }
+    }
+
+    // == Begin repeating part 
+    let mut found: Option<[u8;32]> = None; 
+    let mut found_input: Option<String> = None;
+    let mut nonces_1: Vec<NonceElement>= Vec::with_capacity(BATCH_SIZE as usize);
+    let mut nonces_2: Vec<NonceElement>= Vec::with_capacity(BATCH_SIZE as usize);
+    let mut nonce_bufs = (nonces_1, nonces_2); 
+    let mut prev_time: Instant = Instant::now(); 
+    let mut prev_submission_index = None; 
+    println!("starting!");
+    for x in 0 .. MAX_ITER {
+        let data_upload_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Upload Data"),
+            size: INPUT_BUF_SIZE,
+            usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::MAP_WRITE,
+            mapped_at_creation: true,
+        });
+        let size_upload_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("Upload Size"),
+            size: SIZES_BUF_SIZE,
+            usage: wgpu::BufferUsages::COPY_SRC | wgpu::BufferUsages::MAP_WRITE,
+            mapped_at_creation: true,
+        });
+        let mut input_slice = data_upload_buffer.get_mapped_range_mut(..); 
+        let mut len_slice   = size_upload_buffer.get_mapped_range_mut(..); 
+
+        let nonces = swaptuple_get_mut(&mut nonce_bufs, x);
+        for i in 0 .. BATCH_SIZE {
+          if let Some(element) = iter.next() {
+              let inputdata = element.to_hash_input(&value, &prev).into_bytes(); 
+              let inputlen  = inputdata.len() as u32;
+              let lenbytes  = inputlen.to_le_bytes(); 
+             
+              // println!("{}: {}, {}", i, inputlen*8, str::from_utf8(&inputdata).unwrap());
+              input_slice[(i * INPUT_SIZE) as usize .. (i*INPUT_SIZE + inputlen as u64) as usize].copy_from_slice(&inputdata);
+              len_slice[(i * INSIZE_SIZE) as usize .. ((i+1) * INSIZE_SIZE) as usize].copy_from_slice(&lenbytes);
+              // println!("In data {:?}", &input_slice[(i * INPUT_SIZE) as usize .. (i*INPUT_SIZE + inputlen as u64) as usize]);
+              // println!("In length {:?}", &len_slice[(i * INSIZE_SIZE) as usize .. ((i+1) * INSIZE_SIZE) as usize]);
+              nonces.push(element);
+          } else {
+                panic!("Infinite iter died")
+          }
+        }
+        // println!("Hex slice:\n{}", hex::encode(&input_slice[..]));
+        drop(input_slice);       // The command encoder allows us to record commands that we will later submit to the GPU.
+        drop(len_slice); 
+        data_upload_buffer.unmap();
+        size_upload_buffer.unmap();
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+         
+        // We add a copy operation to the encoder. This will copy the data from the output buffer on the
+        // GPU to the download buffer on the CPU.
+        encoder.copy_buffer_to_buffer(
+            &data_upload_buffer,
+            0,
+            &input_data_buffer,
+            0,
+            data_upload_buffer.size(),
+        );
+        encoder.copy_buffer_to_buffer(
+            &size_upload_buffer,
+            0,
+            &data_size_buffer,
+            0,
+            size_upload_buffer.size(),
+        );
+
+        // A compute pass is a single series of compute operations. While we are recording a compute
+        // pass, we cannot record to the encoder.
+        let mut compute_pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: None,
+            timestamp_writes: None,
+        });
+
+        // Set the pipeline that we want to use
+        compute_pass.set_pipeline(&pipeline);
+        // Set the bind group that we want to use
+        compute_pass.set_bind_group(0, &bind_group, &[]);
+
+        // Now we dispatch a series of workgroups. Each workgroup is a 3D grid of individual programs.
+        //
+        // We defined the workgroup size in the shader as 64x1x1. So in order to process all of our
+        // inputs, we ceiling divide the number of inputs by 64. If the user passes 32 inputs, we will
+        // dispatch 1 workgroups. If the user passes 65 inputs, we will dispatch 2 workgroups, etc.
+        
+        let rest = BATCH_SIZE % 64; 
+        let mut workgroup_count = BATCH_SIZE - rest / 64;
+        if rest > 0  {
+            workgroup_count += 1;
+        }
+        compute_pass.dispatch_workgroups(workgroup_count as u32, 1, 1);
+
+        // Now we drop the compute pass, giving us access to the encoder again.
+        drop(compute_pass);
+
+        // We add a copy operation to the encoder. This will copy the data from the output buffer on the
+        // GPU to the download buffer on the CPU.
+        encoder.copy_buffer_to_buffer(
+            &output_data_buffer,
+            0,
+            swaptuple_get(&dl_bufs, x),
+            0,
+            output_data_buffer.size(),
+        );
+
+        // We finish the encoder, giving us a fully recorded command buffer.
+        let command_buffer = encoder.finish();
+
+
+        
+        let sub_index = queue.submit([command_buffer]);
+        // println!("A pass has been submitted");
+        // Wait for the GPU to finish working on the submitted work. This doesn't work on WebGPU, so we would need
+        // to rely on the callback to know when the buffer is mapped.
+        
+        if x > 0 {
+            let nonces = swaptuple_get_mut(&mut nonce_bufs, x -1);
+            let download_buffer = swaptuple_get(&dl_bufs, x-1); 
+            let buffer_slice = download_buffer.slice(..);
+            buffer_slice.map_async(wgpu::MapMode::Read, |_| {});
+            device.poll(wgpu::PollType::Wait { submission_index: prev_submission_index, timeout: None }).unwrap();
+            // In this case we know exactly when the mapping will be finished,
+            // so we don't need to do anything in the callback.
+            let data = buffer_slice.get_mapped_range();
+            // println!("Out data {:?}", &data[..]);
+            // println!("Out hash {:?}", hex::encode(&data[..]));
+            let results: &[u8] = bytemuck::cast_slice(&data[..]);
+            
+            // println!("Full buffer: {}", hex::encode(results) );
+            for i in 0..(BATCH_SIZE as usize) {
+                let result: &[u8;32] = &results[32*i..32*(i+1)].try_into().unwrap();  
+                // println!("Hash of {}, {}", hex::encode(nonces.get(i).unwrap().bytes.clone()), hex::encode(bytemuck::cast_slice(result)));
+                let mut correct = true; 
+                for n in 0..4 {
+                    if u64::from_be_bytes(result[n*8..(n+1)*8].try_into().unwrap()).leading_zeros() < nibbles[n] as u32{
+                        correct = false; 
+                        break;
+                    }
+                }
+                if correct {
+                    found = Some(*result);
+                    let nonce = nonces.get(i).expect("We pushed all those nonces, right?").bytes.clone(); 
+                    found_input = Some(String::from_utf8(nonce).unwrap());
+                    // break;
+                }
+            }
+            if found.is_some() {
+                break;
+            }
+            drop(data);
+            download_buffer.unmap();
+            nonces.clear();
+            let now = Instant::now(); 
+            let delta = (now - prev_time).as_secs_f64(); 
+            let hashrate = BATCH_SIZE as f64 / delta; 
+             println!("\x1b[A\rFinished batch {}, {} hashes processed ({:.2} H/s) (delta: {}) ", x + 1, x * BATCH_SIZE, hashrate, delta);
+            prev_time = Instant::now(); 
+        }
+        prev_submission_index = Some(sub_index); 
+     
+    }
+
+    match found {
+        Some(f) => {
+            let hashbytes: &[u8] = bytemuck::cast_slice(&f);
+            let nonce  = found_input.unwrap();
+            println!("Found! {:?}", hex::encode(hashbytes));
+            println!("Full block: {}{}{}", value, prev , nonce);
+            println!("Nonce: \"{}\"", nonce);
+        }
+        None => {
+            println!("Could not find a valid hash")
+        }
+    }
+
+
+
+    println!("Done for now")
+
+
+}
diff --git a/src/shader.wgsl b/src/shader.wgsl
new file mode 100644
index 0000000..39b53d5
--- /dev/null
+++ b/src/shader.wgsl
@@ -0,0 +1,208 @@
+struct SHA256_CTX {
+  data : array<u32, 64>,
+  datalen : u32,
+  bitlen : array<u32, 2>,
+  state : array<u32, 8>,
+  info : u32,
+};
+
+@group(0) @binding(0) var<storage, read> inputs : array<u32>;
+@group(0) @binding(1) var<storage, read> inputSizes : array<u32>;
+@group(0) @binding(2) var<storage, read_write> results : array<u32>;
+
+const SHA256_BLOCK_SIZE = 32;
+const INPUT_BLOCK_SIZE  = 64; 
+const INPUT_BLOCKS_MAX  = 8; 
+const INPUT_SLOT_SIZE   = INPUT_BLOCK_SIZE * INPUT_BLOCKS_MAX / 4;
+
+const k = array<u32, 64> (
+  0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
+  0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174,
+  0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da,
+  0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967,
+  0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85,
+  0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070,
+  0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3,
+  0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2
+);
+
+fn ROTLEFT(a : u32, b : u32) -> u32{return (((a) << (b)) | ((a) >> (32-(b))));}
+fn ROTRIGHT(a : u32, b : u32) -> u32{return (((a) >> (b)) | ((a) << (32-(b))));}
+
+fn CH(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ (~(x) & (z)));}
+fn MAJ(x : u32, y : u32, z : u32) -> u32{return (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z)));}
+fn EP0(x : u32) -> u32{return (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22));}
+fn EP1(x : u32) -> u32{return (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25));}
+fn SIG0(x : u32) -> u32{return (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3));}
+fn SIG1(x : u32) -> u32{return (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10));}
+
+fn sha256_transform(ctx : ptr<function, SHA256_CTX>)
+{
+  var a : u32;
+  var b : u32;
+  var c : u32;
+  var d : u32;
+  var e : u32;
+  var f : u32;
+  var g : u32;
+  var h : u32;
+  var i : u32 = 0;
+  var j : u32 = 0;
+  var t1 : u32;
+  var t2 : u32;
+  var m : array<u32, 64> ;
+
+
+  while(i < 16) {
+    m[i] = ((*ctx).data[j] << 24) | ((*ctx).data[j + 1] << 16) | ((*ctx).data[j + 2] << 8) | ((*ctx).data[j + 3]);
+    i++;
+    j += 4;
+  }            
+
+  while(i < 64) {
+        m[i] = SIG1(m[i - 2]) + m[i - 7] + SIG0(m[i - 15]) + m[i - 16];
+    i++;
+  }
+
+  a = (*ctx).state[0];
+  b = (*ctx).state[1];
+  c = (*ctx).state[2];
+  d = (*ctx).state[3];
+  e = (*ctx).state[4];
+  f = (*ctx).state[5];
+  g = (*ctx).state[6];
+  h = (*ctx).state[7];
+
+  i = 0;
+  for (; i < 64; i++) {
+    t1 = h + EP1(e) + CH(e,f,g) + k[i] + m[i];
+    t2 = EP0(a) + MAJ(a,b,c);
+    h = g;
+    g = f;
+    f = e;
+    e = d + t1;
+    d = c;
+    c = b;
+    b = a;
+    a = t1 + t2;
+  }
+
+
+  (*ctx).state[0] += a;
+  (*ctx).state[1] += b;
+  (*ctx).state[2] += c;
+  (*ctx).state[3] += d;
+  (*ctx).state[4] += e;
+  (*ctx).state[5] += f;
+  (*ctx).state[6] += g;
+  (*ctx).state[7] += h;
+}
+
+
+fn sha256_update(ctx : ptr<function, SHA256_CTX>, len : u32, index: u32)
+{
+  for (var i :u32 = (INPUT_SLOT_SIZE * index); i < ((INPUT_SLOT_SIZE*index)+len); i++) {
+    (*ctx).data[(*ctx).datalen] = inputs[i];
+    (*ctx).datalen++;
+    if ((*ctx).datalen == 64) {
+      sha256_transform(ctx);
+      
+      if ((*ctx).bitlen[0] > 0xffffffff - (512)){
+        (*ctx).bitlen[1]++;
+      }
+      (*ctx).bitlen[0] += 512;
+
+
+      (*ctx).datalen = 0;
+    }
+  }
+}
+
+fn sha256_final(ctx : ptr<function, SHA256_CTX>, hash:  ptr<function, array<u32, 8>> )
+{
+  var i : u32 = (*ctx).datalen;
+
+  if ((*ctx).datalen < 56) {
+    (*ctx).data[i] = 0x80;
+      i++;
+    while (i < 56){
+      (*ctx).data[i] = 0x00;
+      i++;
+    }
+  }
+  else {
+    (*ctx).data[i] = 0x80;
+    i++;
+    while (i < 64){
+      (*ctx).data[i] = 0x00;
+      i++;
+    }
+    sha256_transform(ctx);
+    for (var i = 0; i < 56 ; i++) {
+      (*ctx).data[i] = 0;
+    }
+  }
+
+  if ((*ctx).bitlen[0] > 0xffffffff - (*ctx).datalen * 8) {
+    (*ctx).bitlen[1]++;
+  }
+  (*ctx).bitlen[0] += (*ctx).datalen * 8;
+
+
+  (*ctx).data[63] = (*ctx).bitlen[0];
+  (*ctx).data[62] = (*ctx).bitlen[0] >> 8;
+  (*ctx).data[61] = (*ctx).bitlen[0] >> 16;
+  (*ctx).data[60] = (*ctx).bitlen[0] >> 24;
+  (*ctx).data[59] = (*ctx).bitlen[1];
+  (*ctx).data[58] = (*ctx).bitlen[1] >> 8;
+  (*ctx).data[57] = (*ctx).bitlen[1] >> 16;
+  (*ctx).data[56] = (*ctx).bitlen[1] >> 24;
+  sha256_transform(ctx);
+
+  for (i = 0; i < 8; i++) {
+    (*hash)[i] = (*ctx).state[i]; 
+
+  } 
+  // for (i = 0; i < 4; i++) {
+  //   (*hash)[i] = ((*ctx).state[0] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 4] = ((*ctx).state[1] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 8] = ((*ctx).state[2] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 12] = ((*ctx).state[3] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 16] = ((*ctx).state[4] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 20] = ((*ctx).state[5] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 24] = ((*ctx).state[6] >> (24 - i * 8)) & 0x000000ff;
+  //   (*hash)[i + 28] = ((*ctx).state[7] >> (24 - i * 8)) & 0x000000ff;
+  // }
+}
+
+@compute @workgroup_size(1, 1,1)
+fn main(@builtin(global_invocation_id) global_id : vec3<u32>) {
+  var index = global_id.x;
+  let array_length = arrayLength(&inputSizes);
+  if (index >= array_length) {
+        return;
+  }
+  var ctx : SHA256_CTX;
+  var buf : array<u32, 8>;
+
+  // CTX INIT
+  ctx.datalen = 0;
+  ctx.bitlen[0] = 0;
+  ctx.bitlen[1] = 0;
+  ctx.state[0] = 0x6a09e667;
+  ctx.state[1] = 0xbb67ae85;
+  ctx.state[2] = 0x3c6ef372;
+  ctx.state[3] = 0xa54ff53a;
+  ctx.state[4] = 0x510e527f;
+  ctx.state[5] = 0x9b05688c;
+  ctx.state[6] = 0x1f83d9ab;
+  ctx.state[7] = 0x5be0cd19;
+
+  sha256_update(&ctx, inputSizes[index], index);
+  sha256_final(&ctx, &buf);
+
+
+  for (var i=(index * SHA256_BLOCK_SIZE); i < ((index+1) * SHA256_BLOCK_SIZE); i++) {
+    results[i] = buf[i];
+  }
+}
diff --git a/src/shader_own.wgsl b/src/shader_own.wgsl
new file mode 100644
index 0000000..90e1a84
--- /dev/null
+++ b/src/shader_own.wgsl
@@ -0,0 +1,226 @@
+struct SHA256_CTX {
+  data : array<u32, 64>, 
+  datalen: u32, 
+  bitlen: array<u32, 2>, 
+  state: array<u32, 8>, 
+  info: array<u32, 8>,
+}
+
+const SHA256_BLOCK_SIZE = 32; 
+const INPUT_BLOCK_SIZE  = 64; 
+const INPUT_BLOCK_MAX   = 8; 
+const INPUT_PACKED_SIZE = (INPUT_BLOCK_SIZE * INPUT_BLOCK_MAX / 4); 
+
+struct InputBlocks { 
+  data: array<u32, 128>,
+}
+struct OutputHash {
+  data: array<u32, 8>, 
+} 
+
+@group(0) @binding(0) var<storage, read> inputs: array<InputBlocks>; 
+@group(0) @binding(1) var<storage, read> sizes: array<u32>; 
+@group(0) @binding(2) var<storage, read_write> results: array<OutputHash>; 
+
+
+const k = array<u32, 64> (
+  0x428a2f98,0x71374491,0xb5c0fbcf,0xe9b5dba5,0x3956c25b,0x59f111f1,0x923f82a4,0xab1c5ed5,
+  0xd807aa98,0x12835b01,0x243185be,0x550c7dc3,0x72be5d74,0x80deb1fe,0x9bdc06a7,0xc19bf174,
+  0xe49b69c1,0xefbe4786,0x0fc19dc6,0x240ca1cc,0x2de92c6f,0x4a7484aa,0x5cb0a9dc,0x76f988da,
+  0x983e5152,0xa831c66d,0xb00327c8,0xbf597fc7,0xc6e00bf3,0xd5a79147,0x06ca6351,0x14292967,
+  0x27b70a85,0x2e1b2138,0x4d2c6dfc,0x53380d13,0x650a7354,0x766a0abb,0x81c2c92e,0x92722c85,
+  0xa2bfe8a1,0xa81a664b,0xc24b8b70,0xc76c51a3,0xd192e819,0xd6990624,0xf40e3585,0x106aa070,
+  0x19a4c116,0x1e376c08,0x2748774c,0x34b0bcb5,0x391c0cb3,0x4ed8aa4a,0x5b9cca4f,0x682e6ff3,
+  0x748f82ee,0x78a5636f,0x84c87814,0x8cc70208,0x90befffa,0xa4506ceb,0xbef9a3f7,0xc67178f2
+);
+
+const masks = array<u32, 4> (
+  0x000000FFu, 
+  0x0000FF00u,
+  0x00FF0000u, 
+  0xFF000000u, 
+);
+
+const full_masks = array<u32, 5> (
+  0x00000000u, 
+  0x000000FFu, 
+  0x0000FFFFu,
+  0x00FFFFFFu, 
+  0xFFFFFFFFu, 
+); 
+
+fn ROTLEFT(a : u32, b : u32) -> u32{return (((a) << (b)) | ((a) >> (32-(b))));}
+fn ROTRIGHT(a : u32, b : u32) -> u32{return (((a) >> (b)) | ((a) << (32-(b))));}
+
+fn CH(e : u32, f : u32, g : u32) -> u32{return (((e) & (f)) ^ (~(e) & (g)));}
+fn MAJ(a : u32, b : u32, c: u32) -> u32{return (((a) & (b)) ^ ((a) & (c)) ^ ((b) & (c)));}
+fn EP0(x : u32) -> u32{return (ROTRIGHT(x,2) ^ ROTRIGHT(x,13) ^ ROTRIGHT(x,22));}
+fn EP1(x : u32) -> u32{return (ROTRIGHT(x,6) ^ ROTRIGHT(x,11) ^ ROTRIGHT(x,25));}
+fn SIG0(x : u32) -> u32{return (ROTRIGHT(x,7) ^ ROTRIGHT(x,18) ^ ((x) >> 3));}
+fn SIG1(x : u32) -> u32{return (ROTRIGHT(x,17) ^ ROTRIGHT(x,19) ^ ((x) >> 10));}
+
+fn flip_endian(x: u32) -> u32 {
+  var y: u32 = 0;
+  // 11 22 33 44 => 44 33 22 11
+  // 11 >> 3*8 << 0*8;  
+  // 22 >> 2*8 << 1*8;  
+  // 33 >> 1*8 << 2*8;  
+  // 44 >> 0*8 << 3*8;  
+  for (var i: u32; i < 4; i++) {
+    y |= ((x >> (i*8)) & 0xFF) << ((3-i) * 8); 
+  } 
+  return y; 
+}
+
+fn sha256_init(ctx: ptr<function, SHA256_CTX>) {
+  
+  // CTX INIT
+  (*ctx).datalen = 0;
+  (*ctx).bitlen[0] = 0;
+  (*ctx).bitlen[1] = 0;
+  (*ctx).state[0] = 0x6a09e667;
+  (*ctx).state[1] = 0xbb67ae85;
+  (*ctx).state[2] = 0x3c6ef372;
+  (*ctx).state[3] = 0xa54ff53a;
+  (*ctx).state[4] = 0x510e527f;
+  (*ctx).state[5] = 0x9b05688c;
+  (*ctx).state[6] = 0x1f83d9ab;
+  (*ctx).state[7] = 0x5be0cd19;
+} 
+
+fn sha256_transform(ctx: ptr<function, SHA256_CTX>) {
+  var a : u32 = (*ctx).state[0];
+  var b : u32 = (*ctx).state[1];
+  var c : u32 = (*ctx).state[2];
+  var d : u32 = (*ctx).state[3];
+  var e : u32 = (*ctx).state[4];
+  var f : u32 = (*ctx).state[5];
+  var g : u32 = (*ctx).state[6];
+  var h : u32 = (*ctx).state[7];
+  var t1 : u32;
+  var t2 : u32;
+  var m: array<u32,64>;
+  
+
+  for (var x: u32 = 0;  x < 16; x++) {
+    m[x] = flip_endian((*ctx).data[x]); 
+  }
+  for (var x: u32 = 16; x < 64; x++) {
+    // W_16 = w_0 + sig0(w_1) + w_9 + sig1(w_14)
+    m[x]= m[x-16] + SIG0(m[x-15]) + m[x-7] + SIG1(m[x-2]);
+    // if (x >= 16 && x < 24) {
+    //   var z = x - 16;
+    //   (*ctx).info[z] = m[x];
+    // }
+  } 
+
+  
+  for (var x: u32 = 0; x < 64; x++) {
+    t1 = h + EP1(e) + CH(e, f, g) + k[x] + m[x];  
+    t2 = EP0(a) + MAJ(a,b,c); 
+    h = g; 
+    g = f;
+    f = e;
+    e = d + t1; 
+    d = c; 
+    c = b; 
+    b = a; 
+    a = t1 + t2; 
+   
+     }
+
+  (*ctx).state[0] += a;
+  (*ctx).state[1] += b;
+  (*ctx).state[2] += c;
+  (*ctx).state[3] += d;
+  (*ctx).state[4] += e;
+  (*ctx).state[5] += f;
+  (*ctx).state[6] += g;
+  (*ctx).state[7] += h;
+
+  for (var x: u32 = 0; x < 64; x++) {
+    (*ctx).data[x] = m[x]; 
+  }
+
+
+} 
+
+fn sha256_append_len(ctx: ptr<function, SHA256_CTX>, len: u32) {
+  if ((*ctx).bitlen[0] > 0xffffffff - (len)) {
+    (*ctx).bitlen[1]++; 
+  }
+  (*ctx).bitlen[0] += len; 
+} 
+
+
+fn sha256_update(ctx: ptr<function, SHA256_CTX>, index: u32) {
+  var len = sizes[index];
+  var block_index = 0u; 
+  while len > 64 {
+    for (var i : u32 = 0; i < 16; i++) {
+      (*ctx).data[i] = inputs[index].data[16 * block_index + i];
+    }
+    sha256_transform(ctx);
+    sha256_append_len(ctx, 512);
+    block_index += 1; 
+    len -= 64; 
+  } 
+  var imax = len/4;  
+  if  (len % 4 > 0) {imax += 1;}  
+  for (var i : u32 = 0; i < (imax); i++) {
+      (*ctx).data[i] = inputs[index].data[16 * block_index + i];
+  }    
+  for(var i = imax; i < 16; i ++) {
+        (*ctx).data[i] = 0; 
+  }
+
+  var offset   : u32     = (len) % 4; 
+  var last_word: u32     = (len - offset) / 4; 
+  // if (offset == 0) {last_word += 1;}
+  var padding_word: u32  = (0x80u << (8 * offset));
+  (*ctx).data[last_word] = ((*ctx).data[last_word] & full_masks[offset]) | padding_word; 
+
+  for (var i: u32 = 0; i < 8; i++) {
+     (*ctx).info[i] = (*ctx).data[i];  
+  }
+  for (var i: u32 = 14; i < 16; i++) {
+     (*ctx).info[i] = (*ctx).data[i];  
+  }
+
+  if (len > 56) {
+    sha256_transform(ctx);
+    for(var i = 0; i < 14; i ++) {
+       (*ctx).data[i] = 0; 
+    }
+  }  
+
+  var final_len: u32 = (len*8);
+  sha256_append_len(ctx, final_len);
+  var upper = flip_endian((*ctx).bitlen[0]);  
+  var lower = flip_endian((*ctx).bitlen[1]);
+
+  (*ctx).data[15] = upper; 
+  (*ctx).data[14] = lower; 
+
+  sha256_transform(ctx);
+}
+
+@compute @workgroup_size(64,1,1) 
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+  var ctx: SHA256_CTX; 
+  var index: u32 = global_id.x; 
+  sha256_init(&ctx); 
+  sha256_update(&ctx, index); 
+
+  for(var i: u32 = 0; i < 8; i++) {
+      results[index].data[i] = flip_endian(ctx.state[i]);
+  }
+  // for(var i: u32 = 0; i < 8; i++) {
+  //      results[index].data[i] = inputs[index].data[i];
+  // }
+
+  // results[index].data[0] = ctx.info;
+}
+
+
+