diff --git a/Cargo.lock b/Cargo.lock index e5b2eaa..aed3d6f 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -4,9 +4,9 @@ version = 3 [[package]] name = "ab_glyph_rasterizer" -version = "0.1.4" +version = "0.1.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d9fe5e32de01730eb1f6b7f5b51c17e03e2325bf40a74f754f04f130043affff" +checksum = "a13739d7177fbd22bb0ed28badfff9f372f8bef46c863db4e1c6248f6b223b6e" [[package]] name = "adler32" @@ -33,7 +33,7 @@ version = "0.11.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ee49baf6cb617b853aa8d93bf420db2383fab46d314482ca2803b40d5fde979b" dependencies = [ - "winapi 0.3.9", + "winapi", ] [[package]] @@ -70,14 +70,20 @@ checksum = "d9b39be18770d11421cdb1b9947a45dd3f37e93092cbf377614828a319d5fee8" dependencies = [ "hermit-abi", "libc", - "winapi 0.3.9", + "winapi", ] [[package]] -name = "bitflags" -version = "1.2.1" +name = "autocfg" +version = "1.0.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cf1de2fe8c75bc145a2f577add951f8134889b4795d47466a54a5c846d691693" +checksum = "cdb031dd78e28731d87d56cc8ffef4a8f36ca26c38fe2de700543e627f8a464a" + +[[package]] +name = "bitflags" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" [[package]] name = "block" @@ -93,9 +99,9 @@ checksum = "72957246c41db82b8ef88a5486143830adeb8227ef9837740bdec67724cf2c5b" [[package]] name = "byteorder" -version = "1.3.4" +version = "1.4.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "08c48aae112d48ed9f069b33538ea9e3e90aa263cfa3d1c24309612b1f7472de" +checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610" [[package]] name = "calloop" @@ -104,14 +110,14 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0b036167e76041694579972c28cf4877b4f92da222560ddb49008937b6a6727c" dependencies = [ "log", - "nix", + "nix 0.18.0", ] [[package]] name = "cc" -version = "1.0.62" +version = "1.0.72" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f1770ced377336a88a67c473594ccc14eca6f4559217c34f64aac8f83d641b40" +checksum = "22a9137b95ea06864e018375b72adfb7db6e6f68cfc8df5a04d00288050485ee" [[package]] name = "cfg-if" @@ -140,41 +146,17 @@ dependencies = [ "vec_map", ] -[[package]] -name = "cloudabi" -version = "0.1.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4344512281c643ae7638bbabc3af17a11307803ec8f0fcad9fae512a8bf36467" -dependencies = [ - "bitflags", -] - [[package]] name = "cocoa" -version = "0.20.2" +version = "0.24.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0c49e86fc36d5704151f5996b7b3795385f50ce09e3be0f47a0cfde869681cf8" -dependencies = [ - "bitflags", - "block", - "core-foundation 0.7.0", - "core-graphics 0.19.2", - "foreign-types", - "libc", - "objc", -] - -[[package]] -name = "cocoa" -version = "0.23.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c54201c07dcf3a5ca33fececb8042aed767ee4bfd5a0235a8ceabcda956044b2" +checksum = "6f63902e9223530efb4e26ccd0cf55ec30d592d3b42e21a28defc42a9586e832" dependencies = [ "bitflags", "block", "cocoa-foundation", - "core-foundation 0.9.1", - "core-graphics 0.22.1", + "core-foundation 0.9.2", + "core-graphics 0.22.3", "foreign-types", "libc", "objc", @@ -188,7 +170,7 @@ checksum = "7ade49b65d560ca58c403a479bb396592b155c0185eada742ee323d1d68d6318" dependencies = [ "bitflags", "block", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "core-graphics-types", "foreign-types", "libc", @@ -207,11 +189,11 @@ dependencies = [ [[package]] name = "core-foundation" -version = "0.9.1" +version = "0.9.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0a89e2ae426ea83155dccf10c0fa6b1463ef6d5fcb44cee0b224a408fa640a62" +checksum = "6888e10551bb93e424d8df1d07f1a8b4fceb0001a3a4b048bfc47554946f47b3" dependencies = [ - "core-foundation-sys 0.8.2", + "core-foundation-sys 0.8.3", "libc", ] @@ -223,9 +205,9 @@ checksum = "b3a71ab494c0b5b860bdc8407ae08978052417070c2ced38573a9157ad75b8ac" [[package]] name = "core-foundation-sys" -version = "0.8.2" +version = "0.8.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ea221b5284a47e40033bf9b66f35f984ec0ea2931eb03505246cd27a963f981b" +checksum = "5827cebf4670468b8772dd191856768aedcb1b0278a04f989f7766351917b9dc" [[package]] name = "core-graphics" @@ -241,12 +223,12 @@ dependencies = [ [[package]] name = "core-graphics" -version = "0.22.1" +version = "0.22.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fc239bba52bab96649441699533a68de294a101533b0270b2d65aa402b29a7f9" +checksum = "2581bbab3b8ffc6fcbd550bf46c355135d16e9ff2a6ea032ad6b9bf1d7efe4fb" dependencies = [ "bitflags", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "core-graphics-types", "foreign-types", "libc", @@ -259,7 +241,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "3a68b68b3446082644c91ac778bf50cd4104bfb002b5a6a7c44cca5a2c70788b" dependencies = [ "bitflags", - "core-foundation 0.9.1", + "core-foundation 0.9.2", "foreign-types", "libc", ] @@ -286,6 +268,74 @@ dependencies = [ "cfg-if 1.0.0", ] +[[package]] +name = "crossbeam" +version = "0.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4ae5588f6b3c3cb05239e90bd110f257254aecd01e4635400391aeae07497845" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-channel", + "crossbeam-deque", + "crossbeam-epoch", + "crossbeam-queue", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-channel" +version = "0.5.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06ed27e177f16d65f0f0c22a213e17c696ace5dd64b14258b52f9417ccb52db4" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-deque" +version = "0.8.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6455c0ca19f0d2fbf751b908d5c55c1f5cbc65e03c4225427254b46890bdde1e" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-epoch", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-epoch" +version = "0.9.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4ec02e091aa634e2c3ada4a392989e7c3116673ef0ac5b72232439094d73b7fd" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", + "lazy_static", + "memoffset", + "scopeguard", +] + +[[package]] +name = "crossbeam-queue" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9b10ddc024425c88c2ad148c1b0fd53f4c6d38db9697c9f1588381212fa657c9" +dependencies = [ + "cfg-if 1.0.0", + "crossbeam-utils", +] + +[[package]] +name = "crossbeam-utils" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d82cfc11ce7f2c3faef78d8a684447b40d503d9681acebed6cb728d45940c4db" +dependencies = [ + "cfg-if 1.0.0", + "lazy_static", +] + [[package]] name = "darling" version = "0.10.2" @@ -333,15 +383,35 @@ dependencies = [ [[package]] name = "derivative" -version = "2.1.1" +version = "2.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cb582b60359da160a9477ee80f15c8d784c477e69c217ef2cdd4169c24ea380f" +checksum = "fcc3dd5e9e9c0b295d6e1e4d811fb6f157d5ffd784b8d202fc62eac8035a770b" dependencies = [ "proc-macro2", "quote", "syn", ] +[[package]] +name = "dirs" +version = "3.0.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "30baa043103c9d0c2a57cf537cc2f35623889dc0d405e6c3cccfadbc81c71309" +dependencies = [ + "dirs-sys", +] + +[[package]] +name = "dirs-sys" +version = "0.3.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "03d86534ed367a67548dc68113a0f5db55432fdfbb6e6f9d77704397d95d5780" +dependencies = [ + "libc", + "redox_users", + "winapi", +] + [[package]] name = "dispatch" version = "0.2.0" @@ -354,7 +424,16 @@ version = "0.4.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b11f15d1e3268f140f68d390637d5e76d849782d971ae7063e0da69fe9709a76" dependencies = [ - "libloading 0.6.5", + "libloading 0.6.7", +] + +[[package]] +name = "dlib" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ac1b7517328c04c2aa68422fc60a41b92208182142ed04a25879c26c8f878794" +dependencies = [ + "libloading 0.7.1", ] [[package]] @@ -385,43 +464,38 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" [[package]] -name = "fuchsia-zircon" -version = "0.3.3" +name = "getrandom" +version = "0.1.16" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2e9763c69ebaae630ba35f74888db465e49e259ba1bc0eda7d06f4a067615d82" +checksum = "8fc3cb4d91f53b50155bdcfd23f6a4c39ae1969c2ae85982b135750cccaf5fce" dependencies = [ - "bitflags", - "fuchsia-zircon-sys", + "cfg-if 1.0.0", + "libc", + "wasi 0.9.0+wasi-snapshot-preview1", ] -[[package]] -name = "fuchsia-zircon-sys" -version = "0.3.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3dcaa9ae7725d12cdb85b3ad99a434db70b468c09ded17e012d86b5c1010f7a7" - [[package]] name = "getrandom" -version = "0.1.15" +version = "0.2.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fc587bc0ec293155d5bfa6b9891ec18a1e330c234f896ea47fbada4cadbe47e6" +checksum = "7fcd999463524c52659517fe2cea98493cfe485d10565e7b0fb07dbba7ad2753" dependencies = [ - "cfg-if 0.1.10", + "cfg-if 1.0.0", "libc", - "wasi", + "wasi 0.10.2+wasi-snapshot-preview1", ] [[package]] name = "half" -version = "1.6.0" +version = "1.8.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d36fab90f82edc3c747f9d438e06cf0a491055896f2a279638bb5beed6c40177" +checksum = "eabb4a44450da02c90444cf74558da904edde8fb4e9035a9a6a4e15445af0bd7" [[package]] name = "hermit-abi" -version = "0.1.17" +version = "0.1.19" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5aca5565f760fb5b220e499d72710ed156fdb74e631659e99377d9ebfbd13ae8" +checksum = "62b467343b94ba476dcb2500d242dadbb39557df889310ac77c5d99100aaac33" dependencies = [ "libc", ] @@ -434,38 +508,19 @@ checksum = "b9e0384b61958566e926dc50660321d12159025e767c18e043daf26b70104c39" [[package]] name = "instant" -version = "0.1.8" +version = "0.1.12" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cb1fc4429a33e1f80d41dc9fea4d108a88bec1de8053878898ae448a0b52f613" +checksum = "7a5bbe824c507c5da5956355e86a746d82e0e1464f65d862cc5e71da70e94b2c" dependencies = [ "cfg-if 1.0.0", ] -[[package]] -name = "iovec" -version = "0.1.4" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b2b3ea6ff95e175473f8ffe6a7eb7c00d054240321b84c57051175fe3c1e075e" -dependencies = [ - "libc", -] - [[package]] name = "jni-sys" version = "0.3.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" -[[package]] -name = "kernel32-sys" -version = "0.2.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7507624b29483431c0ba2d82aece8ca6cdba9382bff4ddd0f7490560c056098d" -dependencies = [ - "winapi 0.2.8", - "winapi-build", -] - [[package]] name = "kurbo" version = "0.7.1" @@ -481,26 +536,20 @@ version = "1.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e2abad23fbc42b3700f2f279844dc832adb2b2eb069b2df918f455c4e18cc646" -[[package]] -name = "lazycell" -version = "1.3.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "830d08ce1d1d941e6b30645f1a0eb5643013d835ce3779a5fc208261dbe10f55" - [[package]] name = "libc" -version = "0.2.80" +version = "0.2.107" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4d58d1b70b004888f764dfbf6a26a3b0342a1632d33968e4a179d8011c760614" +checksum = "fbe5e23404da5b4f555ef85ebed98fb4083e55a00c317800bc2a50ede9f3d219" [[package]] name = "libloading" -version = "0.6.5" +version = "0.6.7" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1090080fe06ec2648d0da3881d9453d97e71a45f00eb179af7fdd7e3f686fdb0" +checksum = "351a32417a12d5f7e82c368a66781e307834dae04c6ce0cd4456d52989229883" dependencies = [ "cfg-if 1.0.0", - "winapi 0.3.9", + "winapi", ] [[package]] @@ -510,25 +559,25 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c0cf036d15402bea3c5d4de17b3fce76b3e4a56ebc1f577be0e7a72f7c607cf0" dependencies = [ "cfg-if 1.0.0", - "winapi 0.3.9", + "winapi", ] [[package]] name = "lock_api" -version = "0.4.1" +version = "0.4.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "28247cc5a5be2f05fbcd76dd0cf2c7d3b5400cb978a28042abcd4fa0b3f8261c" +checksum = "712a4d093c9976e24e7dbca41db895dabcbac38eb5f4045393d17a95bdfb1109" dependencies = [ "scopeguard", ] [[package]] name = "log" -version = "0.4.11" +version = "0.4.14" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "4fabed175da42fed1fa0746b0ea71f412aa9d35e76e95e59b192c64b9dc2bf8b" +checksum = "51b9bbe6c47d51fc3e1a9b945965946b4c44142ab8792c50835a980d362c2710" dependencies = [ - "cfg-if 0.1.10", + "cfg-if 1.0.0", ] [[package]] @@ -542,30 +591,32 @@ dependencies = [ [[package]] name = "matches" -version = "0.1.8" +version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7ffc5c5338469d4d3ea17d269fa8ea3512ad247247c30bd2df69e68309ed0a08" - -[[package]] -name = "maybe-uninit" -version = "2.0.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "60302e4db3a61da70c0cb7991976248362f30319e88850c487b9b95bbf059e00" +checksum = "a3e378b66a060d48947b590737b30a1be76706c8dd7b8ba0f2fe3989c68a853f" [[package]] name = "memchr" -version = "2.3.4" +version = "2.4.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0ee1c47aaa256ecabcaea351eae4a9b01ef39ed810004e298d2511ed284b1525" +checksum = "308cc39be01b73d0d18f82a0e7b2a3df85245f84af96fdddc5d202d27e47b86a" [[package]] -name = "memmap" -version = "0.7.0" +name = "memmap2" +version = "0.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6585fd95e7bb50d6cc31e20d4cf9afb4e2ba16c5846fc76793f11218da9c475b" +checksum = "d9b70ca2a6103ac8b665dc150b142ef0e4e89df640c9e6cf295d189c3caebe5a" dependencies = [ "libc", - "winapi 0.3.9", +] + +[[package]] +name = "memoffset" +version = "0.6.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "59accc507f1338036a0477ef61afdae33cde60840f4dfe481319ce3ad116ddf9" +dependencies = [ + "autocfg", ] [[package]] @@ -582,6 +633,12 @@ dependencies = [ "objc", ] +[[package]] +name = "minimal-lexical" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a" + [[package]] name = "miniz_oxide" version = "0.3.7" @@ -593,57 +650,36 @@ dependencies = [ [[package]] name = "mio" -version = "0.6.22" +version = "0.7.14" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fce347092656428bc8eaf6201042cb551b8d67855af7374542a92a0fbfcac430" +checksum = "8067b404fe97c70829f082dec8bcf4f71225d7eaea1d8645349cb76fa06205cc" dependencies = [ - "cfg-if 0.1.10", - "fuchsia-zircon", - "fuchsia-zircon-sys", - "iovec", - "kernel32-sys", "libc", "log", "miow", - "net2", - "slab", - "winapi 0.2.8", + "ntapi", + "winapi", ] [[package]] -name = "mio-extras" -version = "2.0.6" +name = "mio-misc" +version = "1.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "52403fe290012ce777c4626790c8951324a2b9e3316b3143779c72b029742f19" +checksum = "0ddf05411bb159cdb5801bb10002afb66cb4572be656044315e363460ce69dc2" dependencies = [ - "lazycell", + "crossbeam", + "crossbeam-queue", "log", "mio", - "slab", ] [[package]] name = "miow" -version = "0.2.1" +version = "0.3.7" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "8c1f2f3b1cf331de6896aabf6e9d55dca90356cc9960cca7eaaf408a355ae919" +checksum = "b9f1c5b025cda876f66ef43a113f91ebc9f4ccef34843000e0adf6ebbab84e21" dependencies = [ - "kernel32-sys", - "net2", - "winapi 0.2.8", - "ws2_32-sys", -] - -[[package]] -name = "ndk" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5eb167c1febed0a496639034d0c76b3b74263636045db5489eee52143c246e73" -dependencies = [ - "jni-sys", - "ndk-sys", - "num_enum 0.4.3", - "thiserror", + "winapi", ] [[package]] @@ -654,24 +690,10 @@ checksum = "8794322172319b972f528bf90c6b467be0079f1fa82780ffb431088e741a73ab" dependencies = [ "jni-sys", "ndk-sys", - "num_enum 0.5.1", + "num_enum", "thiserror", ] -[[package]] -name = "ndk-glue" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "bdf399b8b7a39c6fb153c4ec32c72fd5fe789df24a647f229c239aa7adb15241" -dependencies = [ - "lazy_static", - "libc", - "log", - "ndk 0.2.1", - "ndk-macro", - "ndk-sys", -] - [[package]] name = "ndk-glue" version = "0.3.0" @@ -681,7 +703,7 @@ dependencies = [ "lazy_static", "libc", "log", - "ndk 0.3.0", + "ndk", "ndk-macro", "ndk-sys", ] @@ -693,7 +715,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "05d1c6307dc424d0f65b9b06e94f88248e6305726b14729fd67a5e47b2dc481d" dependencies = [ "darling", - "proc-macro-crate", + "proc-macro-crate 0.1.5", "proc-macro2", "quote", "syn", @@ -705,17 +727,6 @@ version = "0.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c44922cb3dbb1c70b5e5f443d63b64363a898564d739ba5198e3a9138442868d" -[[package]] -name = "net2" -version = "0.2.35" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3ebc3ec692ed7c9a255596c67808dee269f64655d8baf7b4f0638e51ba1d6853" -dependencies = [ - "cfg-if 0.1.10", - "libc", - "winapi 0.3.9", -] - [[package]] name = "nix" version = "0.18.0" @@ -729,54 +740,54 @@ dependencies = [ ] [[package]] -name = "nom" -version = "5.1.2" +name = "nix" +version = "0.20.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ffb4262d26ed83a1c0a33a38fe2bb15797329c85770da05e6b828ddb782627af" +checksum = "fa9b4819da1bc61c0ea48b63b7bc8604064dd43013e7cc325df098d49cd7c18a" +dependencies = [ + "bitflags", + "cc", + "cfg-if 1.0.0", + "libc", +] + +[[package]] +name = "nom" +version = "7.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1b1d11e1ef389c76fe5b81bcaf2ea32cf88b62bc494e19f493d0b30e7a930109" dependencies = [ "memchr", + "minimal-lexical", "version_check", ] [[package]] -name = "num_enum" -version = "0.4.3" +name = "ntapi" +version = "0.3.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ca565a7df06f3d4b485494f25ba05da1435950f4dc263440eda7a6fa9b8e36e4" +checksum = "3f6bb902e437b6d86e03cce10a7e2af662292c5dfef23b65899ea3ac9354ad44" dependencies = [ - "derivative", - "num_enum_derive 0.4.3", + "winapi", ] [[package]] name = "num_enum" -version = "0.5.1" +version = "0.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "226b45a5c2ac4dd696ed30fa6b94b057ad909c7b7fc2e0d0808192bced894066" +checksum = "3f9bd055fb730c4f8f4f57d45d35cd6b3f0980535b056dc7ff119cee6a66ed6f" dependencies = [ "derivative", - "num_enum_derive 0.5.1", + "num_enum_derive", ] [[package]] name = "num_enum_derive" -version = "0.4.3" +version = "0.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ffa5a33ddddfee04c0283a7653987d634e880347e96b5b2ed64de07efb59db9d" +checksum = "486ea01961c4a818096de679a8b740b26d9033146ac5291b1c98557658f8cdd9" dependencies = [ - "proc-macro-crate", - "proc-macro2", - "quote", - "syn", -] - -[[package]] -name = "num_enum_derive" -version = "0.5.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1c0fd9eba1d5db0994a239e09c1be402d35622277e35468ba891aa5e3188ce7e" -dependencies = [ - "proc-macro-crate", + "proc-macro-crate 1.1.0", "proc-macro2", "quote", "syn", @@ -803,9 +814,9 @@ dependencies = [ [[package]] name = "once_cell" -version = "1.5.2" +version = "1.8.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "13bd41f508810a131401606d54ac32a467c97172d74ba7662562ebba5ad07fa0" +checksum = "692fcb63b64b1758029e0a96ee63e049ce8c5948587f2f7208df04625e5f6b56" [[package]] name = "owned_ttf_parser" @@ -818,9 +829,9 @@ dependencies = [ [[package]] name = "parking_lot" -version = "0.11.0" +version = "0.11.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "a4893845fa2ca272e647da5d0e46660a314ead9c2fdd9a883aabc32e481a8733" +checksum = "7d17b78036a60663b797adeaee46f5c9dfebb86948d1255007a1d6be0271ff99" dependencies = [ "instant", "lock_api", @@ -829,17 +840,16 @@ dependencies = [ [[package]] name = "parking_lot_core" -version = "0.8.0" +version = "0.8.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c361aa727dd08437f2f1447be8b59a33b0edd15e0fcee698f935613d9efbca9b" +checksum = "d76e8e1493bcac0d2766c42737f34458f1c8c50c0d23bcb24ea953affb273216" dependencies = [ - "cfg-if 0.1.10", - "cloudabi", + "cfg-if 1.0.0", "instant", "libc", "redox_syscall", "smallvec", - "winapi 0.3.9", + "winapi", ] [[package]] @@ -863,8 +873,8 @@ name = "piet-gpu" version = "0.1.0" dependencies = [ "clap", - "ndk 0.3.0", - "ndk-glue 0.3.0", + "ndk", + "ndk-glue", "ndk-sys", "piet", "piet-gpu-hal", @@ -899,7 +909,7 @@ dependencies = [ "objc", "raw-window-handle", "smallvec", - "winapi 0.3.9", + "winapi", "wio", ] @@ -922,15 +932,15 @@ dependencies = [ [[package]] name = "pkg-config" -version = "0.3.19" +version = "0.3.22" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3831453b3449ceb48b6d9c7ad7c96d5ea673e9b470a1dc578c2ce6521230884c" +checksum = "12295df4f294471248581bc09bef3c38a5e46f1e36d6a37353621a0c6c357e1f" [[package]] name = "png" -version = "0.16.7" +version = "0.16.8" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "dfe7f9f1c730833200b134370e1d5098964231af8450bce9b78ee3ab5278b970" +checksum = "3c3287920cb847dee3de33d301c463fba14dda99db24214ddf93f83d3021f4c6" dependencies = [ "bitflags", "crc32fast", @@ -940,9 +950,9 @@ dependencies = [ [[package]] name = "ppv-lite86" -version = "0.2.10" +version = "0.2.15" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ac74c624d6b2d21f425f752262f42188365d7b8ff1aff74c82e45136510a4857" +checksum = "ed0cfbc8191465bed66e1718596ee0b0b35d5ee1f41c5df2189d0fe8bde535ba" [[package]] name = "proc-macro-crate" @@ -954,19 +964,29 @@ dependencies = [ ] [[package]] -name = "proc-macro2" -version = "1.0.24" +name = "proc-macro-crate" +version = "1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1e0704ee1a7e00d7bb417d0770ea303c1bccbabf0ef1667dae92b5967f5f8a71" +checksum = "1ebace6889caf889b4d3f76becee12e90353f2b8c7d875534a71e5742f8f6f83" +dependencies = [ + "thiserror", + "toml", +] + +[[package]] +name = "proc-macro2" +version = "1.0.32" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ba508cc11742c0dc5c1659771673afbab7a0efab23aa17e854cbab0837ed0b43" dependencies = [ "unicode-xid", ] [[package]] name = "quote" -version = "1.0.7" +version = "1.0.10" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "aa563d17ecb180e500da1cfd2b028310ac758de548efdd203e18f283af693f37" +checksum = "38bc8cc6a5f2e3655e0899c1b848643b2562f853f114bfec7be120678e3ace05" dependencies = [ "proc-macro2", ] @@ -977,7 +997,7 @@ version = "0.7.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6a6b1679d49b24bbfe0c803429aa1874472f50d9b363131f0e89fc356b544d03" dependencies = [ - "getrandom", + "getrandom 0.1.16", "libc", "rand_chacha", "rand_core", @@ -1000,7 +1020,7 @@ version = "0.5.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "90bde5296fc891b0cef12a6d03ddccc162ce7b2aff54160af9338f8d40df6d19" dependencies = [ - "getrandom", + "getrandom 0.1.16", ] [[package]] @@ -1023,27 +1043,40 @@ dependencies = [ [[package]] name = "raw-window-metal" -version = "0.1.0" +version = "0.1.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6b0f43bdc87adef4ce827b07775c9e59716b52f369696e7fb4ec7c4acb4e20b1" +checksum = "2cd21ed1cdef7f1b1579b972148ba6058b5b545959a14d91ea83c4f0ea9f289b" dependencies = [ - "cocoa 0.20.2", - "core-graphics 0.19.2", + "cocoa", + "core-graphics 0.22.3", "objc", "raw-window-handle", ] [[package]] name = "redox_syscall" -version = "0.1.57" +version = "0.2.10" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "41cc0f7e4d5d4544e8861606a285bb08d3e70712ccc7d2b84d7c0ccfaf4b05ce" +checksum = "8383f39639269cde97d255a32bdb68c047337295414940c68bdd30c2e13203ff" +dependencies = [ + "bitflags", +] + +[[package]] +name = "redox_users" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "528532f3d801c87aec9def2add9ca802fe569e44a544afe633765267840abe64" +dependencies = [ + "getrandom 0.2.3", + "redox_syscall", +] [[package]] name = "roxmltree" -version = "0.13.0" +version = "0.13.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "17dfc6c39f846bfc7d2ec442ad12055d79608d501380789b965d22f9354451f2" +checksum = "dbf7d7b1ea646d380d0e8153158063a6da7efe30ddbf3184042848e3f8a6f671" dependencies = [ "xmlparser", ] @@ -1081,37 +1114,30 @@ checksum = "d29ab0c6d3fc0ee92fe66e2d99f700eab17a8d57d1c1d3b748380fb20baa78cd" [[package]] name = "serde" -version = "1.0.117" +version = "1.0.130" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b88fa983de7720629c9387e9f517353ed404164b1e482c970a90c1a4aaf7dc1a" - -[[package]] -name = "slab" -version = "0.4.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c111b5bd5695e56cffe5129854aa230b39c93a305372fdbb2668ca2394eea9f8" +checksum = "f12d06de37cf59146fbdecab66aa99f9fe4f78722e3607577a5375d66bd0c913" [[package]] name = "smallvec" -version = "1.6.1" +version = "1.7.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fe0f37c9e8f3c5a4a66ad655a93c74daac4ad00c441533bf5c6e7990bb42604e" +checksum = "1ecab6c735a6bb4139c0caafd0cc3635748bbb3acf4550e8138122099251f309" [[package]] name = "smithay-client-toolkit" -version = "0.12.0" +version = "0.12.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2ec5c077def8af49f9b5aeeb5fcf8079c638c6615c3a8f9305e2dea601de57f7" +checksum = "4750c76fd5d3ac95fa3ed80fe667d6a3d8590a960e5b575b98eea93339a80b80" dependencies = [ "andrew", "bitflags", - "byteorder", "calloop", - "dlib", + "dlib 0.4.2", "lazy_static", "log", - "memmap", - "nix", + "memmap2", + "nix 0.18.0", "wayland-client", "wayland-cursor", "wayland-protocols", @@ -1141,9 +1167,9 @@ dependencies = [ [[package]] name = "syn" -version = "1.0.48" +version = "1.0.81" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "cc371affeffc477f42a221a1e4297aedcea33d47d19b61455588bd9d8f6b19ac" +checksum = "f2afee18b8beb5a596ecb4a2dce128c719b4ba399d34126b9e4396e3f9860966" dependencies = [ "proc-macro2", "quote", @@ -1161,18 +1187,18 @@ dependencies = [ [[package]] name = "thiserror" -version = "1.0.22" +version = "1.0.30" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0e9ae34b84616eedaaf1e9dd6026dbe00dcafa92aa0c8077cb69df1fcfe5e53e" +checksum = "854babe52e4df1653706b98fcfc05843010039b406875930a70e4d9644e5c417" dependencies = [ "thiserror-impl", ] [[package]] name = "thiserror-impl" -version = "1.0.22" +version = "1.0.30" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "9ba20f23e85b10754cd195504aebf6a27e2e6cbe28c17778a0c930724628dd56" +checksum = "aa32fd3f627f367fe16f893e2597ae3c05020f8bba2666a4e6ea73d377e5714b" dependencies = [ "proc-macro2", "quote", @@ -1181,9 +1207,9 @@ dependencies = [ [[package]] name = "toml" -version = "0.5.7" +version = "0.5.8" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "75cf45bb0bef80604d001caaec0d09da99611b3c0fd39d3080468875cdb65645" +checksum = "a31142970826733df8241ef35dc040ef98c679ab14d7c3e54d827099b3acecaa" dependencies = [ "serde", ] @@ -1247,15 +1273,15 @@ dependencies = [ [[package]] name = "unicode-width" -version = "0.1.8" +version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "9337591893a19b88d8d87f2cec1e73fad5cdfd10e5a6f349f498ad6ea2ffb1e3" +checksum = "3ed742d4ea2bd1176e236172c8429aaf54486e7ac098db29ffe6529e0ce50973" [[package]] name = "unicode-xid" -version = "0.2.1" +version = "0.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f7fe0bb3479651439c9112f72b6c505038574c9fbb575ed1bf3b797fa39dd564" +checksum = "8ccb82d61f80a663efe1f787a51b16b5a51e3314d6ac365b08639f52387b33f3" [[package]] name = "vec_map" @@ -1265,18 +1291,18 @@ checksum = "f1bddf1187be692e79c5ffeab891132dfb0f236ed36a43c7ed39f1165ee20191" [[package]] name = "version_check" -version = "0.9.2" +version = "0.9.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b5a972e5669d67ba988ce3dc826706fb0a8b01471c088cb0b6110b805cc36aed" +checksum = "5fecdca9a5291cc2b8dcf7dc02453fee791a280f3743cb0905f8822ae463b3fe" [[package]] name = "walkdir" -version = "2.3.1" +version = "2.3.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "777182bc735b6424e1a57516d35ed72cb8019d85c8c9bf536dccb3445c1a2f7d" +checksum = "808cf2735cd4b6866113f648b791c6adc5714537bc222d9347bb203386ffda56" dependencies = [ "same-file", - "winapi 0.3.9", + "winapi", "winapi-util", ] @@ -1287,15 +1313,21 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "cccddf32554fecc6acb585f82a32a72e28b48f8c4c1883ddfeeeaa96f7d8e519" [[package]] -name = "wayland-client" -version = "0.28.2" +name = "wasi" +version = "0.10.2+wasi-snapshot-preview1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "222b227f47871e47d657c1c5e5360b4af9a877aa9c892716787be1c192c78c42" +checksum = "fd6fbd9a79829dd1ad0cc20627bf1ed606756a7f77edff7b66b7064f9cb327c6" + +[[package]] +name = "wayland-client" +version = "0.28.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e3ab332350e502f159382201394a78e3cc12d0f04db863429260164ea40e0355" dependencies = [ "bitflags", "downcast-rs", "libc", - "nix", + "nix 0.20.0", "scoped-tls", "wayland-commons", "wayland-scanner", @@ -1304,11 +1336,11 @@ dependencies = [ [[package]] name = "wayland-commons" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "230b3ffeda101f877ff8ecb8573f5d26e7beb345b197807c4df34ec06879a3e6" +checksum = "a21817947c7011bbd0a27e11b17b337bfd022e8544b071a2641232047966fbda" dependencies = [ - "nix", + "nix 0.20.0", "once_cell", "smallvec", "wayland-sys", @@ -1316,20 +1348,20 @@ dependencies = [ [[package]] name = "wayland-cursor" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0aad1b4301cdccfb5f64056a4736e8155a5f4734bac41fdbca80b1fdbe1ab3e1" +checksum = "be610084edd1586d45e7bdd275fe345c7c1873598caa464c4fb835dee70fa65a" dependencies = [ - "nix", + "nix 0.20.0", "wayland-client", "xcursor", ] [[package]] name = "wayland-protocols" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "dc16a9db803cae58b45f9a84a6cf364434cc49a95c8b1ef98ffeb467d228bdc9" +checksum = "286620ea4d803bacf61fa087a4242ee316693099ee5a140796aaba02b29f861f" dependencies = [ "bitflags", "wayland-client", @@ -1339,9 +1371,9 @@ dependencies = [ [[package]] name = "wayland-scanner" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5ee5bd43a1d746efc486515fec561e47205f328b74802b959f10f5500f7e56cc" +checksum = "ce923eb2deb61de332d1f356ec7b6bf37094dc5573952e1c8936db03b54c03f1" dependencies = [ "proc-macro2", "quote", @@ -1350,21 +1382,15 @@ dependencies = [ [[package]] name = "wayland-sys" -version = "0.28.2" +version = "0.28.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0814adbecc7ea97869971e1d1c1b657e31863dda6fd768f119ad3dc408a01e58" +checksum = "d841fca9aed7febf9bed2e9796c49bf58d4152ceda8ac949ebe00868d8f0feb8" dependencies = [ - "dlib", + "dlib 0.5.0", "lazy_static", "pkg-config", ] -[[package]] -name = "winapi" -version = "0.2.8" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "167dc9d6949a9b857f3451275e911c3f44255842c1f7a76f33c55103a909087a" - [[package]] name = "winapi" version = "0.3.9" @@ -1375,12 +1401,6 @@ dependencies = [ "winapi-x86_64-pc-windows-gnu", ] -[[package]] -name = "winapi-build" -version = "0.1.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2d315eee3b34aca4797b2da6b13ed88266e6d612562a0c46390af8299fc699bc" - [[package]] name = "winapi-i686-pc-windows-gnu" version = "0.4.0" @@ -1393,7 +1413,7 @@ version = "0.1.5" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "70ec6ce85bb158151cae5e5c87f95a8e97d2c0c4b001223f33a334e3ce5de178" dependencies = [ - "winapi 0.3.9", + "winapi", ] [[package]] @@ -1404,14 +1424,14 @@ checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" [[package]] name = "winit" -version = "0.23.0" +version = "0.25.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b5bc559da567d8aa671bbcd08304d49e982c7bf2cb91e10288b9188931c1b772" +checksum = "79610794594d5e86be473ef7763f604f2159cbac8c94debd00df8fb41e86c2f8" dependencies = [ "bitflags", - "cocoa 0.23.0", - "core-foundation 0.9.1", - "core-graphics 0.22.1", + "cocoa", + "core-foundation 0.9.2", + "core-graphics 0.22.3", "core-video-sys", "dispatch", "instant", @@ -1419,17 +1439,18 @@ dependencies = [ "libc", "log", "mio", - "mio-extras", - "ndk 0.2.1", - "ndk-glue 0.2.1", + "mio-misc", + "ndk", + "ndk-glue", "ndk-sys", "objc", "parking_lot", "percent-encoding", "raw-window-handle", + "scopeguard", "smithay-client-toolkit", "wayland-client", - "winapi 0.3.9", + "winapi", "x11-dl", ] @@ -1439,51 +1460,43 @@ version = "0.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "5d129932f4644ac2396cb456385cbf9e63b5b30c6e8dc4820bdca4eb082037a5" dependencies = [ - "winapi 0.3.9", -] - -[[package]] -name = "ws2_32-sys" -version = "0.2.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d59cefebd0c892fa2dd6de581e937301d8552cb44489cdff035c6187cb63fa5e" -dependencies = [ - "winapi 0.2.8", - "winapi-build", + "winapi", ] [[package]] name = "x11-dl" -version = "2.18.5" +version = "2.19.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2bf981e3a5b3301209754218f962052d4d9ee97e478f4d26d4a6eced34c1fef8" +checksum = "ea26926b4ce81a6f5d9d0f3a0bc401e5a37c6ae14a1bfaa8ff6099ca80038c59" dependencies = [ "lazy_static", "libc", - "maybe-uninit", "pkg-config", ] [[package]] name = "xcursor" -version = "0.3.2" +version = "0.3.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d3a481cfdefd35e1c50073ae33a8000d695c98039544659f5dc5dd71311b0d01" +checksum = "463705a63313cd4301184381c5e8042f0a7e9b4bb63653f216311d4ae74690b7" dependencies = [ "nom", ] [[package]] name = "xdg" -version = "2.2.0" +version = "2.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d089681aa106a86fade1b0128fb5daf07d5867a509ab036d99988dec80429a57" +checksum = "3a23fe958c70412687039c86f578938b4a0bb50ec788e96bce4d6ab00ddd5803" +dependencies = [ + "dirs", +] [[package]] name = "xml-rs" -version = "0.8.3" +version = "0.8.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b07db065a5cf61a7e4ba64f29e67db906fb1787316516c4e6e5ff0fea1efcd8a" +checksum = "d2d7d3948613f75c98fd9328cfdcc45acc4d360655289d0a7d4ec931392200a3" [[package]] name = "xmlparser" diff --git a/piet-gpu-hal/examples/collatz.rs b/piet-gpu-hal/examples/collatz.rs index cad508e..dae5b31 100644 --- a/piet-gpu-hal/examples/collatz.rs +++ b/piet-gpu-hal/examples/collatz.rs @@ -1,8 +1,8 @@ -use piet_gpu_hal::include_shader; -use piet_gpu_hal::{BufferUsage, Instance, Session}; +use piet_gpu_hal::{include_shader, BindType}; +use piet_gpu_hal::{BufferUsage, Instance, InstanceFlags, Session}; fn main() { - let (instance, _) = Instance::new(None).unwrap(); + let (instance, _) = Instance::new(None, InstanceFlags::empty()).unwrap(); unsafe { let device = instance.device(None).unwrap(); let session = Session::new(device); @@ -10,7 +10,9 @@ fn main() { let src = (0..256).map(|x| x + 1).collect::>(); let buffer = session.create_buffer_init(&src, usage).unwrap(); let code = include_shader!(&session, "./shader/gen/collatz"); - let pipeline = session.create_simple_compute_pipeline(code, 1).unwrap(); + let pipeline = session + .create_compute_pipeline(code, &[BindType::Buffer]) + .unwrap(); let descriptor_set = session .create_simple_descriptor_set(&pipeline, &[&buffer]) .unwrap(); diff --git a/piet-gpu-hal/examples/shader/build.ninja b/piet-gpu-hal/examples/shader/build.ninja index f1c6328..3b9cf3f 100644 --- a/piet-gpu-hal/examples/shader/build.ninja +++ b/piet-gpu-hal/examples/shader/build.ninja @@ -4,6 +4,7 @@ glslang_validator = glslangValidator spirv_cross = spirv-cross +dxc = dxc rule glsl command = $glslang_validator -V -o $out $in @@ -11,9 +12,13 @@ rule glsl rule hlsl command = $spirv_cross --hlsl $in --output $out +rule dxil + command = $dxc -T cs_6_0 $in -Fo $out + rule msl command = $spirv_cross --msl $in --output $out build gen/collatz.spv: glsl collatz.comp build gen/collatz.hlsl: hlsl gen/collatz.spv +build gen/collatz.dxil: dxil gen/collatz.hlsl build gen/collatz.msl: msl gen/collatz.spv diff --git a/piet-gpu-hal/examples/shader/gen/collatz.dxil b/piet-gpu-hal/examples/shader/gen/collatz.dxil new file mode 100644 index 0000000..a03f96a Binary files /dev/null and b/piet-gpu-hal/examples/shader/gen/collatz.dxil differ diff --git a/piet-gpu-hal/src/backend.rs b/piet-gpu-hal/src/backend.rs index 0fc3920..8df7354 100644 --- a/piet-gpu-hal/src/backend.rs +++ b/piet-gpu-hal/src/backend.rs @@ -16,7 +16,7 @@ //! The generic trait for backends to implement. -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; pub trait Device: Sized { type Buffer: 'static; @@ -27,7 +27,6 @@ pub trait Device: Sized { type CmdBuf: CmdBuf; type Fence; type Semaphore; - type PipelineBuilder: PipelineBuilder; type DescriptorSetBuilder: DescriptorSetBuilder; type Sampler; type ShaderSource: ?Sized; @@ -60,33 +59,21 @@ pub trait Device: Sized { /// Maybe doesn't need result return? unsafe fn destroy_image(&self, image: &Self::Image) -> Result<(), Error>; - /// Start building a pipeline. + /// Build a compute pipeline. /// /// A pipeline is a bit of shader IR plus a signature for what kinds of resources /// it expects. - unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder; + unsafe fn create_compute_pipeline( + &self, + code: &Self::ShaderSource, + bind_types: &[BindType], + ) -> Result; /// Start building a descriptor set. /// /// A descriptor set is a binding of resources for a given pipeline. unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder; - /// Create a simple compute pipeline that operates on buffers and storage images. - /// - /// This is provided as a convenience but will probably go away, as the functionality - /// is subsumed by the builder. - unsafe fn create_simple_compute_pipeline( - &self, - code: &Self::ShaderSource, - n_buffers: u32, - n_images: u32, - ) -> Result { - let mut builder = self.pipeline_builder(); - builder.add_buffers(n_buffers); - builder.add_images(n_images); - builder.create_compute_pipeline(self, code) - } - /// Create a descriptor set for a given pipeline, binding buffers and images. /// /// This is provided as a convenience but will probably go away, as the functionality @@ -236,21 +223,6 @@ pub trait CmdBuf { unsafe fn finish_timestamps(&mut self, _pool: &D::QueryPool) {} } -/// A builder for pipelines with more complex layouts. -pub trait PipelineBuilder { - /// Add buffers to the pipeline. Each has its own binding. - fn add_buffers(&mut self, n_buffers: u32); - /// Add storage images to the pipeline. Each has its own binding. - fn add_images(&mut self, n_images: u32); - /// Add a binding with a variable-size array of textures. - fn add_textures(&mut self, max_textures: u32); - unsafe fn create_compute_pipeline( - self, - device: &D, - code: &D::ShaderSource, - ) -> Result; -} - /// A builder for descriptor sets with more complex layouts. /// /// Note: the order needs to match the pipeline building, and it also needs to diff --git a/piet-gpu-hal/src/dx12.rs b/piet-gpu-hal/src/dx12.rs index 557df83..66befa5 100644 --- a/piet-gpu-hal/src/dx12.rs +++ b/piet-gpu-hal/src/dx12.rs @@ -3,17 +3,19 @@ mod error; mod wrappers; -use std::{cell::Cell, convert::TryInto, mem, ptr}; +use std::{cell::Cell, convert::{TryFrom, TryInto}, mem, ptr}; use winapi::shared::minwindef::TRUE; -use winapi::shared::{dxgi, dxgi1_2, dxgi1_3, dxgitype}; +use winapi::shared::{dxgi, dxgi1_2, dxgitype}; +#[allow(unused)] +use winapi::shared::dxgi1_3; // for error reporting in debug mode use winapi::um::d3d12; use raw_window_handle::{HasRawWindowHandle, RawWindowHandle}; use smallvec::SmallVec; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, WorkgroupLimits}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, WorkgroupLimits}; use self::wrappers::{CommandAllocator, CommandQueue, Device, Factory4, Resource, ShaderByteCode}; @@ -83,13 +85,6 @@ pub struct Fence { /// semaphore is needed for presentation on DX12. pub struct Semaphore; -#[derive(Default)] -pub struct PipelineBuilder { - ranges: Vec, - n_uav: u32, - // TODO: add counters for other resource types -} - // TODO #[derive(Default)] pub struct DescriptorSetBuilder { @@ -239,14 +234,13 @@ impl crate::backend::Device for Dx12Device { type Semaphore = Semaphore; - type PipelineBuilder = PipelineBuilder; - type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = (); - // Currently this is HLSL source, but we'll probably change it to IR. - type ShaderSource = str; + // Currently due to type inflexibility this is hardcoded to either HLSL or + // DXIL, but it would be nice to be able to handle both at runtime. + type ShaderSource = [u8]; fn create_buffer(&self, size: u64, usage: BufferUsage) -> Result { // TODO: consider supporting BufferUsage::QUERY_RESOLVE here rather than @@ -289,9 +283,7 @@ impl crate::backend::Device for Dx12Device { fn create_cmd_buf(&self) -> Result { let list_type = d3d12::D3D12_COMMAND_LIST_TYPE_DIRECT; - let allocator = - unsafe { self.device.create_command_allocator(list_type)? } - ; + let allocator = unsafe { self.device.create_command_allocator(list_type)? }; let node_mask = 0; unsafe { let c = self @@ -420,8 +412,94 @@ impl crate::backend::Device for Dx12Device { self.gpu_info.clone() } - unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { - PipelineBuilder::default() + unsafe fn create_compute_pipeline( + &self, + code: &Self::ShaderSource, + bind_types: &[BindType], + ) -> Result { + if u32::try_from(bind_types.len()).is_err() { + panic!("bind type length overflow"); + } + let mut ranges = Vec::new(); + let mut i = 0; + fn map_range_type(bind_type: BindType) -> d3d12::D3D12_DESCRIPTOR_RANGE_TYPE { + match bind_type { + BindType::Buffer | BindType::Image | BindType::ImageRead => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV, + BindType::BufReadOnly => d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_SRV, + } + } + while i < bind_types.len() { + let range_type = map_range_type(bind_types[i]); + let mut end = i + 1; + while end < bind_types.len() && map_range_type(bind_types[end]) == range_type { + end += 1; + } + let n_descriptors = (end - i) as u32; + ranges.push(d3d12::D3D12_DESCRIPTOR_RANGE { + RangeType: range_type, + NumDescriptors: n_descriptors, + BaseShaderRegister: i as u32, + RegisterSpace: 0, + OffsetInDescriptorsFromTableStart: d3d12::D3D12_DESCRIPTOR_RANGE_OFFSET_APPEND, + }); + i = end; + } + + // We could always have ShaderSource as [u8] even when it's HLSL, and use the + // magic number to distinguish. In any case, for now it's hardcoded as one or + // the other. + /* + // HLSL code path + #[cfg(debug_assertions)] + let flags = winapi::um::d3dcompiler::D3DCOMPILE_DEBUG + | winapi::um::d3dcompiler::D3DCOMPILE_SKIP_OPTIMIZATION; + #[cfg(not(debug_assertions))] + let flags = 0; + let shader_blob = ShaderByteCode::compile(code, "cs_5_1", "main", flags)?; + let shader = ShaderByteCode::from_blob(shader_blob); + */ + + // DXIL code path + let shader = ShaderByteCode::from_slice(code); + + let mut root_parameter = d3d12::D3D12_ROOT_PARAMETER { + ParameterType: d3d12::D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE, + ShaderVisibility: d3d12::D3D12_SHADER_VISIBILITY_ALL, + ..mem::zeroed() + }; + *root_parameter.u.DescriptorTable_mut() = d3d12::D3D12_ROOT_DESCRIPTOR_TABLE { + NumDescriptorRanges: ranges.len() as u32, + pDescriptorRanges: ranges.as_ptr(), + }; + let root_signature_desc = d3d12::D3D12_ROOT_SIGNATURE_DESC { + NumParameters: 1, + pParameters: &root_parameter, + NumStaticSamplers: 0, + pStaticSamplers: ptr::null(), + Flags: d3d12::D3D12_ROOT_SIGNATURE_FLAG_NONE, + }; + let root_signature_blob = wrappers::RootSignature::serialize_description( + &root_signature_desc, + d3d12::D3D_ROOT_SIGNATURE_VERSION_1, + )?; + let root_signature = self + .device + .create_root_signature(0, root_signature_blob)?; + let desc = d3d12::D3D12_COMPUTE_PIPELINE_STATE_DESC { + pRootSignature: root_signature.0.as_raw(), + CS: shader.bytecode, + NodeMask: 0, + CachedPSO: d3d12::D3D12_CACHED_PIPELINE_STATE { + pCachedBlob: ptr::null(), + CachedBlobSizeInBytes: 0, + }, + Flags: d3d12::D3D12_PIPELINE_STATE_FLAG_NONE, + }; + let pipeline_state = self.device.create_compute_pipeline_state(&desc)?; + Ok(Pipeline { + pipeline_state, + root_signature, + }) } unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder { @@ -451,8 +529,7 @@ impl Dx12Device { impl crate::backend::CmdBuf for CmdBuf { unsafe fn begin(&mut self) { - if self.needs_reset { - } + if self.needs_reset {} } unsafe fn finish(&mut self) { @@ -559,86 +636,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder for PipelineBuilder { - fn add_buffers(&mut self, n_buffers: u32) { - // Note: if the buffer is readonly, then it needs to be bound - // as an SRV, not a UAV. I think that requires distinguishing - // readonly and read-write cases in pipeline and descriptor set - // creation. For now we punt. - if n_buffers != 0 { - self.ranges.push(d3d12::D3D12_DESCRIPTOR_RANGE { - RangeType: d3d12::D3D12_DESCRIPTOR_RANGE_TYPE_UAV, - NumDescriptors: n_buffers, - BaseShaderRegister: self.n_uav, - RegisterSpace: 0, - OffsetInDescriptorsFromTableStart: d3d12::D3D12_DESCRIPTOR_RANGE_OFFSET_APPEND, - }); - self.n_uav += n_buffers; - } - } - - fn add_images(&mut self, n_images: u32) { - // These are UAV images, so the descriptor type is the same as buffers. - self.add_buffers(n_images); - } - - fn add_textures(&mut self, _max_textures: u32) { - todo!() - } - - unsafe fn create_compute_pipeline( - self, - device: &Dx12Device, - code: &str, - ) -> Result { - #[cfg(debug_assertions)] - let flags = winapi::um::d3dcompiler::D3DCOMPILE_DEBUG - | winapi::um::d3dcompiler::D3DCOMPILE_SKIP_OPTIMIZATION; - #[cfg(not(debug_assertions))] - let flags = 0; - let shader_blob = ShaderByteCode::compile(code, "cs_5_1", "main", flags)?; - let shader = ShaderByteCode::from_blob(shader_blob); - let mut root_parameter = d3d12::D3D12_ROOT_PARAMETER { - ParameterType: d3d12::D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE, - ShaderVisibility: d3d12::D3D12_SHADER_VISIBILITY_ALL, - ..mem::zeroed() - }; - *root_parameter.u.DescriptorTable_mut() = d3d12::D3D12_ROOT_DESCRIPTOR_TABLE { - NumDescriptorRanges: self.ranges.len().try_into()?, - pDescriptorRanges: self.ranges.as_ptr(), - }; - let root_signature_desc = d3d12::D3D12_ROOT_SIGNATURE_DESC { - NumParameters: 1, - pParameters: &root_parameter, - NumStaticSamplers: 0, - pStaticSamplers: ptr::null(), - Flags: d3d12::D3D12_ROOT_SIGNATURE_FLAG_NONE, - }; - let root_signature_blob = wrappers::RootSignature::serialize_description( - &root_signature_desc, - d3d12::D3D_ROOT_SIGNATURE_VERSION_1, - )?; - let root_signature = device - .device - .create_root_signature(0, root_signature_blob)?; - let desc = d3d12::D3D12_COMPUTE_PIPELINE_STATE_DESC { - pRootSignature: root_signature.0.as_raw(), - CS: shader.bytecode, - NodeMask: 0, - CachedPSO: d3d12::D3D12_CACHED_PIPELINE_STATE { - pCachedBlob: ptr::null(), - CachedBlobSizeInBytes: 0, - }, - Flags: d3d12::D3D12_PIPELINE_STATE_FLAG_NONE, - }; - let pipeline_state = device.device.create_compute_pipeline_state(&desc)?; - Ok(Pipeline { - pipeline_state, - root_signature, - }) - } -} - impl crate::backend::DescriptorSetBuilder for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { // Note: we could get rid of the clone here (which is an AddRef) diff --git a/piet-gpu-hal/src/dx12/wrappers.rs b/piet-gpu-hal/src/dx12/wrappers.rs index edec3e4..dd834fa 100644 --- a/piet-gpu-hal/src/dx12/wrappers.rs +++ b/piet-gpu-hal/src/dx12/wrappers.rs @@ -10,9 +10,7 @@ use crate::dx12::error::{self, error_if_failed_else_unit, explain_error, Error}; use std::convert::{TryFrom, TryInto}; use std::sync::atomic::{AtomicPtr, Ordering}; use std::{ffi, mem, ptr}; -use winapi::shared::{ - dxgi, dxgi1_2, dxgi1_3, dxgi1_4, dxgiformat, dxgitype, minwindef, windef, -}; +use winapi::shared::{dxgi, dxgi1_2, dxgi1_3, dxgi1_4, dxgiformat, dxgitype, minwindef, windef}; use winapi::um::d3dcommon::ID3DBlob; use winapi::um::{ d3d12, d3d12sdklayers, d3dcommon, d3dcompiler, dxgidebug, handleapi, synchapi, winnt, @@ -198,7 +196,7 @@ impl Factory4 { error_if_failed_else_unit(self.0.EnumAdapters1(id, &mut adapter))?; let mut desc = mem::zeroed(); (*adapter).GetDesc(&mut desc); - println!("desc: {:?}", desc.Description); + //println!("desc: {:?}", desc.Description); Ok(Adapter1(ComPtr::from_raw(adapter))) } @@ -278,6 +276,7 @@ impl SwapChain3 { } impl Blob { + #[allow(unused)] pub unsafe fn print_to_console(blob: &Blob) { println!("==SHADER COMPILE MESSAGES=="); let message = { @@ -563,7 +562,6 @@ impl Device { Ok(QueryHeap(ComPtr::from_raw(query_heap))) } - pub unsafe fn create_buffer( &self, buffer_size_in_bytes: u32, @@ -717,13 +715,13 @@ impl RootSignature { let hresult = d3d12::D3D12SerializeRootSignature(desc, version, &mut blob, &mut error_blob_ptr); - let error_blob = if error_blob_ptr.is_null() { - None - } else { - Some(Blob(ComPtr::from_raw(error_blob_ptr))) - }; #[cfg(debug_assertions)] { + let error_blob = if error_blob_ptr.is_null() { + None + } else { + Some(Blob(ComPtr::from_raw(error_blob_ptr))) + }; if let Some(error_blob) = &error_blob { Blob::print_to_console(error_blob); } @@ -739,6 +737,7 @@ impl ShaderByteCode { // `blob` may not be null. // TODO: this is not super elegant, maybe want to move the get // operations closer to where they're used. + #[allow(unused)] pub unsafe fn from_blob(blob: Blob) -> ShaderByteCode { ShaderByteCode { bytecode: d3d12::D3D12_SHADER_BYTECODE { @@ -752,6 +751,7 @@ impl ShaderByteCode { /// Compile a shader from raw HLSL. /// /// * `target`: example format: `ps_5_1`. + #[allow(unused)] pub unsafe fn compile( source: &str, target: &str, @@ -798,6 +798,24 @@ impl ShaderByteCode { Ok(Blob(ComPtr::from_raw(shader_blob_ptr))) } + + /// Create bytecode from a slice. + /// + /// # Safety + /// + /// This call elides the lifetime from the slice. The caller is responsible + /// for making sure the reference remains valid for the lifetime of this + /// object. + #[allow(unused)] + pub unsafe fn from_slice(bytecode: &[u8]) -> ShaderByteCode { + ShaderByteCode { + bytecode: d3d12::D3D12_SHADER_BYTECODE { + BytecodeLength: bytecode.len(), + pShaderBytecode: bytecode.as_ptr() as *const _, + }, + blob: None, + } + } } impl Fence { @@ -864,7 +882,11 @@ impl GraphicsCommandList { explain_error(self.0.Close(), "error closing command list") } - pub unsafe fn reset(&self, allocator: &CommandAllocator, initial_pso: Option<&PipelineState>) -> Result<(), Error> { + pub unsafe fn reset( + &self, + allocator: &CommandAllocator, + initial_pso: Option<&PipelineState>, + ) -> Result<(), Error> { let p_initial_state = initial_pso.map(|p| p.0.as_raw()).unwrap_or(ptr::null_mut()); error::error_if_failed_else_unit(self.0.Reset(allocator.0.as_raw(), p_initial_state)) } @@ -1072,9 +1094,8 @@ pub unsafe fn create_transition_resource_barrier( resource_barrier } +#[allow(unused)] pub unsafe fn enable_debug_layer() -> Result<(), Error> { - println!("enabling debug layer."); - let mut debug_controller: *mut d3d12sdklayers::ID3D12Debug1 = ptr::null_mut(); explain_error( d3d12::D3D12GetDebugInterface( diff --git a/piet-gpu-hal/src/hub.rs b/piet-gpu-hal/src/hub.rs index d79e955..2acfee0 100644 --- a/piet-gpu-hal/src/hub.rs +++ b/piet-gpu-hal/src/hub.rs @@ -11,9 +11,9 @@ use std::sync::{Arc, Mutex, Weak}; use smallvec::SmallVec; -use crate::mux; +use crate::{mux, BackendType}; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; +use crate::{BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams}; pub use crate::mux::{DescriptorSet, Fence, Pipeline, QueryPool, Sampler, Semaphore, ShaderCode}; @@ -100,12 +100,6 @@ struct BufferInner { session: Weak, } -/// A builder for creating pipelines. -/// -/// Configure the signature (buffers and images accessed) for a pipeline, -/// which is essentially compiled shader code, ready to be dispatched. -pub struct PipelineBuilder(mux::PipelineBuilder); - /// A builder for creating descriptor sets. /// /// Add bindings to the descriptor set before dispatching a shader. @@ -316,26 +310,16 @@ impl Session { self.0.device.create_semaphore() } - /// This creates a pipeline that operates on some buffers and images. - /// - /// The descriptor set layout is just some number of storage buffers - /// and storage images (this might change). - pub unsafe fn create_simple_compute_pipeline<'a>( - &self, - code: ShaderCode<'a>, - n_buffers: u32, - ) -> Result { - self.pipeline_builder() - .add_buffers(n_buffers) - .create_compute_pipeline(self, code) - } - - /// Start building a pipeline. + /// Create a compute shader pipeline. /// /// A pipeline is essentially a compiled shader, with more specific /// details about what resources may be bound to it. - pub unsafe fn pipeline_builder(&self) -> PipelineBuilder { - PipelineBuilder(self.0.device.pipeline_builder()) + pub unsafe fn create_compute_pipeline<'a>( + &self, + code: ShaderCode<'a>, + bind_types: &[BindType], + ) -> Result { + self.0.device.create_compute_pipeline(code, bind_types) } /// Create a descriptor set for a simple pipeline that just references buffers. @@ -385,8 +369,13 @@ impl Session { } /// Choose shader code from the available choices. - pub fn choose_shader<'a>(&self, spv: &'a [u8], hlsl: &'a str, msl: &'a str) -> ShaderCode<'a> { - self.0.device.choose_shader(spv, hlsl, msl) + pub fn choose_shader<'a>(&self, spv: &'a [u8], hlsl: &'a str, dxil: &'a [u8], msl: &'a str) -> ShaderCode<'a> { + self.0.device.choose_shader(spv, hlsl, dxil, msl) + } + + /// Report the backend type that was chosen. + pub fn backend_type(&self) -> BackendType { + self.0.device.backend_type() } } @@ -729,38 +718,6 @@ impl Buffer { } } -impl PipelineBuilder { - /// Add buffers to the pipeline. Each has its own binding. - pub fn add_buffers(mut self, n_buffers: u32) -> Self { - self.0.add_buffers(n_buffers); - self - } - - /// Add storage images to the pipeline. Each has its own binding. - pub fn add_images(mut self, n_images: u32) -> Self { - self.0.add_images(n_images); - self - } - - /// Add a binding with a variable-size array of textures. - pub fn add_textures(mut self, max_textures: u32) -> Self { - self.0.add_textures(max_textures); - self - } - - /// Create the compute pipeline. - /// - /// The shader code must be given in an appropriate format for - /// the back-end. See [`Session::choose_shader`] for a helper. - pub unsafe fn create_compute_pipeline<'a>( - self, - session: &Session, - code: ShaderCode<'a>, - ) -> Result { - self.0.create_compute_pipeline(&session.0.device, code) - } -} - impl DescriptorSetBuilder { pub fn add_buffers<'a>(mut self, buffers: impl IntoRefs<'a, Buffer>) -> Self { let mux_buffers = buffers diff --git a/piet-gpu-hal/src/lib.rs b/piet-gpu-hal/src/lib.rs index 2dd0eff..05e2394 100644 --- a/piet-gpu-hal/src/lib.rs +++ b/piet-gpu-hal/src/lib.rs @@ -1,7 +1,8 @@ -/// The cross-platform abstraction for a GPU device. -/// -/// This abstraction is inspired by gfx-hal, but is specialized to the needs of piet-gpu. -/// In time, it may go away and be replaced by either gfx-hal or wgpu. +//! The cross-platform abstraction for a GPU device. +//! +//! This abstraction is inspired by gfx-hal, but is specialized to the needs of piet-gpu. +//! In time, it may go away and be replaced by either gfx-hal or wgpu. + use bitflags::bitflags; mod backend; @@ -17,8 +18,8 @@ pub use crate::mux::{ Swapchain, }; pub use hub::{ - Buffer, CmdBuf, DescriptorSetBuilder, Image, PipelineBuilder, PlainData, RetainResource, - Session, SubmittedCmdBuf, + Buffer, CmdBuf, DescriptorSetBuilder, Image, PlainData, RetainResource, Session, + SubmittedCmdBuf, }; // TODO: because these are conditionally included, "cargo fmt" does not @@ -36,9 +37,27 @@ mod metal; /// The common error type for the crate. /// -/// This keeps things imple and can be expanded later. +/// This keeps things simple and can be expanded later. pub type Error = Box; +bitflags! { + /// Options when creating an instance. + #[derive(Default)] + pub struct InstanceFlags: u32 { + /// Prefer DX12 over Vulkan. + const DX12 = 0x1; + // TODO: discrete vs integrated selection + } +} + +/// The GPU backend that was selected. +#[derive(Clone, Copy, PartialEq, Eq, Debug)] +pub enum BackendType { + Vulkan, + Dx12, + Metal, +} + /// An image layout state. /// /// An image must be in a particular layout state to be used for @@ -84,10 +103,31 @@ bitflags! { const STORAGE = 0x80; /// The buffer can be used to store the results of queries. const QUERY_RESOLVE = 0x200; + /// The buffer may be cleared. + const CLEAR = 0x8000; // May add other types. } } +/// The type of resource that will be bound to a slot in a shader. +#[derive(Clone, Copy, PartialEq, Eq)] +pub enum BindType { + /// A storage buffer with read/write access. + Buffer, + /// A storage buffer with read only access. + BufReadOnly, + /// A storage image. + Image, + /// A storage image with read only access. + /// + /// A note on this. None of the backends are currently making a + /// distinction between Image and ImageRead as far as bindings go, + /// but the `--hlsl-nonwritable-uav-texture-as-srv` option to + /// spirv-cross (marked as unstable) would do so. + ImageRead, + // TODO: Uniform, Sampler, maybe others +} + #[derive(Clone, Debug)] /// Information about the GPU. pub struct GpuInfo { diff --git a/piet-gpu-hal/src/macros.rs b/piet-gpu-hal/src/macros.rs index 38897a8..a4a441e 100644 --- a/piet-gpu-hal/src/macros.rs +++ b/piet-gpu-hal/src/macros.rs @@ -198,6 +198,7 @@ macro_rules! include_shader { $device.choose_shader( include_bytes!(concat!($path_base, ".spv")), include_str!(concat!($path_base, ".hlsl")), + include_bytes!(concat!($path_base, ".dxil")), include_str!(concat!($path_base, ".msl")), ) }; diff --git a/piet-gpu-hal/src/metal.rs b/piet-gpu-hal/src/metal.rs index 4da8491..78c0682 100644 --- a/piet-gpu-hal/src/metal.rs +++ b/piet-gpu-hal/src/metal.rs @@ -82,8 +82,6 @@ pub struct CmdBuf { pub struct QueryPool; -pub struct PipelineBuilder; - pub struct Pipeline(metal::ComputePipelineState); #[derive(Default)] @@ -220,8 +218,6 @@ impl crate::backend::Device for MtlDevice { type Semaphore = Semaphore; - type PipelineBuilder = PipelineBuilder; - type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = (); @@ -273,8 +269,18 @@ impl crate::backend::Device for MtlDevice { todo!() } - unsafe fn pipeline_builder(&self) -> Self::PipelineBuilder { - PipelineBuilder + unsafe fn create_compute_pipeline( + &self, + code: &Self::ShaderSource, + _bind_types: &[crate::BindType], + ) -> Result { + let options = metal::CompileOptions::new(); + let library = self.device.new_library_with_source(code, &options)?; + let function = library.get_function("main0", None)?; + let pipeline = self + .device + .new_compute_pipeline_state_with_function(&function)?; + Ok(Pipeline(pipeline)) } unsafe fn descriptor_set_builder(&self) -> Self::DescriptorSetBuilder { @@ -552,33 +558,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder for PipelineBuilder { - fn add_buffers(&mut self, _n_buffers: u32) { - // My understanding is that Metal infers the pipeline layout from - // the source. - } - - fn add_images(&mut self, _n_images: u32) {} - - fn add_textures(&mut self, _max_textures: u32) {} - - unsafe fn create_compute_pipeline( - self, - device: &MtlDevice, - code: &str, - ) -> Result { - let options = metal::CompileOptions::new(); - // Probably want to set MSL version here. - let library = device.device.new_library_with_source(code, &options)?; - // This seems to be the default name from spirv-cross, but we may need to tweak. - let function = library.get_function("main0", None)?; - let pipeline = device - .device - .new_compute_pipeline_state_with_function(&function)?; - Ok(Pipeline(pipeline)) - } -} - impl crate::backend::DescriptorSetBuilder for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { self.0.buffers.extend(buffers.iter().copied().cloned()); diff --git a/piet-gpu-hal/src/mux.rs b/piet-gpu-hal/src/mux.rs index 4af5b3e..a0ea28a 100644 --- a/piet-gpu-hal/src/mux.rs +++ b/piet-gpu-hal/src/mux.rs @@ -33,8 +33,9 @@ mux_cfg! { use crate::backend::CmdBuf as CmdBufTrait; use crate::backend::DescriptorSetBuilder as DescriptorSetBuilderTrait; use crate::backend::Device as DeviceTrait; -use crate::backend::PipelineBuilder as PipelineBuilderTrait; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout}; +use crate::BackendType; +use crate::BindType; +use crate::{BufferUsage, Error, GpuInfo, ImageLayout, InstanceFlags}; mux_enum! { /// An instance, selected from multiple backends. @@ -84,7 +85,6 @@ mux_device_enum! { /// presentation by the back-end, this may or may not be a "real" /// semaphore. Semaphore } -mux_device_enum! { PipelineBuilder } mux_device_enum! { /// A pipeline object; basically a compiled shader. Pipeline } @@ -104,6 +104,8 @@ pub enum ShaderCode<'a> { Spv(&'a [u8]), /// HLSL (source) Hlsl(&'a str), + /// DXIL (DX12 intermediate language) + Dxil(&'a [u8]), /// Metal Shading Language (source) Msl(&'a str), } @@ -118,22 +120,33 @@ impl Instance { /// work. pub fn new( window_handle: Option<&dyn raw_window_handle::HasRawWindowHandle>, + flags: InstanceFlags, ) -> Result<(Instance, Option), Error> { - mux_cfg! { - #[cfg(vk)] - { - let result = vulkan::VkInstance::new(window_handle); - if let Ok((instance, surface)) = result { - return Ok((Instance::Vk(instance), surface.map(Surface::Vk))); + let mut backends = [BackendType::Vulkan, BackendType::Dx12]; + if flags.contains(InstanceFlags::DX12) { + backends.swap(0, 1); + } + for backend in backends { + if backend == BackendType::Vulkan { + mux_cfg! { + #[cfg(vk)] + { + let result = vulkan::VkInstance::new(window_handle); + if let Ok((instance, surface)) = result { + return Ok((Instance::Vk(instance), surface.map(Surface::Vk))); + } + } } } - } - mux_cfg! { - #[cfg(dx12)] - { - let result = dx12::Dx12Instance::new(window_handle); - if let Ok((instance, surface)) = result { - return Ok((Instance::Dx12(instance), surface.map(Surface::Dx12))); + if backend == BackendType::Dx12 { + mux_cfg! { + #[cfg(dx12)] + { + let result = dx12::Dx12Instance::new(window_handle); + if let Ok((instance, surface)) = result { + return Ok((Instance::Dx12(instance), surface.map(Surface::Dx12))); + } + } } } } @@ -293,11 +306,40 @@ impl Device { } } - pub unsafe fn pipeline_builder(&self) -> PipelineBuilder { + pub unsafe fn create_compute_pipeline<'a>( + &self, + code: ShaderCode<'a>, + bind_types: &[BindType], + ) -> Result { mux_match! { self; - Device::Vk(d) => PipelineBuilder::Vk(d.pipeline_builder()), - Device::Dx12(d) => PipelineBuilder::Dx12(d.pipeline_builder()), - Device::Mtl(d) => PipelineBuilder::Mtl(d.pipeline_builder()), + Device::Vk(d) => { + let shader_code = match code { + ShaderCode::Spv(spv) => spv, + // Panic or return "incompatible shader" error here? + _ => panic!("Vulkan backend requires shader code in SPIR-V format"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Vk) + } + Device::Dx12(d) => { + let shader_code = match code { + //ShaderCode::Hlsl(hlsl) => hlsl, + ShaderCode::Dxil(dxil) => dxil, + // Panic or return "incompatible shader" error here? + _ => panic!("DX12 backend requires shader code in DXIL format"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Dx12) + } + Device::Mtl(d) => { + let shader_code = match code { + ShaderCode::Msl(msl) => msl, + // Panic or return "incompatible shader" error here? + _ => panic!("Metal backend requires shader code in MSL format"), + }; + d.create_compute_pipeline(shader_code, bind_types) + .map(Pipeline::Mtl) + } } } @@ -436,74 +478,21 @@ impl Device { &self, _spv: &'a [u8], _hlsl: &'a str, + _dxil: &'a [u8], _msl: &'a str, ) -> ShaderCode<'a> { mux_match! { self; Device::Vk(_d) => ShaderCode::Spv(_spv), - Device::Dx12(_d) => ShaderCode::Hlsl(_hlsl), + Device::Dx12(_d) => ShaderCode::Dxil(_dxil), Device::Mtl(_d) => ShaderCode::Msl(_msl), } } -} -impl PipelineBuilder { - pub fn add_buffers(&mut self, n_buffers: u32) { + pub fn backend_type(&self) -> BackendType { mux_match! { self; - PipelineBuilder::Vk(x) => x.add_buffers(n_buffers), - PipelineBuilder::Dx12(x) => x.add_buffers(n_buffers), - PipelineBuilder::Mtl(x) => x.add_buffers(n_buffers), - } - } - - pub fn add_images(&mut self, n_buffers: u32) { - mux_match! { self; - PipelineBuilder::Vk(x) => x.add_images(n_buffers), - PipelineBuilder::Dx12(x) => x.add_images(n_buffers), - PipelineBuilder::Mtl(x) => x.add_images(n_buffers), - } - } - - pub fn add_textures(&mut self, n_buffers: u32) { - mux_match! { self; - PipelineBuilder::Vk(x) => x.add_textures(n_buffers), - PipelineBuilder::Dx12(x) => x.add_textures(n_buffers), - PipelineBuilder::Mtl(x) => x.add_textures(n_buffers), - } - } - - pub unsafe fn create_compute_pipeline<'a>( - self, - device: &Device, - code: ShaderCode<'a>, - ) -> Result { - mux_match! { self; - PipelineBuilder::Vk(x) => { - let shader_code = match code { - ShaderCode::Spv(spv) => spv, - // Panic or return "incompatible shader" error here? - _ => panic!("Vulkan backend requires shader code in SPIR-V format"), - }; - x.create_compute_pipeline(device.vk(), shader_code) - .map(Pipeline::Vk) - } - PipelineBuilder::Dx12(x) => { - let shader_code = match code { - ShaderCode::Hlsl(hlsl) => hlsl, - // Panic or return "incompatible shader" error here? - _ => panic!("DX12 backend requires shader code in HLSL format"), - }; - x.create_compute_pipeline(device.dx12(), shader_code) - .map(Pipeline::Dx12) - } - PipelineBuilder::Mtl(x) => { - let shader_code = match code { - ShaderCode::Msl(msl) => msl, - // Panic or return "incompatible shader" error here? - _ => panic!("Metal backend requires shader code in MSL format"), - }; - x.create_compute_pipeline(device.mtl(), shader_code) - .map(Pipeline::Mtl) - } + Device::Vk(_d) => BackendType::Vulkan, + Device::Dx12(_d) => BackendType::Dx12, + Device::Mtl(_d) => BackendType::Metal, } } } diff --git a/piet-gpu-hal/src/vulkan.rs b/piet-gpu-hal/src/vulkan.rs index 26e095f..34b6109 100644 --- a/piet-gpu-hal/src/vulkan.rs +++ b/piet-gpu-hal/src/vulkan.rs @@ -11,9 +11,11 @@ use ash::{vk, Device, Entry, Instance}; use smallvec::SmallVec; -use crate::{BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams, SubgroupSize, WorkgroupLimits}; use crate::backend::Device as DeviceTrait; - +use crate::{ + BindType, BufferUsage, Error, GpuInfo, ImageLayout, SamplerParams, SubgroupSize, + WorkgroupLimits, +}; pub struct VkInstance { /// Retain the dynamic lib. @@ -98,12 +100,6 @@ pub struct QueryPool { #[derive(Clone, Copy)] pub struct MemFlags(vk::MemoryPropertyFlags); -pub struct PipelineBuilder { - bindings: Vec, - binding_flags: Vec, - max_textures: u32, -} - pub struct DescriptorSetBuilder { buffers: Vec, images: Vec, @@ -262,9 +258,9 @@ impl VkInstance { if vk1_1 { let mut descriptor_indexing_features = vk::PhysicalDeviceDescriptorIndexingFeatures::builder(); - features2 = features2 - .push_next(&mut descriptor_indexing_features); - self.instance.get_physical_device_features2(pdevice, &mut features2); + features2 = features2.push_next(&mut descriptor_indexing_features); + self.instance + .get_physical_device_features2(pdevice, &mut features2); set_features2 = set_features2.features(features2.features); has_descriptor_indexing = descriptor_indexing_features .shader_storage_image_array_non_uniform_indexing @@ -296,14 +292,13 @@ impl VkInstance { extensions.try_add(vk::KhrMaintenance3Fn::name()); extensions.try_add(vk::ExtDescriptorIndexingFn::name()); } - let has_subgroup_size = vk1_1 - && extensions.try_add(vk::ExtSubgroupSizeControlFn::name()); - let has_memory_model = vk1_1 - && extensions.try_add(vk::KhrVulkanMemoryModelFn::name()); + let has_subgroup_size = vk1_1 && extensions.try_add(vk::ExtSubgroupSizeControlFn::name()); + let has_memory_model = vk1_1 && extensions.try_add(vk::KhrVulkanMemoryModelFn::name()); let mut create_info = vk::DeviceCreateInfo::builder() .queue_create_infos(&queue_create_infos) .enabled_extension_names(extensions.as_ptrs()); - let mut set_memory_model_features = vk::PhysicalDeviceVulkanMemoryModelFeatures::builder(); if vk1_1 { + let mut set_memory_model_features = vk::PhysicalDeviceVulkanMemoryModelFeatures::builder(); + if vk1_1 { create_info = create_info.push_next(&mut set_features2); if has_memory_model { set_memory_model_features = set_memory_model_features @@ -422,7 +417,8 @@ impl VkInstance { 0 => u32::MAX, x => x, }; - let image_count = PREFERRED_IMAGE_COUNT.clamp(capabilities.min_image_count, max_image_count); + let image_count = + PREFERRED_IMAGE_COUNT.clamp(capabilities.min_image_count, max_image_count); let mut extent = capabilities.current_extent; if extent.width == u32::MAX || extent.height == u32::MAX { // We're deciding the size. @@ -475,7 +471,6 @@ impl crate::backend::Device for VkDevice { type QueryPool = QueryPool; type Fence = vk::Fence; type Semaphore = vk::Semaphore; - type PipelineBuilder = PipelineBuilder; type DescriptorSetBuilder = DescriptorSetBuilder; type Sampler = vk::Sampler; type ShaderSource = [u8]; @@ -649,12 +644,65 @@ impl crate::backend::Device for VkDevice { Ok(device.get_fence_status(*fence)?) } - unsafe fn pipeline_builder(&self) -> PipelineBuilder { - PipelineBuilder { - bindings: Vec::new(), - binding_flags: Vec::new(), - max_textures: 0, - } + unsafe fn create_compute_pipeline( + &self, + code: &[u8], + bind_types: &[BindType], + ) -> Result { + let device = &self.device.device; + let bindings = bind_types + .iter() + .enumerate() + .map(|(i, bind_type)| { + let descriptor_type = match bind_type { + BindType::Buffer | BindType::BufReadOnly => vk::DescriptorType::STORAGE_BUFFER, + BindType::Image | BindType::ImageRead => vk::DescriptorType::STORAGE_IMAGE, + }; + vk::DescriptorSetLayoutBinding::builder() + .binding(i.try_into().unwrap()) + .descriptor_type(descriptor_type) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::COMPUTE) + .build() + }) + .collect::>(); + let descriptor_set_layout = device.create_descriptor_set_layout( + &vk::DescriptorSetLayoutCreateInfo::builder().bindings(&bindings), + None, + )?; + let descriptor_set_layouts = [descriptor_set_layout]; + + // Create compute pipeline. + let code_u32 = convert_u32_vec(code); + let compute_shader_module = device + .create_shader_module(&vk::ShaderModuleCreateInfo::builder().code(&code_u32), None)?; + let entry_name = CString::new("main").unwrap(); + let pipeline_layout = device.create_pipeline_layout( + &vk::PipelineLayoutCreateInfo::builder().set_layouts(&descriptor_set_layouts), + None, + )?; + + let pipeline = device + .create_compute_pipelines( + vk::PipelineCache::null(), + &[vk::ComputePipelineCreateInfo::builder() + .stage( + vk::PipelineShaderStageCreateInfo::builder() + .stage(vk::ShaderStageFlags::COMPUTE) + .module(compute_shader_module) + .name(&entry_name) + .build(), + ) + .layout(pipeline_layout) + .build()], + None, + ) + .map_err(|(_pipeline, err)| err)?[0]; + Ok(Pipeline { + pipeline, + pipeline_layout, + descriptor_set_layout, + }) } unsafe fn descriptor_set_builder(&self) -> DescriptorSetBuilder { @@ -715,13 +763,7 @@ impl crate::backend::Device for VkDevice { // fence should make the query available, but otherwise we get sporadic NOT_READY // results (Windows 10, AMD 5700 XT). let flags = vk::QueryResultFlags::TYPE_64 | vk::QueryResultFlags::WAIT; - device.get_query_pool_results( - pool.pool, - 0, - pool.n_queries, - &mut buf, - flags, - )?; + device.get_query_pool_results(pool.pool, 0, pool.n_queries, &mut buf, flags)?; let ts0 = buf[0]; let tsp = self.timestamp_period as f64 * 1e-9; let result = buf[1..] @@ -1080,109 +1122,6 @@ impl crate::backend::CmdBuf for CmdBuf { } } -impl crate::backend::PipelineBuilder for PipelineBuilder { - fn add_buffers(&mut self, n_buffers: u32) { - let start = self.bindings.len() as u32; - for i in 0..n_buffers { - self.bindings.push( - vk::DescriptorSetLayoutBinding::builder() - .binding(start + i) - .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) - .descriptor_count(1) - .stage_flags(vk::ShaderStageFlags::COMPUTE) - .build(), - ); - self.binding_flags - .push(vk::DescriptorBindingFlags::default()); - } - } - - fn add_images(&mut self, n_images: u32) { - let start = self.bindings.len() as u32; - for i in 0..n_images { - self.bindings.push( - vk::DescriptorSetLayoutBinding::builder() - .binding(start + i) - .descriptor_type(vk::DescriptorType::STORAGE_IMAGE) - .descriptor_count(1) - .stage_flags(vk::ShaderStageFlags::COMPUTE) - .build(), - ); - self.binding_flags - .push(vk::DescriptorBindingFlags::default()); - } - } - - fn add_textures(&mut self, n_images: u32) { - let start = self.bindings.len() as u32; - for i in 0..n_images { - self.bindings.push( - vk::DescriptorSetLayoutBinding::builder() - .binding(start + i) - .descriptor_type(vk::DescriptorType::STORAGE_IMAGE) - .descriptor_count(1) - .stage_flags(vk::ShaderStageFlags::COMPUTE) - .build(), - ); - self.binding_flags - .push(vk::DescriptorBindingFlags::default()); - } - self.max_textures += n_images; - } - - unsafe fn create_compute_pipeline( - self, - device: &VkDevice, - code: &[u8], - ) -> Result { - let device = &device.device.device; - let descriptor_set_layout = device.create_descriptor_set_layout( - &vk::DescriptorSetLayoutCreateInfo::builder() - .bindings(&self.bindings) - // It might be a slight optimization not to push this if max_textures = 0 - .push_next( - &mut vk::DescriptorSetLayoutBindingFlagsCreateInfo::builder() - .binding_flags(&self.binding_flags) - .build(), - ), - None, - )?; - let descriptor_set_layouts = [descriptor_set_layout]; - - // Create compute pipeline. - let code_u32 = convert_u32_vec(code); - let compute_shader_module = device - .create_shader_module(&vk::ShaderModuleCreateInfo::builder().code(&code_u32), None)?; - let entry_name = CString::new("main").unwrap(); - let pipeline_layout = device.create_pipeline_layout( - &vk::PipelineLayoutCreateInfo::builder().set_layouts(&descriptor_set_layouts), - None, - )?; - - let pipeline = device - .create_compute_pipelines( - vk::PipelineCache::null(), - &[vk::ComputePipelineCreateInfo::builder() - .stage( - vk::PipelineShaderStageCreateInfo::builder() - .stage(vk::ShaderStageFlags::COMPUTE) - .module(compute_shader_module) - .name(&entry_name) - .build(), - ) - .layout(pipeline_layout) - .build()], - None, - ) - .map_err(|(_pipeline, err)| err)?[0]; - Ok(Pipeline { - pipeline, - pipeline_layout, - descriptor_set_layout, - }) - } -} - impl crate::backend::DescriptorSetBuilder for DescriptorSetBuilder { fn add_buffers(&mut self, buffers: &[&Buffer]) { self.buffers.extend(buffers.iter().map(|b| b.buffer)); diff --git a/piet-gpu/Cargo.toml b/piet-gpu/Cargo.toml index cc9684a..f8f5c0a 100644 --- a/piet-gpu/Cargo.toml +++ b/piet-gpu/Cargo.toml @@ -30,7 +30,7 @@ piet = "0.2.0" png = "0.16.2" rand = "0.7.3" roxmltree = "0.13" -winit = "0.23" +winit = "0.25" clap = "2.33" swash = "0.1.4" diff --git a/piet-gpu/bin/android.rs b/piet-gpu/bin/android.rs index eb7fb02..8254cc0 100644 --- a/piet-gpu/bin/android.rs +++ b/piet-gpu/bin/android.rs @@ -56,7 +56,7 @@ fn my_main() -> Result<(), Error> { let width = window.width() as usize; let height = window.height() as usize; let handle = get_handle(window); - let (instance, surface) = Instance::new(Some(&handle))?; + let (instance, surface) = Instance::new(Some(&handle), Default::default())?; gfx_state = Some(GfxState::new(&instance, surface.as_ref(), width, height)?); } else { diff --git a/piet-gpu/bin/cli.rs b/piet-gpu/bin/cli.rs index 837bd55..c48f65f 100644 --- a/piet-gpu/bin/cli.rs +++ b/piet-gpu/bin/cli.rs @@ -226,7 +226,7 @@ fn main() -> Result<(), Error> { .takes_value(true), ) .get_matches(); - let (instance, _) = Instance::new(None)?; + let (instance, _) = Instance::new(None, Default::default())?; unsafe { let device = instance.device(None)?; let session = Session::new(device); diff --git a/piet-gpu/bin/winit.rs b/piet-gpu/bin/winit.rs index ef41b31..bff0f70 100644 --- a/piet-gpu/bin/winit.rs +++ b/piet-gpu/bin/winit.rs @@ -38,7 +38,7 @@ fn main() -> Result<(), Error> { .with_resizable(false) // currently not supported .build(&event_loop)?; - let (instance, surface) = Instance::new(Some(&window))?; + let (instance, surface) = Instance::new(Some(&window), Default::default())?; let mut info_string = "info".to_string(); unsafe { let device = instance.device(surface.as_ref())?; diff --git a/piet-gpu/src/lib.rs b/piet-gpu/src/lib.rs index 30fcf8f..bee07aa 100644 --- a/piet-gpu/src/lib.rs +++ b/piet-gpu/src/lib.rs @@ -14,8 +14,8 @@ use piet::{ImageFormat, RenderContext}; use piet_gpu_types::encoder::Encode; use piet_gpu_hal::{ - Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline, QueryPool, - Session, ShaderCode, + BindType, Buffer, BufferUsage, CmdBuf, DescriptorSet, Error, Image, ImageLayout, Pipeline, + QueryPool, Session, ShaderCode, }; use pico_svg::PicoSvg; @@ -140,7 +140,15 @@ impl Renderer { let memory_buf_dev = session.create_buffer(128 * 1024 * 1024, dev)?; let el_code = ShaderCode::Spv(include_bytes!("../shader/elements.spv")); - let el_pipeline = session.create_simple_compute_pipeline(el_code, 4)?; + let el_pipeline = session.create_compute_pipeline( + el_code, + &[ + BindType::Buffer, + BindType::Buffer, + BindType::Buffer, + BindType::Buffer, + ], + )?; let mut el_ds = Vec::with_capacity(n_bufs); for scene_buf in &scene_bufs { el_ds.push(session.create_simple_descriptor_set( @@ -150,12 +158,14 @@ impl Renderer { } let tile_alloc_code = ShaderCode::Spv(include_bytes!("../shader/tile_alloc.spv")); - let tile_pipeline = session.create_simple_compute_pipeline(tile_alloc_code, 2)?; + let tile_pipeline = session + .create_compute_pipeline(tile_alloc_code, &[BindType::Buffer, BindType::Buffer])?; let tile_ds = session .create_simple_descriptor_set(&tile_pipeline, &[&memory_buf_dev, &config_buf])?; let path_alloc_code = ShaderCode::Spv(include_bytes!("../shader/path_coarse.spv")); - let path_pipeline = session.create_simple_compute_pipeline(path_alloc_code, 2)?; + let path_pipeline = session + .create_compute_pipeline(path_alloc_code, &[BindType::Buffer, BindType::Buffer])?; let path_ds = session .create_simple_descriptor_set(&path_pipeline, &[&memory_buf_dev, &config_buf])?; @@ -165,18 +175,21 @@ impl Renderer { println!("using small workgroup backdrop kernel"); ShaderCode::Spv(include_bytes!("../shader/backdrop.spv")) }; - let backdrop_pipeline = session.create_simple_compute_pipeline(backdrop_code, 2)?; + let backdrop_pipeline = session + .create_compute_pipeline(backdrop_code, &[BindType::Buffer, BindType::Buffer])?; let backdrop_ds = session .create_simple_descriptor_set(&backdrop_pipeline, &[&memory_buf_dev, &config_buf])?; // TODO: constants let bin_code = ShaderCode::Spv(include_bytes!("../shader/binning.spv")); - let bin_pipeline = session.create_simple_compute_pipeline(bin_code, 2)?; + let bin_pipeline = + session.create_compute_pipeline(bin_code, &[BindType::Buffer, BindType::Buffer])?; let bin_ds = session.create_simple_descriptor_set(&bin_pipeline, &[&memory_buf_dev, &config_buf])?; let coarse_code = ShaderCode::Spv(include_bytes!("../shader/coarse.spv")); - let coarse_pipeline = session.create_simple_compute_pipeline(coarse_code, 2)?; + let coarse_pipeline = + session.create_compute_pipeline(coarse_code, &[BindType::Buffer, BindType::Buffer])?; let coarse_ds = session .create_simple_descriptor_set(&coarse_pipeline, &[&memory_buf_dev, &config_buf])?; @@ -194,17 +207,16 @@ impl Renderer { let gradients = Self::make_gradient_image(&session); let k4_code = ShaderCode::Spv(include_bytes!("../shader/kernel4.spv")); - // This is a bit of a stand-in for future development. For now, we assume one - // atlas image for all images, and another image for the gradients. In the future, - // on GPUs that support it, we will probably want to go to descriptor indexing in - // order to cut down on allocation and copying for the atlas image. - let max_textures = 2; - let k4_pipeline = session - .pipeline_builder() - .add_buffers(2) - .add_images(1) - .add_textures(max_textures) - .create_compute_pipeline(&session, k4_code)?; + let k4_pipeline = session.create_compute_pipeline( + k4_code, + &[ + BindType::Buffer, + BindType::Buffer, + BindType::Image, + BindType::ImageRead, + BindType::ImageRead, + ], + )?; let k4_ds = session .descriptor_set_builder() .add_buffers(&[&memory_buf_dev, &config_buf]) diff --git a/tests/shader/build.ninja b/tests/shader/build.ninja index 93a0b66..19297c9 100644 --- a/tests/shader/build.ninja +++ b/tests/shader/build.ninja @@ -4,6 +4,11 @@ glslang_validator = glslangValidator spirv_cross = spirv-cross +dxc = dxc + +# See https://github.com/KhronosGroup/SPIRV-Cross/issues/1248 for +# why we set this. +msl_flags = --msl-decoration-binding rule glsl command = $glslang_validator $flags -V -o $out $in @@ -11,22 +16,34 @@ rule glsl rule hlsl command = $spirv_cross --hlsl $in --output $out +rule dxil + command = $dxc -T cs_6_0 $in -Fo $out + rule msl - command = $spirv_cross --msl $in --output $out + command = $spirv_cross --msl $in --output $out $msl_flags + +build gen/clear.spv: glsl clear.comp +build gen/clear.hlsl: hlsl gen/clear.spv +build gen/clear.dxil: dxil gen/clear.hlsl +build gen/clear.msl: msl gen/clear.spv build gen/prefix.spv: glsl prefix.comp build gen/prefix.hlsl: hlsl gen/prefix.spv +build gen/prefix.dxil: dxil gen/prefix.hlsl build gen/prefix.msl: msl gen/prefix.spv build gen/prefix_reduce.spv: glsl prefix_reduce.comp build gen/prefix_reduce.hlsl: hlsl gen/prefix_reduce.spv +build gen/prefix_reduce.dxil: dxil gen/prefix_reduce.hlsl build gen/prefix_reduce.msl: msl gen/prefix_reduce.spv build gen/prefix_root.spv: glsl prefix_scan.comp flags = -DROOT build gen/prefix_root.hlsl: hlsl gen/prefix_root.spv +build gen/prefix_root.dxil: dxil gen/prefix_root.hlsl build gen/prefix_root.msl: msl gen/prefix_root.spv build gen/prefix_scan.spv: glsl prefix_scan.comp build gen/prefix_scan.hlsl: hlsl gen/prefix_scan.spv +build gen/prefix_scan.dxil: dxil gen/prefix_scan.hlsl build gen/prefix_scan.msl: msl gen/prefix_scan.spv diff --git a/tests/shader/clear.comp b/tests/shader/clear.comp new file mode 100644 index 0000000..62a5fb2 --- /dev/null +++ b/tests/shader/clear.comp @@ -0,0 +1,26 @@ +// SPDX-License-Identifier: Apache-2.0 OR MIT OR Unlicense + +// Clear a buffer. + +#version 450 + +layout(local_size_x = 256) in; + +// This should probably be uniform rather than readonly, +// but we haven't done the binding work yet. +layout(binding = 0) readonly buffer ConfigBuf { + // size is in uint (4 byte) units + uint size; + uint value; +}; + +layout(binding = 1) buffer TargetBuf { + uint[] data; +}; + +void main() { + uint ix = gl_GlobalInvocationID.x; + if (ix < size) { + data[ix] = value; + } +} diff --git a/tests/shader/gen/clear.dxil b/tests/shader/gen/clear.dxil new file mode 100644 index 0000000..a79182a Binary files /dev/null and b/tests/shader/gen/clear.dxil differ diff --git a/tests/shader/gen/clear.hlsl b/tests/shader/gen/clear.hlsl new file mode 100644 index 0000000..f6a576c --- /dev/null +++ b/tests/shader/gen/clear.hlsl @@ -0,0 +1,26 @@ +static const uint3 gl_WorkGroupSize = uint3(256u, 1u, 1u); + +ByteAddressBuffer _19 : register(t0); +RWByteAddressBuffer _32 : register(u1); + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + uint ix = gl_GlobalInvocationID.x; + if (ix < _19.Load(0)) + { + _32.Store(ix * 4 + 0, _19.Load(4)); + } +} + +[numthreads(256, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/tests/shader/gen/clear.msl b/tests/shader/gen/clear.msl new file mode 100644 index 0000000..d89853b --- /dev/null +++ b/tests/shader/gen/clear.msl @@ -0,0 +1,27 @@ +#include +#include + +using namespace metal; + +struct ConfigBuf +{ + uint size; + uint value; +}; + +struct TargetBuf +{ + uint data[1]; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u); + +kernel void main0(const device ConfigBuf& _19 [[buffer(0)]], device TargetBuf& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint ix = gl_GlobalInvocationID.x; + if (ix < _19.size) + { + _32.data[ix] = _19.value; + } +} + diff --git a/tests/shader/gen/clear.spv b/tests/shader/gen/clear.spv new file mode 100644 index 0000000..0e8d1d7 Binary files /dev/null and b/tests/shader/gen/clear.spv differ diff --git a/tests/shader/gen/prefix.dxil b/tests/shader/gen/prefix.dxil new file mode 100644 index 0000000..34f3d6a Binary files /dev/null and b/tests/shader/gen/prefix.dxil differ diff --git a/tests/shader/gen/prefix.hlsl b/tests/shader/gen/prefix.hlsl index c0600e2..3af5a96 100644 --- a/tests/shader/gen/prefix.hlsl +++ b/tests/shader/gen/prefix.hlsl @@ -12,11 +12,11 @@ struct State static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _187 = { 0u }; +static const Monoid _185 = { 0u }; globallycoherent RWByteAddressBuffer _43 : register(u2); ByteAddressBuffer _67 : register(t0); -RWByteAddressBuffer _374 : register(u1); +RWByteAddressBuffer _372 : register(u1); static uint3 gl_LocalInvocationID; struct SPIRV_Cross_Input @@ -64,9 +64,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -92,7 +92,7 @@ void comp_main() } _43.Store(part_ix * 12 + 4, flag); } - Monoid exclusive = _187; + Monoid exclusive = _185; if (part_ix != 0u) { uint look_back_ix = part_ix - 1u; @@ -113,9 +113,9 @@ void comp_main() { if (gl_LocalInvocationID.x == 511u) { - Monoid _225; - _225.element = _43.Load(look_back_ix * 12 + 12); - their_prefix.element = _225.element; + Monoid _223; + _223.element = _43.Load(look_back_ix * 12 + 12); + their_prefix.element = _223.element; Monoid param_4 = their_prefix; Monoid param_5 = exclusive; exclusive = combine_monoid(param_4, param_5); @@ -128,9 +128,9 @@ void comp_main() { if (gl_LocalInvocationID.x == 511u) { - Monoid _247; - _247.element = _43.Load(look_back_ix * 12 + 8); - their_agg.element = _247.element; + Monoid _245; + _245.element = _43.Load(look_back_ix * 12 + 8); + their_agg.element = _245.element; Monoid param_6 = their_agg; Monoid param_7 = exclusive; exclusive = combine_monoid(param_6, param_7); @@ -142,9 +142,9 @@ void comp_main() } if (gl_LocalInvocationID.x == 511u) { - Monoid _269; - _269.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0); - m.element = _269.element; + Monoid _267; + _267.element = _67.Load(((look_back_ix * 8192u) + their_ix) * 4 + 0); + m.element = _267.element; if (their_ix == 0u) { their_agg = m; @@ -211,7 +211,7 @@ void comp_main() Monoid param_16 = row; Monoid param_17 = local[i_2]; Monoid m_1 = combine_monoid(param_16, param_17); - _374.Store((ix + i_2) * 4 + 0, m_1.element); + _372.Store((ix + i_2) * 4 + 0, m_1.element); } } diff --git a/tests/shader/gen/prefix.msl b/tests/shader/gen/prefix.msl index ecdf8bd..8e402a9 100644 --- a/tests/shader/gen/prefix.msl +++ b/tests/shader/gen/prefix.msl @@ -87,7 +87,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBuf& _67 [[buffer(1)]], device OutBuf& _374 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) +kernel void main0(const device InBuf& _67 [[buffer(0)]], device OutBuf& _372 [[buffer(1)]], volatile device StateBuf& _43 [[buffer(2)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { threadgroup uint sh_part_ix; threadgroup Monoid sh_scratch[512]; @@ -115,9 +115,9 @@ kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBu for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -256,7 +256,7 @@ kernel void main0(volatile device StateBuf& _43 [[buffer(0)]], const device InBu Monoid param_16 = row; Monoid param_17 = local[i_2]; Monoid m_1 = combine_monoid(param_16, param_17); - _374.outbuf[ix + i_2].element = m_1.element; + _372.outbuf[ix + i_2].element = m_1.element; } } diff --git a/tests/shader/gen/prefix.spv b/tests/shader/gen/prefix.spv index 170a569..b934189 100644 Binary files a/tests/shader/gen/prefix.spv and b/tests/shader/gen/prefix.spv differ diff --git a/tests/shader/gen/prefix_reduce.dxil b/tests/shader/gen/prefix_reduce.dxil new file mode 100644 index 0000000..0ee28e8 Binary files /dev/null and b/tests/shader/gen/prefix_reduce.dxil differ diff --git a/tests/shader/gen/prefix_reduce.hlsl b/tests/shader/gen/prefix_reduce.hlsl index 837a75a..f2de539 100644 --- a/tests/shader/gen/prefix_reduce.hlsl +++ b/tests/shader/gen/prefix_reduce.hlsl @@ -6,7 +6,7 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); ByteAddressBuffer _40 : register(t0); -RWByteAddressBuffer _129 : register(u1); +RWByteAddressBuffer _127 : register(u1); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -46,9 +46,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u) + if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)]; Monoid param_2 = agg; Monoid param_3 = other; agg = combine_monoid(param_2, param_3); @@ -58,7 +58,7 @@ void comp_main() } if (gl_LocalInvocationID.x == 0u) { - _129.Store(gl_WorkGroupID.x * 4 + 0, agg.element); + _127.Store(gl_WorkGroupID.x * 4 + 0, agg.element); } } diff --git a/tests/shader/gen/prefix_reduce.msl b/tests/shader/gen/prefix_reduce.msl index e1ed0ce..3a3125d 100644 --- a/tests/shader/gen/prefix_reduce.msl +++ b/tests/shader/gen/prefix_reduce.msl @@ -33,7 +33,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _127 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) { threadgroup Monoid sh_scratch[512]; uint ix = gl_GlobalInvocationID.x * 8u; @@ -50,9 +50,9 @@ kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[b for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u) + if ((gl_LocalInvocationID.x + (1u << i_1)) < 512u) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i_1)]; Monoid param_2 = agg; Monoid param_3 = other; agg = combine_monoid(param_2, param_3); @@ -62,7 +62,7 @@ kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[b } if (gl_LocalInvocationID.x == 0u) { - _129.outbuf[gl_WorkGroupID.x].element = agg.element; + _127.outbuf[gl_WorkGroupID.x].element = agg.element; } } diff --git a/tests/shader/gen/prefix_reduce.spv b/tests/shader/gen/prefix_reduce.spv index d1db3aa..b2e35fc 100644 Binary files a/tests/shader/gen/prefix_reduce.spv and b/tests/shader/gen/prefix_reduce.spv differ diff --git a/tests/shader/gen/prefix_root.dxil b/tests/shader/gen/prefix_root.dxil new file mode 100644 index 0000000..03fe2d1 Binary files /dev/null and b/tests/shader/gen/prefix_root.dxil differ diff --git a/tests/shader/gen/prefix_root.hlsl b/tests/shader/gen/prefix_root.hlsl index 2ad617c..adf6bf8 100644 --- a/tests/shader/gen/prefix_root.hlsl +++ b/tests/shader/gen/prefix_root.hlsl @@ -5,7 +5,7 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _133 = { 0u }; +static const Monoid _131 = { 0u }; RWByteAddressBuffer _42 : register(u0); @@ -46,9 +46,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -57,7 +57,7 @@ void comp_main() sh_scratch[gl_LocalInvocationID.x] = agg; } GroupMemoryBarrierWithGroupSync(); - Monoid row = _133; + Monoid row = _131; if (gl_LocalInvocationID.x > 0u) { row = sh_scratch[gl_LocalInvocationID.x - 1u]; diff --git a/tests/shader/gen/prefix_root.msl b/tests/shader/gen/prefix_root.msl index ff02287..897a6a4 100644 --- a/tests/shader/gen/prefix_root.msl +++ b/tests/shader/gen/prefix_root.msl @@ -85,9 +85,9 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], uint3 gl_GlobalInvocationID for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); diff --git a/tests/shader/gen/prefix_root.spv b/tests/shader/gen/prefix_root.spv index 70ba31c..3e04224 100644 Binary files a/tests/shader/gen/prefix_root.spv and b/tests/shader/gen/prefix_root.spv differ diff --git a/tests/shader/gen/prefix_scan.dxil b/tests/shader/gen/prefix_scan.dxil new file mode 100644 index 0000000..427f14d Binary files /dev/null and b/tests/shader/gen/prefix_scan.dxil differ diff --git a/tests/shader/gen/prefix_scan.hlsl b/tests/shader/gen/prefix_scan.hlsl index feeff2e..d9e74ea 100644 --- a/tests/shader/gen/prefix_scan.hlsl +++ b/tests/shader/gen/prefix_scan.hlsl @@ -5,10 +5,10 @@ struct Monoid static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u); -static const Monoid _133 = { 0u }; +static const Monoid _131 = { 0u }; RWByteAddressBuffer _42 : register(u0); -RWByteAddressBuffer _143 : register(u1); +ByteAddressBuffer _141 : register(t1); static uint3 gl_WorkGroupID; static uint3 gl_LocalInvocationID; @@ -49,9 +49,9 @@ void comp_main() for (uint i_1 = 0u; i_1 < 9u; i_1++) { GroupMemoryBarrierWithGroupSync(); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -60,12 +60,12 @@ void comp_main() sh_scratch[gl_LocalInvocationID.x] = agg; } GroupMemoryBarrierWithGroupSync(); - Monoid row = _133; + Monoid row = _131; if (gl_WorkGroupID.x > 0u) { - Monoid _148; - _148.element = _143.Load((gl_WorkGroupID.x - 1u) * 4 + 0); - row.element = _148.element; + Monoid _146; + _146.element = _141.Load((gl_WorkGroupID.x - 1u) * 4 + 0); + row.element = _146.element; } if (gl_LocalInvocationID.x > 0u) { diff --git a/tests/shader/gen/prefix_scan.msl b/tests/shader/gen/prefix_scan.msl index c1efb22..5be4e65 100644 --- a/tests/shader/gen/prefix_scan.msl +++ b/tests/shader/gen/prefix_scan.msl @@ -72,7 +72,7 @@ Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b) return Monoid{ a.element + b.element }; } -kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +kernel void main0(device DataBuf& _42 [[buffer(0)]], const device ParentBuf& _141 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) { threadgroup Monoid sh_scratch[512]; uint ix = gl_GlobalInvocationID.x * 8u; @@ -90,9 +90,9 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[bu for (uint i_1 = 0u; i_1 < 9u; i_1++) { threadgroup_barrier(mem_flags::mem_threadgroup); - if (gl_LocalInvocationID.x >= uint(1 << int(i_1))) + if (gl_LocalInvocationID.x >= (1u << i_1)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))]; + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i_1)]; Monoid param_2 = other; Monoid param_3 = agg; agg = combine_monoid(param_2, param_3); @@ -104,7 +104,7 @@ kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[bu Monoid row = Monoid{ 0u }; if (gl_WorkGroupID.x > 0u) { - row.element = _143.parent[gl_WorkGroupID.x - 1u].element; + row.element = _141.parent[gl_WorkGroupID.x - 1u].element; } if (gl_LocalInvocationID.x > 0u) { diff --git a/tests/shader/gen/prefix_scan.spv b/tests/shader/gen/prefix_scan.spv index d4216e9..6d8fe0a 100644 Binary files a/tests/shader/gen/prefix_scan.spv and b/tests/shader/gen/prefix_scan.spv differ diff --git a/tests/shader/prefix.comp b/tests/shader/prefix.comp index ed5bcbc..3ca1509 100644 --- a/tests/shader/prefix.comp +++ b/tests/shader/prefix.comp @@ -71,8 +71,8 @@ void main() { sh_scratch[gl_LocalInvocationID.x] = agg; for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); - if (gl_LocalInvocationID.x >= (1 << i)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - (1 << i)]; + if (gl_LocalInvocationID.x >= (1u << i)) { + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i)]; agg = combine_monoid(other, agg); } barrier(); diff --git a/tests/shader/prefix_reduce.comp b/tests/shader/prefix_reduce.comp index 378da88..36750e9 100644 --- a/tests/shader/prefix_reduce.comp +++ b/tests/shader/prefix_reduce.comp @@ -40,8 +40,8 @@ void main() { for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); // We could make this predicate tighter, but would it help? - if (gl_LocalInvocationID.x + (1 << i) < WG_SIZE) { - Monoid other = sh_scratch[gl_LocalInvocationID.x + (1 << i)]; + if (gl_LocalInvocationID.x + (1u << i) < WG_SIZE) { + Monoid other = sh_scratch[gl_LocalInvocationID.x + (1u << i)]; agg = combine_monoid(agg, other); } barrier(); diff --git a/tests/shader/prefix_scan.comp b/tests/shader/prefix_scan.comp index 59903ab..82ac847 100644 --- a/tests/shader/prefix_scan.comp +++ b/tests/shader/prefix_scan.comp @@ -20,7 +20,7 @@ layout(set = 0, binding = 0) buffer DataBuf { }; #ifndef ROOT -layout(set = 0, binding = 1) buffer ParentBuf { +layout(set = 0, binding = 1) readonly buffer ParentBuf { Monoid[] parent; }; #endif @@ -45,8 +45,8 @@ void main() { sh_scratch[gl_LocalInvocationID.x] = agg; for (uint i = 0; i < LG_WG_SIZE; i++) { barrier(); - if (gl_LocalInvocationID.x >= (1 << i)) { - Monoid other = sh_scratch[gl_LocalInvocationID.x - (1 << i)]; + if (gl_LocalInvocationID.x >= (1u << i)) { + Monoid other = sh_scratch[gl_LocalInvocationID.x - (1u << i)]; agg = combine_monoid(other, agg); } barrier(); diff --git a/tests/src/clear.rs b/tests/src/clear.rs new file mode 100644 index 0000000..f691928 --- /dev/null +++ b/tests/src/clear.rs @@ -0,0 +1,141 @@ +// Copyright 2021 The piet-gpu authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// https://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. +// +// Also licensed under MIT license, at your choice. + +//! Utilities (and a benchmark) for clearing buffers with compute shaders. + +use piet_gpu_hal::{include_shader, BindType, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{Buffer, Pipeline}; + +use crate::config::Config; +use crate::runner::{Commands, Runner}; +use crate::test_result::TestResult; + +const WG_SIZE: u64 = 256; + +/// The shader code for clearing buffers. +pub struct ClearCode { + pipeline: Pipeline, +} + +/// The stage resources for clearing buffers. +pub struct ClearStage { + n_elements: u64, + config_buf: Buffer, +} + +/// The binding for clearing buffers. +pub struct ClearBinding { + descriptor_set: DescriptorSet, +} + +pub unsafe fn run_clear_test(runner: &mut Runner, config: &Config) -> TestResult { + let mut result = TestResult::new("clear buffers"); + let n_elements: u64 = config.size.choose(1 << 12, 1 << 20, 1 << 24); + let out_buf = runner.buf_down(n_elements * 4); + let code = ClearCode::new(runner); + let stage = ClearStage::new_with_value(runner, n_elements, 0x42); + let binding = stage.bind(runner, &code, &out_buf.dev_buf); + let n_iter = config.n_iter; + let mut total_elapsed = 0.0; + for i in 0..n_iter { + let mut commands = runner.commands(); + commands.write_timestamp(0); + stage.record(&mut commands, &code, &binding); + commands.write_timestamp(1); + if i == 0 { + commands.cmd_buf.memory_barrier(); + commands.download(&out_buf); + } + total_elapsed += runner.submit(commands); + if i == 0 { + let mut dst: Vec = Default::default(); + out_buf.read(&mut dst); + if let Some(failure) = verify(&dst) { + result.fail(format!("failure at {}", failure)); + } + } + } + result.timing(total_elapsed, n_elements * n_iter); + result +} + +impl ClearCode { + pub unsafe fn new(runner: &mut Runner) -> ClearCode { + let code = include_shader!(&runner.session, "../shader/gen/Clear"); + let pipeline = runner + .session + .create_compute_pipeline(code, &[BindType::BufReadOnly, BindType::Buffer]) + .unwrap(); + ClearCode { pipeline } + } +} + +impl ClearStage { + pub unsafe fn new(runner: &mut Runner, n_elements: u64) -> ClearStage { + Self::new_with_value(runner, n_elements, 0) + } + + pub unsafe fn new_with_value(runner: &mut Runner, n_elements: u64, value: u32) -> ClearStage { + let config = [n_elements as u32, value]; + let config_buf = runner + .session + .create_buffer_init(&config, BufferUsage::STORAGE) + .unwrap(); + ClearStage { + n_elements, + config_buf, + } + } + + pub unsafe fn bind( + &self, + runner: &mut Runner, + code: &ClearCode, + out_buf: &Buffer, + ) -> ClearBinding { + let descriptor_set = runner + .session + .create_simple_descriptor_set(&code.pipeline, &[&self.config_buf, out_buf]) + .unwrap(); + ClearBinding { descriptor_set } + } + + pub unsafe fn record( + &self, + commands: &mut Commands, + code: &ClearCode, + bindings: &ClearBinding, + ) { + let n_workgroups = (self.n_elements + WG_SIZE - 1) / WG_SIZE; + // An issue: for clearing large buffers (>16M), we need to check the + // number of workgroups against the (dynamically detected) limit, and + // potentially issue multiple dispatches. + commands.cmd_buf.dispatch( + &code.pipeline, + &bindings.descriptor_set, + (n_workgroups as u32, 1, 1), + (WG_SIZE as u32, 1, 1), + ); + // One thing that's missing here is registering the buffers so + // they can be safely dropped by Rust code before the execution + // of the command buffer completes. + } +} + +// Verify that the data is cleared. +fn verify(data: &[u32]) -> Option { + data.iter().position(|val| *val != 0x42) +} diff --git a/tests/src/config.rs b/tests/src/config.rs index 50bd3be..edc1140 100644 --- a/tests/src/config.rs +++ b/tests/src/config.rs @@ -21,6 +21,7 @@ use clap::ArgMatches; pub struct Config { pub groups: Groups, pub size: Size, + pub n_iter: u64, } pub struct Groups(String); @@ -35,8 +36,14 @@ impl Config { pub fn from_matches(matches: &ArgMatches) -> Config { let groups = Groups::from_str(matches.value_of("groups").unwrap_or("all")); let size = Size::from_str(matches.value_of("size").unwrap_or("m")); + let n_iter = matches + .value_of("n_iter") + .and_then(|s| s.parse().ok()) + .unwrap_or(1000); Config { - groups, size + groups, + size, + n_iter, } } } diff --git a/tests/src/main.rs b/tests/src/main.rs index b7bc1d9..adefa7f 100644 --- a/tests/src/main.rs +++ b/tests/src/main.rs @@ -16,6 +16,7 @@ //! Tests for piet-gpu shaders and GPU capabilities. +mod clear; mod config; mod prefix; mod prefix_tree; @@ -23,6 +24,7 @@ mod runner; mod test_result; use clap::{App, Arg}; +use piet_gpu_hal::InstanceFlags; use crate::config::Config; use crate::runner::Runner; @@ -41,21 +43,26 @@ fn main() { .short("g") .long("groups") .help("Groups to run") - .takes_value(true) + .takes_value(true), ) .arg( Arg::with_name("size") .short("s") .long("size") .help("Size of tests") - .takes_value(true) + .takes_value(true), ) .arg( Arg::with_name("n_iter") .short("n") .long("n_iter") .help("Number of iterations") - .takes_value(true) + .takes_value(true), + ) + .arg( + Arg::with_name("dx12") + .long("dx12") + .help("Prefer DX12 backend"), ) .get_matches(); let style = if matches.is_present("verbose") { @@ -68,7 +75,16 @@ fn main() { let report = |test_result: &TestResult| { test_result.report(style); }; - let mut runner = Runner::new(); + let mut flags = InstanceFlags::empty(); + if matches.is_present("dx12") { + flags |= InstanceFlags::DX12; + } + let mut runner = Runner::new(flags); + if style == ReportStyle::Verbose { + // TODO: get adapter name in here too + println!("Backend: {:?}", runner.backend_type()); + } + report(&clear::run_clear_test(&mut runner, &config)); if config.groups.matches("prefix") { report(&prefix::run_prefix_test(&mut runner, &config)); report(&prefix_tree::run_prefix_test(&mut runner, &config)); diff --git a/tests/src/prefix.rs b/tests/src/prefix.rs index adc58b4..a2e52c3 100644 --- a/tests/src/prefix.rs +++ b/tests/src/prefix.rs @@ -14,9 +14,10 @@ // // Also licensed under MIT license, at your choice. -use piet_gpu_hal::{include_shader, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{include_shader, BackendType, BindType, BufferUsage, DescriptorSet}; use piet_gpu_hal::{Buffer, Pipeline}; +use crate::clear::{ClearBinding, ClearCode, ClearStage}; use crate::config::Config; use crate::runner::{Commands, Runner}; use crate::test_result::TestResult; @@ -30,6 +31,7 @@ const ELEMENTS_PER_WG: u64 = WG_SIZE * N_ROWS; /// A code struct can be created once and reused any number of times. struct PrefixCode { pipeline: Pipeline, + clear_code: Option, } /// The stage resources for the prefix sum example. @@ -41,6 +43,7 @@ struct PrefixStage { // treat it as a capacity. n_elements: u64, state_buf: Buffer, + clear_stage: Option<(ClearStage, ClearBinding)>, } /// The binding for the prefix sum example. @@ -50,7 +53,13 @@ struct PrefixBinding { pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResult { let mut result = TestResult::new("prefix sum, decoupled look-back"); - // This will be configurable. + /* + // We're good if we're using DXC. + if runner.backend_type() == BackendType::Dx12 { + result.skip("Shader won't compile on FXC"); + return result; + } + */ let n_elements: u64 = config.size.choose(1 << 12, 1 << 24, 1 << 25); let data: Vec = (0..n_elements as u32).collect(); let data_buf = runner @@ -59,10 +68,9 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul .unwrap(); let out_buf = runner.buf_down(data_buf.size()); let code = PrefixCode::new(runner); - let stage = PrefixStage::new(runner, n_elements); + let stage = PrefixStage::new(runner, &code, n_elements); let binding = stage.bind(runner, &code, &data_buf, &out_buf.dev_buf); - // Also will be configurable of course. - let n_iter = 1000; + let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { let mut commands = runner.commands(); @@ -91,23 +99,44 @@ impl PrefixCode { let code = include_shader!(&runner.session, "../shader/gen/prefix"); let pipeline = runner .session - .create_simple_compute_pipeline(code, 3) + .create_compute_pipeline( + code, + &[BindType::BufReadOnly, BindType::Buffer, BindType::Buffer], + ) .unwrap(); - PrefixCode { pipeline } + // Currently, DX12 and Metal backends don't support buffer clearing, so use a + // compute shader as a workaround. + let clear_code = if runner.backend_type() != BackendType::Vulkan { + Some(ClearCode::new(runner)) + } else { + None + }; + PrefixCode { + pipeline, + clear_code, + } } } impl PrefixStage { - unsafe fn new(runner: &mut Runner, n_elements: u64) -> PrefixStage { + unsafe fn new(runner: &mut Runner, code: &PrefixCode, n_elements: u64) -> PrefixStage { let n_workgroups = (n_elements + ELEMENTS_PER_WG - 1) / ELEMENTS_PER_WG; let state_buf_size = 4 + 12 * n_workgroups; let state_buf = runner .session .create_buffer(state_buf_size, BufferUsage::STORAGE | BufferUsage::COPY_DST) .unwrap(); + let clear_stage = if let Some(clear_code) = &code.clear_code { + let stage = ClearStage::new(runner, state_buf_size / 4); + let binding = stage.bind(runner, clear_code, &state_buf); + Some((stage, binding)) + } else { + None + }; PrefixStage { n_elements, state_buf, + clear_stage, } } @@ -127,7 +156,11 @@ impl PrefixStage { unsafe fn record(&self, commands: &mut Commands, code: &PrefixCode, bindings: &PrefixBinding) { let n_workgroups = (self.n_elements + ELEMENTS_PER_WG - 1) / ELEMENTS_PER_WG; - commands.cmd_buf.clear_buffer(&self.state_buf, None); + if let Some((stage, binding)) = &self.clear_stage { + stage.record(commands, code.clear_code.as_ref().unwrap(), binding); + } else { + commands.cmd_buf.clear_buffer(&self.state_buf, None); + } commands.cmd_buf.memory_barrier(); commands.cmd_buf.dispatch( &code.pipeline, diff --git a/tests/src/prefix_tree.rs b/tests/src/prefix_tree.rs index 1f78202..80a332f 100644 --- a/tests/src/prefix_tree.rs +++ b/tests/src/prefix_tree.rs @@ -14,7 +14,7 @@ // // Also licensed under MIT license, at your choice. -use piet_gpu_hal::{include_shader, BufferUsage, DescriptorSet}; +use piet_gpu_hal::{include_shader, BindType, BufferUsage, DescriptorSet}; use piet_gpu_hal::{Buffer, Pipeline}; use crate::config::Config; @@ -57,7 +57,7 @@ pub unsafe fn run_prefix_test(runner: &mut Runner, config: &Config) -> TestResul let stage = PrefixTreeStage::new(runner, n_elements); let binding = stage.bind(runner, &code, &out_buf.dev_buf); // Also will be configurable of course. - let n_iter = 1000; + let n_iter = config.n_iter; let mut total_elapsed = 0.0; for i in 0..n_iter { let mut commands = runner.commands(); @@ -88,17 +88,17 @@ impl PrefixTreeCode { let reduce_code = include_shader!(&runner.session, "../shader/gen/prefix_reduce"); let reduce_pipeline = runner .session - .create_simple_compute_pipeline(reduce_code, 2) + .create_compute_pipeline(reduce_code, &[BindType::BufReadOnly, BindType::Buffer]) .unwrap(); let scan_code = include_shader!(&runner.session, "../shader/gen/prefix_scan"); let scan_pipeline = runner .session - .create_simple_compute_pipeline(scan_code, 2) + .create_compute_pipeline(scan_code, &[BindType::Buffer, BindType::BufReadOnly]) .unwrap(); let root_code = include_shader!(&runner.session, "../shader/gen/prefix_root"); let root_pipeline = runner .session - .create_simple_compute_pipeline(root_code, 1) + .create_compute_pipeline(root_code, &[BindType::Buffer]) .unwrap(); PrefixTreeCode { reduce_pipeline, diff --git a/tests/src/runner.rs b/tests/src/runner.rs index 9bfde3b..ed57c29 100644 --- a/tests/src/runner.rs +++ b/tests/src/runner.rs @@ -16,7 +16,10 @@ //! Test runner intended to make it easy to write tests. -use piet_gpu_hal::{Buffer, BufferUsage, CmdBuf, Instance, PlainData, QueryPool, Session}; +use piet_gpu_hal::{ + BackendType, Buffer, BufferUsage, CmdBuf, Instance, InstanceFlags, PlainData, QueryPool, + Session, +}; pub struct Runner { #[allow(unused)] @@ -45,8 +48,8 @@ pub struct BufDown { } impl Runner { - pub unsafe fn new() -> Runner { - let (instance, _) = Instance::new(None).unwrap(); + pub unsafe fn new(flags: InstanceFlags) -> Runner { + let (instance, _) = Instance::new(None, flags).unwrap(); let device = instance.device(None).unwrap(); let session = Session::new(device); let cmd_buf_pool = Vec::new(); @@ -82,7 +85,7 @@ impl Runner { let submitted = self.session.run_cmd_buf(cmd_buf, &[], &[]).unwrap(); self.cmd_buf_pool.extend(submitted.wait().unwrap()); let timestamps = self.session.fetch_query_pool(&query_pool).unwrap(); - timestamps[0] + timestamps.get(0).copied().unwrap_or_default() } #[allow(unused)] @@ -114,6 +117,10 @@ impl Runner { .unwrap(); BufDown { stage_buf, dev_buf } } + + pub fn backend_type(&self) -> BackendType { + self.session.backend_type() + } } impl Commands { diff --git a/tests/src/test_result.rs b/tests/src/test_result.rs index 84bbc85..e582c63 100644 --- a/tests/src/test_result.rs +++ b/tests/src/test_result.rs @@ -21,10 +21,17 @@ pub struct TestResult { // TODO: statistics. We're lean and mean for now. total_time: f64, n_elements: u64, - failure: Option, + status: Status, } -#[derive(Clone, Copy)] +pub enum Status { + Pass, + Fail(String), + #[allow(unused)] + Skipped(String), +} + +#[derive(Clone, Copy, PartialEq, Eq)] pub enum ReportStyle { Short, Verbose, @@ -36,14 +43,15 @@ impl TestResult { name: name.to_string(), total_time: 0.0, n_elements: 0, - failure: None, + status: Status::Pass, } } pub fn report(&self, style: ReportStyle) { - let fail_string = match &self.failure { - None => "pass".into(), - Some(s) => format!("fail ({})", s), + let fail_string = match &self.status { + Status::Pass => "pass".into(), + Status::Fail(s) => format!("fail ({})", s), + Status::Skipped(s) => format!("skipped ({})", s), }; match style { ReportStyle::Short => { @@ -73,8 +81,13 @@ impl TestResult { } } - pub fn fail(&mut self, explanation: String) { - self.failure = Some(explanation); + pub fn fail(&mut self, explanation: impl Into) { + self.status = Status::Fail(explanation.into()); + } + + #[allow(unused)] + pub fn skip(&mut self, explanation: impl Into) { + self.status = Status::Skipped(explanation.into()); } pub fn timing(&mut self, total_time: f64, n_elements: u64) {