diff --git a/Cargo.lock b/Cargo.lock index 76c7ead9..62e22bc8 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -49,6 +49,21 @@ dependencies = [ "num-traits", ] +[[package]] +name = "arrayvec" +version = "0.7.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8da52d66c7071e2e3fa2a1e5c6d088fec47b593032b254f5e980de8ea54454d6" + +[[package]] +name = "ash" +version = "0.34.0+1.2.203" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b0f780da53d0063880d45554306489f09dd8d1bda47688b4a57bc579119356df" +dependencies = [ + "libloading", +] + [[package]] name = "atty" version = "0.2.14" @@ -85,12 +100,33 @@ dependencies = [ "shlex", ] +[[package]] +name = "bit-set" +version = "0.5.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e11e16035ea35e4e5997b393eacbf6f63983188f7a2ad25bfb13465f5ad59de" +dependencies = [ + "bit-vec", +] + +[[package]] +name = "bit-vec" +version = "0.6.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "349f9b6a179ed607305526ca489b34ad0a41aed5f7980fa90eb03160b69598fb" + [[package]] name = "bitflags" version = "1.3.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" +[[package]] +name = "block" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" + [[package]] name = "bstr" version = "0.2.17" @@ -151,6 +187,12 @@ version = "1.0.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd" +[[package]] +name = "cfg_aliases" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fd16c4719339c4530435d38e511904438d07cce7950afa3718a84ac36c10e89e" + [[package]] name = "clang-sys" version = "1.3.1" @@ -172,16 +214,62 @@ dependencies = [ "unicode-width", ] +[[package]] +name = "codespan-reporting" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" +dependencies = [ + "termcolor", + "unicode-width", +] + +[[package]] +name = "copyless" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2df960f5d869b2dd8532793fde43eb5427cceb126c929747a26823ab0eeb536" + +[[package]] +name = "core-foundation" +version = "0.9.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "194a7a9e6de53fa55116934067c844d9d749312f75c6f6d0980e8c252f8c2146" +dependencies = [ + "core-foundation-sys", + "libc", +] + +[[package]] +name = "core-foundation-sys" +version = "0.8.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5827cebf4670468b8772dd191856768aedcb1b0278a04f989f7766351917b9dc" + +[[package]] +name = "core-graphics-types" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b" +dependencies = [ + "bitflags", + "core-foundation", + "foreign-types", + "libc", +] + [[package]] name = "coupe" version = "0.1.0" dependencies = [ "approx", "criterion", + "futures-lite", "gnuplot", "itertools", "nalgebra", "num", + "once_cell", "proptest", "rand", "rayon", @@ -190,6 +278,7 @@ dependencies = [ "tracing-chrome", "tracing-subscriber", "tracing-tree", + "wgpu", ] [[package]] @@ -328,6 +417,23 @@ dependencies = [ "memchr", ] +[[package]] +name = "cty" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b365fabc795046672053e29c954733ec3b05e4be654ab130fe8f1f94d7051f35" + +[[package]] +name = "d3d12" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2daefd788d1e96e0a9d66dee4b828b883509bc3ea9ce30665f04c3246372690c" +dependencies = [ + "bitflags", + "libloading", + "winapi", +] + [[package]] name = "either" version = "1.6.1" @@ -367,12 +473,72 @@ version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7360491ce676a36bf9bb3c56c1aa791658183a54d2744120f27285738d90465a" +[[package]] +name = "fastrand" +version = "1.7.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c3fcf0cee53519c866c09b5de1f6c56ff9d647101f81c1964fa632e148896cdf" +dependencies = [ + "instant", +] + [[package]] name = "fnv" version = "1.0.7" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "3f9eec918d3f24069decb9af1554cad7c880e2da24a9afd88aca000531ab82c1" +[[package]] +name = "foreign-types" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f6f339eb8adc052cd2ca78910fda869aefa38d22d5cb648e6485e4d3fc06f3b1" +dependencies = [ + "foreign-types-shared", +] + +[[package]] +name = "foreign-types-shared" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" + +[[package]] +name = "futures-core" +version = "0.3.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0c09fd04b7e4073ac7156a9539b57a484a8ea920f79c7c675d05d289ab6110d3" + +[[package]] +name = "futures-io" +version = "0.3.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fc4045962a5a5e935ee2fdedaa4e08284547402885ab326734432bed5d12966b" + +[[package]] +name = "futures-lite" +version = "1.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7694489acd39452c77daa48516b894c153f192c3578d5a839b62c58099fcbf48" +dependencies = [ + "fastrand", + "futures-core", + "futures-io", + "memchr", + "parking", + "pin-project-lite", + "waker-fn", +] + +[[package]] +name = "fxhash" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c" +dependencies = [ + "byteorder", +] + [[package]] name = "getopts" version = "0.2.21" @@ -399,6 +565,18 @@ version = "0.3.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "9b919933a397b79c37e33b77bb2aa3dc8eb6e165ad809e58ff75bc7db2e34574" +[[package]] +name = "glow" +version = "0.11.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d8bd5877156a19b8ac83a29b2306fe20537429d318f3ff0a1a2119f8d9c61919" +dependencies = [ + "js-sys", + "slotmap", + "wasm-bindgen", + "web-sys", +] + [[package]] name = "gnuplot" version = "0.0.35" @@ -408,6 +586,45 @@ dependencies = [ "byteorder", ] +[[package]] +name = "gpu-alloc" +version = "0.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7fc59e5f710e310e76e6707f86c561dd646f69a8876da9131703b2f717de818d" +dependencies = [ + "bitflags", + "gpu-alloc-types", +] + +[[package]] +name = "gpu-alloc-types" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "54804d0d6bc9d7f26db4eaec1ad10def69b599315f487d32c334a80d1efe67a5" +dependencies = [ + "bitflags", +] + +[[package]] +name = "gpu-descriptor" +version = "0.2.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a538f217be4d405ff4719a283ca68323cc2384003eca5baaa87501e821c81dda" +dependencies = [ + "bitflags", + "gpu-descriptor-types", + "hashbrown", +] + +[[package]] +name = "gpu-descriptor-types" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "363e3677e55ad168fef68cf9de3a4a310b53124c5e784c53a1d70e92d23f2126" +dependencies = [ + "bitflags", +] + [[package]] name = "half" version = "1.7.1" @@ -441,6 +658,40 @@ dependencies = [ "libc", ] +[[package]] +name = "hexf-parse" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" + +[[package]] +name = "indexmap" +version = "1.8.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e6012d540c5baa3589337a98ce73408de9b5a25ec9fc2c6fd6be8f0d39e0ca5a" +dependencies = [ + "autocfg", + "hashbrown", +] + +[[package]] +name = "inplace_it" +version = "0.3.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "90953f308a79fe6d62a4643e51f848fbfddcd05975a38e69fdf4ab86a7baf7ca" + +[[package]] +name = "instant" +version = "0.1.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c" +dependencies = [ + "cfg-if", + "js-sys", + "wasm-bindgen", + "web-sys", +] + [[package]] name = "itertools" version = "0.10.1" @@ -471,6 +722,16 @@ version = "0.12.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "078e285eafdfb6c4b434e0d31e8cfcb5115b651496faca5749b88fafd4f23bfd" +[[package]] +name = "khronos-egl" +version = "4.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8c2352bd1d0bceb871cb9d40f24360c8133c11d7486b68b5381c1dd1a32015e3" +dependencies = [ + "libc", + "libloading", +] + [[package]] name = "lazy_static" version = "1.4.0" @@ -489,6 +750,16 @@ version = "0.2.103" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "dd8f7255a17a627354f321ef0055d63b898c6fb27eff628af4d1b66b7331edf6" +[[package]] +name = "libloading" +version = "0.7.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "efbc0f03f9a775e9f6aed295c6a1ba2253c5757a9e03d55c6caa46a681abcddd" +dependencies = [ + "cfg-if", + "winapi", +] + [[package]] name = "libm" version = "0.2.1" @@ -505,6 +776,15 @@ dependencies = [ "vcpkg", ] +[[package]] +name = "lock_api" +version = "0.4.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "88943dd7ef4a2e5a4bfa2753aaab3013e34ce2533d1996fb18ef591e315e2b3b" +dependencies = [ + "scopeguard", +] + [[package]] name = "log" version = "0.4.14" @@ -514,6 +794,15 @@ dependencies = [ "cfg-if", ] +[[package]] +name = "malloc_buf" +version = "0.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb" +dependencies = [ + "libc", +] + [[package]] name = "matchers" version = "0.1.0" @@ -562,6 +851,20 @@ dependencies = [ "mesh-io", ] +[[package]] +name = "metal" +version = "0.23.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e0514f491f4cc03632ab399ee01e2c1c1b12d3e1cf2d667c1ff5f87d6dcd2084" +dependencies = [ + "bitflags", + "block", + "core-graphics-types", + "foreign-types", + "log", + "objc", +] + [[package]] name = "metis" version = "0.1.1" @@ -580,6 +883,24 @@ dependencies = [ "bindgen", ] +[[package]] +name = "naga" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3012f2dbcc79e8e0b5825a4836a7106a75dd9b2fe42c528163be0f572538c705" +dependencies = [ + "bit-set", + "bitflags", + "codespan-reporting", + "hexf-parse", + "indexmap", + "log", + "num-traits", + "rustc-hash", + "spirv", + "thiserror", +] + [[package]] name = "nalgebra" version = "0.29.0" @@ -721,6 +1042,25 @@ dependencies = [ "libc", ] +[[package]] +name = "objc" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" +dependencies = [ + "malloc_buf", + "objc_exception", +] + +[[package]] +name = "objc_exception" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad970fb455818ad6cba4c122ad012fae53ae8b4795f86378bce65e4f6bab2ca4" +dependencies = [ + "cc", +] + [[package]] name = "once_cell" version = "1.10.0" @@ -733,6 +1073,37 @@ version = "11.1.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0ab1bc2a289d34bd04a330323ac98a1b4bc82c9d9fcb1e66b63caa84da26b575" +[[package]] +name = "parking" +version = "2.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "427c3892f9e783d91cc128285287e70a59e206ca452770ece88a76f7a3eddd72" + +[[package]] +name = "parking_lot" +version = "0.11.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99" +dependencies = [ + "instant", + "lock_api", + "parking_lot_core", +] + +[[package]] +name = "parking_lot_core" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d76e8e1493bcac0d2766c42737f34458f1c8c50c0d23bcb24ea953affb273216" +dependencies = [ + "cfg-if", + "instant", + "libc", + "redox_syscall", + "smallvec", + "winapi", +] + [[package]] name = "paste" version = "1.0.5" @@ -800,6 +1171,12 @@ dependencies = [ "unicode-xid", ] +[[package]] +name = "profiling" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2f61dcf0b917cd75d4521d7343d1ffff3d1583054133c9b5cbea3375c703c40d" + [[package]] name = "proptest" version = "1.0.0" @@ -908,6 +1285,21 @@ dependencies = [ "rand_core", ] +[[package]] +name = "range-alloc" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "63e935c45e09cc6dcf00d2f0b2d630a58f4095320223d47fc68918722f0538b6" + +[[package]] +name = "raw-window-handle" +version = "0.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b800beb9b6e7d2df1fe337c9e3d04e3af22a124460fb4c30fcc22c9117cefb41" +dependencies = [ + "cty", +] + [[package]] name = "rawpointer" version = "0.2.1" @@ -981,6 +1373,12 @@ dependencies = [ "winapi", ] +[[package]] +name = "renderdoc-sys" +version = "0.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f1382d1f0a252c4bf97dc20d979a2fdd05b024acd7c2ed0f7595d7817666a157" + [[package]] name = "rusqlite" version = "0.26.3" @@ -1143,12 +1541,31 @@ dependencies = [ "wide", ] +[[package]] +name = "slotmap" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e1e08e261d0e8f5c43123b7adf3e4ca1690d655377ac93a03b2c9d3e98de1342" +dependencies = [ + "version_check", +] + [[package]] name = "smallvec" version = "1.6.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "fe0f37c9e8f3c5a4a66ad655a93c74daac4ad00c441533bf5c6e7990bb42604e" +[[package]] +name = "spirv" +version = "0.2.0+1.5.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "246bfa38fe3db3f1dfc8ca5a2cdeb7348c78be2112740cc0ec8ef18b6d94f830" +dependencies = [ + "bitflags", + "num-traits", +] + [[package]] name = "sprs" version = "0.11.0" @@ -1188,6 +1605,15 @@ dependencies = [ "winapi", ] +[[package]] +name = "termcolor" +version = "1.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bab24d30b911b2376f3a13cc2cd443142f0c81dda04c118693e35b3835757755" +dependencies = [ + "winapi-util", +] + [[package]] name = "textwrap" version = "0.11.0" @@ -1197,6 +1623,26 @@ dependencies = [ "unicode-width", ] +[[package]] +name = "thiserror" +version = "1.0.31" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bd829fe32373d27f76265620b5309d0340cb8550f523c1dda251d6298069069a" +dependencies = [ + "thiserror-impl", +] + +[[package]] +name = "thiserror-impl" +version = "1.0.31" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0396bc89e626244658bef819e22d0cc459e795a5ebe878e6ec336d1674a8d79a" +dependencies = [ + "proc-macro2", + "quote", + "syn", +] + [[package]] name = "thread_local" version = "1.1.4" @@ -1347,6 +1793,12 @@ dependencies = [ "libc", ] +[[package]] +name = "waker-fn" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9d5b2c62b4012a3e1eca5a7e077d13b3bf498c4073e33ccd58626607748ceeca" + [[package]] name = "walkdir" version = "2.3.2" @@ -1389,6 +1841,18 @@ dependencies = [ "wasm-bindgen-shared", ] +[[package]] +name = "wasm-bindgen-futures" +version = "0.4.28" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8e8d7523cb1f2a4c96c1317ca690031b714a51cc14e05f712446691f413f5d39" +dependencies = [ + "cfg-if", + "js-sys", + "wasm-bindgen", + "web-sys", +] + [[package]] name = "wasm-bindgen-macro" version = "0.2.78" @@ -1428,6 +1892,97 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "wgpu" +version = "0.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b97cd781ff044d6d697b632a2e212032c2e957d1afaa21dbf58069cbb8f78567" +dependencies = [ + "arrayvec", + "js-sys", + "log", + "naga", + "parking_lot", + "raw-window-handle", + "smallvec", + "wasm-bindgen", + "wasm-bindgen-futures", + "web-sys", + "wgpu-core", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-core" +version = "0.12.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c4688c000eb841ca55f7b35db659b78d6e1cd77d7caf8fb929f4e181f754047d" +dependencies = [ + "arrayvec", + "bitflags", + "cfg_aliases", + "codespan-reporting", + "copyless", + "fxhash", + "log", + "naga", + "parking_lot", + "profiling", + "raw-window-handle", + "smallvec", + "thiserror", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-hal" +version = "0.12.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d684ea6a34974a2fc19f1dfd183d11a62e22d75c4f187a574bb1224df8e056c2" +dependencies = [ + "arrayvec", + "ash", + "bit-set", + "bitflags", + "block", + "core-graphics-types", + "d3d12", + "foreign-types", + "fxhash", + "glow", + "gpu-alloc", + "gpu-descriptor", + "inplace_it", + "js-sys", + "khronos-egl", + "libloading", + "log", + "metal", + "naga", + "objc", + "parking_lot", + "profiling", + "range-alloc", + "raw-window-handle", + "renderdoc-sys", + "thiserror", + "wasm-bindgen", + "web-sys", + "wgpu-types", + "winapi", +] + +[[package]] +name = "wgpu-types" +version = "0.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "549533d9e1cdd4b4cda7718d33ff500fc4c34b5467b71d76b547ae0324f3b2a2" +dependencies = [ + "bitflags", +] + [[package]] name = "wide" version = "0.7.1" diff --git a/Cargo.toml b/Cargo.toml index b405198b..668b2fb1 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -35,6 +35,9 @@ itertools = "0.10" nalgebra = { version = "0.29", default-features = false, features = ["rand", "std"] } num = "0.4" rayon = "1.3" +futures-lite = "1.12" +wgpu = "0.12" +once_cell = "1.10" tracing = { version = "0.1", default-features = false, features = ["std"] } rand = "0.8" sprs = { version = "0.11", default-features = false, features = ["multi_thread"] } diff --git a/src/algorithms/recursive_bisection.wgsl b/src/algorithms/recursive_bisection.wgsl new file mode 100644 index 00000000..7eacb9bf --- /dev/null +++ b/src/algorithms/recursive_bisection.wgsl @@ -0,0 +1,35 @@ +struct Numbers { + data: [[stride(4)]] array; +}; + +[[group(0), binding(0)]] +var numbers: Numbers; + +[[override]] let blockSize: u32; +var workgroup_data: array; + +[[stage(compute), workgroup_size(16)]] +fn main( + [[builtin(global_invocation_id)]] global_id: vec3, + [[builtin(local_invocation_id)]] local_id: vec3, + [[builtin(workgroup_id)]] workgroup_id: vec3, +) { + var n: u32 = arrayLength(&numbers.data); + + if (global_id.x < n) { + workgroup_data[local_id.x] = numbers.data[global_id.x]; + } else { + workgroup_data[local_id.x] = 0u; + } + + workgroupBarrier(); + + for (var stride: u32 = blockSize / 2u; stride > 0u; stride = stride >> 1u) { + var flag: u32 = u32(local_id.x < stride); + workgroup_data[local_id.x] = workgroup_data[local_id.x] + flag * workgroup_data[local_id.x + flag * stride]; + } + + if (local_id.x == 0u) { + numbers.data[workgroup_id.x] = workgroup_data[0]; + } +} diff --git a/src/gpu.rs b/src/gpu.rs new file mode 100644 index 00000000..809319e0 --- /dev/null +++ b/src/gpu.rs @@ -0,0 +1,46 @@ +use once_cell::sync::Lazy; + +#[derive(Debug, Copy, Clone)] +pub enum InitError { + NoAdapter, + CannotRequestDevice, +} + +struct Context { + instance: wgpu::Instance, + adapter: wgpu::Adapter, + device: wgpu::Device, + queue: wgpu::Queue, +} + +async fn init_context() -> Result { + let instance = wgpu::Instance::new(wgpu::Backends::all()); + let adapter = instance + .request_adapter(&wgpu::RequestAdapterOptions { + power_preference: wgpu::PowerPreference::HighPerformance, + force_fallback_adapter: false, + compatible_surface: None, + }) + .await + .ok_or(InitError::NoAdapter)?; + let (device, queue) = adapter + .request_device(&wgpu::DeviceDescriptor::default(), None) + .await + .map_err(|_| InitError::CannotRequestDevice)?; + Ok(Context { + instance, + adapter, + device, + queue, + }) +} + +static CONTEXT: Lazy> = + Lazy::new(|| futures_lite::future::block_on(init_context())); + +pub fn ensure_available() -> Result<(), InitError> { + match &*CONTEXT { + Ok(_) => Ok(()), + Err(err) => Err(*err), + } +} diff --git a/src/lib.rs b/src/lib.rs index 1bf28ee5..17b73ab8 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -37,6 +37,7 @@ mod algorithms; mod geometry; +pub mod gpu; pub mod imbalance; mod nextafter; mod real; diff --git a/src/main.rs b/src/main.rs new file mode 100644 index 00000000..9de09cc1 --- /dev/null +++ b/src/main.rs @@ -0,0 +1,116 @@ +use std::mem; + +use wgpu::util::DeviceExt as _; + +fn slice_as_bytes(v: &[T]) -> &[u8] { + unsafe { std::slice::from_raw_parts(v.as_ptr() as *const u8, v.len() * mem::size_of::()) } +} + +fn bytes_as_slice(v: &[u8]) -> &[T] { + assert_eq!(v.len() % mem::size_of::(), 0); + unsafe { std::slice::from_raw_parts(v.as_ptr() as *const T, v.len() / mem::size_of::()) } +} + +async fn steps_many(numbers: &[u32]) -> u32 { + let instance = wgpu::Instance::new(wgpu::Backends::all()); + let adapter = instance + .request_adapter(&wgpu::RequestAdapterOptions { + power_preference: wgpu::PowerPreference::HighPerformance, + force_fallback_adapter: false, + compatible_surface: None, + }) + .await + .unwrap(); + let (device, queue) = adapter + .request_device(&wgpu::DeviceDescriptor::default(), None) + .await + .unwrap(); + + let number_buf = slice_as_bytes(numbers); + let number_buf_size = number_buf.len() as wgpu::BufferAddress; + let staging_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("my staging buffer"), + size: number_buf_size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + let storage_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: Some("my storage buffer"), + contents: number_buf, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_SRC + | wgpu::BufferUsages::COPY_DST, + }); + + let cs_module = + device.create_shader_module(&wgpu::include_wgsl!("algorithms/recursive_bisection.wgsl")); + let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("my compute pipeline"), + layout: None, + module: &cs_module, + entry_point: "main", + }); + + let bind_group_layout = compute_pipeline.get_bind_group_layout(0); + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("my bindgroup"), + layout: &bind_group_layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: storage_buffer.as_entire_binding(), + }], + }); + + const BLOCK_SIZE: u32 = 16; + let mut n = numbers.len() as u32; + while n >= BLOCK_SIZE { + println!("Iter: {n}"); + + let mut command_encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("my command encoder"), + }); + { + let mut compute_pass = + command_encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("my compute pass"), + }); + compute_pass.set_pipeline(&compute_pipeline); + compute_pass.set_bind_group(0, &bind_group, &[]); + compute_pass.insert_debug_marker("compute this mn yeah"); + compute_pass.dispatch(n, 1, 1); + } + queue.submit(Some(command_encoder.finish())); + + device.poll(wgpu::Maintain::Wait); + n /= BLOCK_SIZE; + } + + let mut command_encoder = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("my command encoder"), + }); + command_encoder.copy_buffer_to_buffer(&storage_buffer, 0, &staging_buffer, 0, n as u64 * 4); + queue.submit(Some(command_encoder.finish())); + + let buffer_slice = staging_buffer.slice(..n as u64 * 4); + let buffer_future = buffer_slice.map_async(wgpu::MapMode::Read); + + device.poll(wgpu::Maintain::Wait); + + if buffer_future.await.is_err() { + panic!("oh noes..."); + } + + let data = buffer_slice.get_mapped_range(); + let result: u32 = bytes_as_slice::(&data).iter().sum(); + mem::drop(data); + staging_buffer.unmap(); + + result +} + +fn main() { + let numbers: Vec = (1..=128).collect(); + let sum = futures_lite::future::block_on(steps_many(&numbers)); + + println!("{sum}"); +}