diff --git a/.gitattributes b/.gitattributes index 9d3b6f8..7b48404 100644 --- a/.gitattributes +++ b/.gitattributes @@ -1,3 +1,3 @@ **/shader/* linguist-language=glsl **/shader/gen/* linguist-generated - +piet-wgpu/shader/* linguist-language=wgsl diff --git a/Cargo.lock b/Cargo.lock index acc4ca4..bcfd743 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -8,6 +8,35 @@ version = "1.0.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "f26201604c87b1e01bd3d98f8d5d9a8fcbb815e8cedb41ffccbeb4bf593a35fe" +[[package]] +name = "ahash" +version = "0.7.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fcb51a0695d8f838b1ee009b3fbf66bda078cd64590202a864a8f3e8c4315c47" +dependencies = [ + "getrandom 0.2.7", + "once_cell", + "version_check", +] + +[[package]] +name = "aho-corasick" +version = "0.7.19" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b4f55bd91a0978cbfd91c457a164bab8b4001c833b7f323132c0a4e1922dd44e" +dependencies = [ + "memchr", +] + +[[package]] +name = "android_system_properties" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "819e7219dbd41043ac279b19830f2efc897156490d7fd6ea916720117ee66311" +dependencies = [ + "libc", +] + [[package]] name = "ansi_term" version = "0.12.1" @@ -66,6 +95,21 @@ version = "1.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d468802bab17cbc0cc575e9b053f41e72aa36bfa6b7f55e3529ffa43161b97fa" +[[package]] +name = "bit-set" +version = "0.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0700ddab506f33b20a03b13996eccd309a48e5ff77d0d95926aa0210fb4e95f1" +dependencies = [ + "bit-vec", +] + +[[package]] +name = "bit-vec" +version = "0.6.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "349f9b6a179ed607305526ca489b34ad0a41aed5f7980fa90eb03160b69598fb" + [[package]] name = "bitflags" version = "1.3.2" @@ -78,6 +122,15 @@ version = "0.1.6" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" +[[package]] +name = "block-buffer" +version = "0.10.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "69cce20737498f97b993470a6e536b8523f0af7892a4f928cceb1ac5e52ebe7e" +dependencies = [ + "generic-array", +] + [[package]] name = "bumpalo" version = "3.11.1" @@ -104,6 +157,12 @@ dependencies = [ "syn", ] +[[package]] +name = "byteorder" +version = "1.4.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "14c189c53d098945499cdfa7ecc63567cf3886b3332b312a5b4585d8d3a6a610" + [[package]] name = "calloop" version = "0.10.1" @@ -148,6 +207,12 @@ version = "1.0.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "baf1de4339761588bc0619e3cbc0120ee582ebb74b53b4efbf79117bd2da40fd" +[[package]] +name = "cfg_aliases" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fd16c4719339c4530435d38e511904438d07cce7950afa3718a84ac36c10e89e" + [[package]] name = "clap" version = "2.34.0" @@ -218,6 +283,16 @@ dependencies = [ "objc", ] +[[package]] +name = "codespan-reporting" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" +dependencies = [ + "termcolor", + "unicode-width", +] + [[package]] name = "core-foundation" version = "0.9.3" @@ -259,6 +334,15 @@ dependencies = [ "libc", ] +[[package]] +name = "cpufeatures" +version = "0.2.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "28d997bd5e24a5928dd43e46dc529867e207907fe0b239c3477d924f7f2ca320" +dependencies = [ + "libc", +] + [[package]] name = "crc32fast" version = "1.3.2" @@ -268,12 +352,33 @@ dependencies = [ "cfg-if", ] +[[package]] +name = "crypto-common" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1bfb12502f3fc46cca1bb51ac28df9d618d813cdc3d2f25b9fe775a34af26bb3" +dependencies = [ + "generic-array", + "typenum", +] + [[package]] name = "cty" version = "0.2.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b365fabc795046672053e29c954733ec3b05e4be654ab130fe8f1f94d7051f35" +[[package]] +name = "d3d12" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "827914e1f53b1e0e025ecd3d967a7836b7bcb54520f90e21ef8df7b4d88a2759" +dependencies = [ + "bitflags", + "libloading", + "winapi", +] + [[package]] name = "darling" version = "0.10.2" @@ -344,6 +449,16 @@ dependencies = [ "syn", ] +[[package]] +name = "digest" +version = "0.10.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "adfbc57365a37acbd2ebf2b64d7e69bb766e2fea813521ed536f5d0520dcf86c" +dependencies = [ + "block-buffer", + "crypto-common", +] + [[package]] name = "dispatch" version = "0.2.0" @@ -365,6 +480,19 @@ version = "1.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "9ea835d29036a4087793836fa931b08837ad5e957da9e23886b29586fb9b6650" +[[package]] +name = "env_logger" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c90bf5f19754d10198ccb95b70664fc925bd1fc090a0fd9a6ebc54acc8cd6272" +dependencies = [ + "atty", + "humantime", + "log", + "regex", + "termcolor", +] + [[package]] name = "fastrand" version = "1.8.0" @@ -405,6 +533,42 @@ version = "0.1.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "00b0228411908ca8685dba7fc2cdd70ec9990a6e753e89b6ac91a84c40fbaf4b" +[[package]] +name = "futures-core" +version = "0.3.25" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "04909a7a7e4633ae6c4a9ab280aeb86da1236243a77b694a49eacd659a4bd3ac" + +[[package]] +name = "futures-intrusive" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1b6bdbb8c5a42b2bb5ee8dd9dc2c7d73ce3e15d26dfe100fb347ffa3f58c672b" +dependencies = [ + "futures-core", + "lock_api", + "parking_lot", +] + +[[package]] +name = "fxhash" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c31b6d751ae2c7f11320402d34e41349dd1016f8d5d45e48c4312bc8625af50c" +dependencies = [ + "byteorder", +] + +[[package]] +name = "generic-array" +version = "0.14.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bff49e947297f3312447abdca79f45f4738097cc82b06e72054d2223f601f1b9" +dependencies = [ + "typenum", + "version_check", +] + [[package]] name = "getrandom" version = "0.1.16" @@ -427,17 +591,86 @@ dependencies = [ "wasi 0.11.0+wasi-snapshot-preview1", ] +[[package]] +name = "glow" +version = "0.11.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d8bd5877156a19b8ac83a29b2306fe20537429d318f3ff0a1a2119f8d9c61919" +dependencies = [ + "js-sys", + "slotmap", + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "gpu-alloc" +version = "0.5.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7fc59e5f710e310e76e6707f86c561dd646f69a8876da9131703b2f717de818d" +dependencies = [ + "bitflags", + "gpu-alloc-types", +] + +[[package]] +name = "gpu-alloc-types" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "54804d0d6bc9d7f26db4eaec1ad10def69b599315f487d32c334a80d1efe67a5" +dependencies = [ + "bitflags", +] + +[[package]] +name = "gpu-descriptor" +version = "0.2.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0b0c02e1ba0bdb14e965058ca34e09c020f8e507a760df1121728e0aef68d57a" +dependencies = [ + "bitflags", + "gpu-descriptor-types", + "hashbrown", +] + +[[package]] +name = "gpu-descriptor-types" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "363e3677e55ad168fef68cf9de3a4a310b53124c5e784c53a1d70e92d23f2126" +dependencies = [ + "bitflags", +] + [[package]] name = "half" version = "1.8.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "eabb4a44450da02c90444cf74558da904edde8fb4e9035a9a6a4e15445af0bd7" +[[package]] +name = "handlebars" +version = "4.3.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "433e4ab33f1213cdc25b5fa45c76881240cfe79284cf2b395e8b9e312a30a2fd" +dependencies = [ + "log", + "pest", + "pest_derive", + "serde", + "serde_json", + "thiserror", + "walkdir", +] + [[package]] name = "hashbrown" version = "0.12.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8a9ee70c43aaf417c914396645a0fa852624801b24ebb7ae78fe8272889ac888" +dependencies = [ + "ahash", +] [[package]] name = "heck" @@ -457,6 +690,18 @@ dependencies = [ "libc", ] +[[package]] +name = "hexf-parse" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" + +[[package]] +name = "humantime" +version = "2.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9a3a5bfb195931eeb336b2a7b4d761daec841b97f947d34394601737a7bba5e4" + [[package]] name = "ident_case" version = "1.0.1" @@ -506,6 +751,17 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "khronos-egl" +version = "4.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8c2352bd1d0bceb871cb9d40f24360c8133c11d7486b68b5381c1dd1a32015e3" +dependencies = [ + "libc", + "libloading", + "pkg-config", +] + [[package]] name = "kurbo" version = "0.7.1" @@ -647,6 +903,25 @@ dependencies = [ "pinot", ] +[[package]] +name = "naga" +version = "0.10.0" +source = "git+https://github.com/gfx-rs/naga?rev=ddcd5d3121150b2b1beee6e54e9125ff31aaa9a2#ddcd5d3121150b2b1beee6e54e9125ff31aaa9a2" +dependencies = [ + "bit-set", + "bitflags", + "codespan-reporting", + "hexf-parse", + "indexmap", + "log", + "num-traits", + "rustc-hash", + "spirv", + "termcolor", + "thiserror", + "unicode-xid", +] + [[package]] name = "ndk" version = "0.3.0" @@ -772,6 +1047,15 @@ dependencies = [ "minimal-lexical", ] +[[package]] +name = "num-traits" +version = "0.2.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "578ede34cf02f8924ab9447f50c28075b4d3e5b269972345e7e0372b38c6cdcd" +dependencies = [ + "autocfg", +] + [[package]] name = "num_enum" version = "0.5.7" @@ -853,6 +1137,50 @@ version = "2.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "478c572c3d73181ff3c2539045f6eb99e5491218eae919370993b890cdbdd98e" +[[package]] +name = "pest" +version = "2.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dbc7bc69c062e492337d74d59b120c274fd3d261b6bf6d3207d499b4b379c41a" +dependencies = [ + "thiserror", + "ucd-trie", +] + +[[package]] +name = "pest_derive" +version = "2.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "60b75706b9642ebcb34dab3bc7750f811609a0eb1dd8b88c2d15bf628c1c65b2" +dependencies = [ + "pest", + "pest_generator", +] + +[[package]] +name = "pest_generator" +version = "2.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f4f9272122f5979a6511a749af9db9bfc810393f63119970d7085fed1c4ea0db" +dependencies = [ + "pest", + "pest_meta", + "proc-macro2", + "quote", + "syn", +] + +[[package]] +name = "pest_meta" +version = "2.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4c8717927f9b79515e565a64fe46c38b8cd0427e64c40680b14a7365ab09ac8d" +dependencies = [ + "once_cell", + "pest", + "sha1", +] + [[package]] name = "pgpu-render" version = "0.1.0" @@ -946,6 +1274,23 @@ dependencies = [ "smallvec", ] +[[package]] +name = "piet-wgsl" +version = "0.1.0" +dependencies = [ + "bytemuck", + "env_logger", + "futures-intrusive", + "handlebars", + "parking_lot", + "piet-scene", + "png", + "pollster", + "serde", + "serde_json", + "wgpu", +] + [[package]] name = "pinot" version = "0.1.5" @@ -970,6 +1315,12 @@ dependencies = [ "miniz_oxide", ] +[[package]] +name = "pollster" +version = "0.2.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5da3b0203fd7ee5720aa0b5e790b591aa5d3f41c3ed2c34a3a393382198af2f7" + [[package]] name = "ppv-lite86" version = "0.2.16" @@ -1005,6 +1356,12 @@ dependencies = [ "unicode-ident", ] +[[package]] +name = "profiling" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "74605f360ce573babfe43964cbe520294dcb081afbf8c108fc6e23036b4da2df" + [[package]] name = "quote" version = "1.0.21" @@ -1085,6 +1442,12 @@ dependencies = [ "rand_core 0.5.1", ] +[[package]] +name = "range-alloc" +version = "0.1.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "63e935c45e09cc6dcf00d2f0b2d630a58f4095320223d47fc68918722f0538b6" + [[package]] name = "raw-window-handle" version = "0.3.4" @@ -1134,6 +1497,23 @@ dependencies = [ "bitflags", ] +[[package]] +name = "regex" +version = "1.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4c4eb3267174b8c6c2f654116623910a0fef09c4753f8dd83db29c48a0df988b" +dependencies = [ + "aho-corasick", + "memchr", + "regex-syntax", +] + +[[package]] +name = "regex-syntax" +version = "0.6.27" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a3f87b73ce11b1619a3c6332f45341e0047173771e8b8b73f87bfeefb7b56244" + [[package]] name = "remove_dir_all" version = "0.5.3" @@ -1143,6 +1523,12 @@ dependencies = [ "winapi", ] +[[package]] +name = "renderdoc-sys" +version = "0.7.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f1382d1f0a252c4bf97dc20d979a2fdd05b024acd7c2ed0f7595d7817666a157" + [[package]] name = "roxmltree" version = "0.13.1" @@ -1152,12 +1538,27 @@ dependencies = [ "xmlparser", ] +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + [[package]] name = "ryu" version = "1.0.11" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "4501abdff3ae82a1c1b477a17252eb69cee9e66eb915c1abaa4f44d873df9f09" +[[package]] +name = "same-file" +version = "1.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "93fc1dc3aaa9bfed95e02e6eadabb4baf7e3078b0bd1b4d7b6b0b68378900502" +dependencies = [ + "winapi-util", +] + [[package]] name = "scoped-tls" version = "1.0.0" @@ -1201,6 +1602,17 @@ dependencies = [ "serde", ] +[[package]] +name = "sha1" +version = "0.10.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f04293dc80c3993519f2d7f6f511707ee7094fe0c6d3406feb330cdb3540eba3" +dependencies = [ + "cfg-if", + "cpufeatures", + "digest", +] + [[package]] name = "slotmap" version = "1.0.6" @@ -1235,6 +1647,22 @@ dependencies = [ "wayland-protocols", ] +[[package]] +name = "spirv" +version = "0.2.0+1.5.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "246bfa38fe3db3f1dfc8ca5a2cdeb7348c78be2112740cc0ec8ef18b6d94f830" +dependencies = [ + "bitflags", + "num-traits", +] + +[[package]] +name = "static_assertions" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" + [[package]] name = "strsim" version = "0.8.0" @@ -1331,6 +1759,18 @@ dependencies = [ "serde", ] +[[package]] +name = "typenum" +version = "1.15.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dcf81ac59edc17cc8697ff311e8f5ef2d99fcbd9817b34cec66f90b6c3dfd987" + +[[package]] +name = "ucd-trie" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9e79c4d996edb816c91e4308506774452e55e95c3c9de07b6729e17e15a5ef81" + [[package]] name = "unicode-ident" version = "1.0.5" @@ -1349,6 +1789,12 @@ version = "0.1.10" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c0edd1e5b14653f783770bce4a4dabb4a5108a5370a5f5d8cfe8710c361f6c8b" +[[package]] +name = "unicode-xid" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f962df74c8c05a667b5ee8bcf162993134c104e96440b663c8daa176dc772d8c" + [[package]] name = "vec_map" version = "0.8.2" @@ -1361,6 +1807,17 @@ version = "0.9.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "49874b5167b65d7193b8aba1567f5c7d93d001cafc34600cee003eda787e483f" +[[package]] +name = "walkdir" +version = "2.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "808cf2735cd4b6866113f648b791c6adc5714537bc222d9347bb203386ffda56" +dependencies = [ + "same-file", + "winapi", + "winapi-util", +] + [[package]] name = "wasi" version = "0.9.0+wasi-snapshot-preview1" @@ -1398,6 +1855,18 @@ dependencies = [ "wasm-bindgen-shared", ] +[[package]] +name = "wasm-bindgen-futures" +version = "0.4.33" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "23639446165ca5a5de86ae1d8896b737ae80319560fbaa4c2887b7da6e7ebd7d" +dependencies = [ + "cfg-if", + "js-sys", + "wasm-bindgen", + "web-sys", +] + [[package]] name = "wasm-bindgen-macro" version = "0.2.83" @@ -1510,6 +1979,100 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "wgpu" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c2272b17bffc8a0c7d53897435da7c1db587c87d3a14e8dae9cdb8d1d210fc0f" +dependencies = [ + "arrayvec 0.7.2", + "js-sys", + "log", + "naga", + "parking_lot", + "raw-window-handle 0.5.0", + "smallvec", + "static_assertions", + "wasm-bindgen", + "wasm-bindgen-futures", + "web-sys", + "wgpu-core", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-core" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "73d14cad393054caf992ee02b7da6a372245d39a484f7461c1f44f6f6359bd28" +dependencies = [ + "arrayvec 0.7.2", + "bit-vec", + "bitflags", + "cfg_aliases", + "codespan-reporting", + "fxhash", + "log", + "naga", + "parking_lot", + "profiling", + "raw-window-handle 0.5.0", + "smallvec", + "thiserror", + "web-sys", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-hal" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cdae6a80dbc725343f02f854b310b37190be946aeea65e9d83afaa7d840ebaac" +dependencies = [ + "android_system_properties", + "arrayvec 0.7.2", + "ash", + "bit-set", + "bitflags", + "block", + "core-graphics-types", + "d3d12", + "foreign-types", + "fxhash", + "glow", + "gpu-alloc", + "gpu-descriptor", + "js-sys", + "khronos-egl", + "libloading", + "log", + "metal", + "naga", + "objc", + "parking_lot", + "profiling", + "range-alloc", + "raw-window-handle 0.5.0", + "renderdoc-sys", + "smallvec", + "thiserror", + "wasm-bindgen", + "web-sys", + "wgpu-types", + "winapi", +] + +[[package]] +name = "wgpu-types" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "28fb86c1909233c804aa79b7dd1ad06ebd979b2a465e3e980582db0ea9e69f3f" +dependencies = [ + "bitflags", +] + [[package]] name = "winapi" version = "0.3.9" diff --git a/Cargo.toml b/Cargo.toml index ab8c423..dc40744 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -8,5 +8,11 @@ members = [ "piet-gpu-hal", "piet-gpu-types", "piet-scene", + "piet-wgsl", "tests", ] + +[patch.crates-io] +# Required for metal support to work on wgpu +# TODO: remove when wgpu is upgraded to 0.15 +naga = { git = "https://github.com/gfx-rs/naga", rev="ddcd5d3121150b2b1beee6e54e9125ff31aaa9a2" } diff --git a/piet-wgsl/Cargo.toml b/piet-wgsl/Cargo.toml new file mode 100644 index 0000000..020ba92 --- /dev/null +++ b/piet-wgsl/Cargo.toml @@ -0,0 +1,20 @@ +[package] +name = "piet-wgsl" +version = "0.1.0" +authors = ["Raph Levien "] +license = "MIT/Apache-2.0" +edition = "2021" + +[dependencies] +wgpu = "0.14" +env_logger = "0.9.1" +pollster = "0.2.5" +futures-intrusive = "0.4.1" +parking_lot = "0.12" +bytemuck = { version = "1.12.1", features = ["derive"] } +handlebars = { version = "4.3.5", features = ["dir_source"] } +serde_json = "1" +serde = { version = "1", features = ["derive"] } +png = "0.17.6" + +piet-scene = { path = "../piet-scene" } diff --git a/piet-wgsl/README.md b/piet-wgsl/README.md new file mode 100644 index 0000000..cc7f0e0 --- /dev/null +++ b/piet-wgsl/README.md @@ -0,0 +1,11 @@ +# piet-wgsl + +This crate is currently a highly experimental proof-of-concept port of the piet-gpu renderer to the WGSL shader language, so it could be run on WebGPU. Depending on how well it works out, it may become the authoritative source for piet-gpu. + +The shaders are actually handlebars templates over WGSL, as it's important to share common data structures; it's likely we'll use the template mechanism to supply various parameters which are not supported by the WGSL language, for example to specify grayscale or RGBA buffer output for fine rasterization. + +This crate also uses a very different approach to the GPU abstraction than piet-gpu. That is essentially a HAL that supports an immediate mode approach to creating resources and submitting commands. Here, we generate a `Recording`, which is basically a simple value type, then an `Engine` plays that recording to the actual GPU. The idea is that this can abstract easily over multiple GPU back-ends, without either the render logic needing to be polymorphic or having dynamic dispatch at the GPU abstraction. The goal is to be more agile. + +Scene encoding is shared with piet-gpu, and currently uses piet-scene in the same repo with no changes. + +This module is still an experimental work in progress. Contributions can be made with the same policy as the root repo, but expect things to change quickly. diff --git a/piet-wgsl/shader/backdrop.twgsl b/piet-wgsl/shader/backdrop.twgsl new file mode 100644 index 0000000..1618ae4 --- /dev/null +++ b/piet-wgsl/shader/backdrop.twgsl @@ -0,0 +1,61 @@ +// Copyright 2022 Google LLC +// +// 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. + +// Note: this is the non-atomic version +struct Tile { + backdrop: i32, + segments: u32, +} + +{{> config}} + +@group(0) @binding(0) +var config: Config; + +@group(0) @binding(1) +var tiles: array; + +let WG_SIZE = 64u; + +var sh_backdrop: array; + +// Each workgroup computes the inclusive prefix sum of the backdrops +// in one row of tiles. +@compute @workgroup_size(64) +fn main( + @builtin(local_invocation_id) local_id: vec3, + @builtin(workgroup_id) wg_id: vec3, +) { + let width_in_tiles = config.width_in_tiles; + let ix = wg_id.x * width_in_tiles + local_id.x; + var backdrop = 0; + if local_id.x < width_in_tiles { + backdrop = tiles[ix].backdrop; + } + sh_backdrop[local_id.x] = backdrop; + // iterate log2(WG_SIZE) times + for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) { + workgroupBarrier(); + if local_id.x >= (1u << i) { + backdrop += sh_backdrop[local_id.x - (1u << i)]; + } + workgroupBarrier(); + sh_backdrop[local_id.x] = backdrop; + } + if local_id.x < width_in_tiles { + tiles[ix].backdrop = backdrop; + } +} diff --git a/piet-wgsl/shader/config.twgsl b/piet-wgsl/shader/config.twgsl new file mode 100644 index 0000000..704a608 --- /dev/null +++ b/piet-wgsl/shader/config.twgsl @@ -0,0 +1,20 @@ +// Copyright 2022 Google LLC +// +// 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. + +struct Config { + width_in_tiles: u32, + height_in_tiles: u32, +} diff --git a/piet-wgsl/shader/fine.twgsl b/piet-wgsl/shader/fine.twgsl new file mode 100644 index 0000000..203c320 --- /dev/null +++ b/piet-wgsl/shader/fine.twgsl @@ -0,0 +1,96 @@ +// Copyright 2022 Google LLC +// +// 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. + +// This is a cut'n'paste w/ backdrop. +struct Tile { + backdrop: i32, + segments: u32, +} + +{{> segment}} +{{> config}} + +@group(0) @binding(0) +var config: Config; + +@group(0) @binding(1) +var tiles: array; + +@group(0) @binding(2) +var segments: array; + +// This will become a texture, but keeping things simple for now +@group(0) @binding(3) +var output: array; + +let PIXELS_PER_THREAD = 4u; + +@compute @workgroup_size(4, 16) +fn main( + @builtin(global_invocation_id) global_id: vec3, + @builtin(local_invocation_id) local_id: vec3, + @builtin(workgroup_id) wg_id: vec3, +) { + let tile_ix = wg_id.y * config.width_in_tiles + wg_id.x; + let xy = vec2(f32(global_id.x * PIXELS_PER_THREAD), f32(global_id.y)); + let tile = tiles[tile_ix]; + var area: array; + let backdrop_f = f32(tile.backdrop); + for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { + area[i] = backdrop_f; + } + var segment_ix = tile.segments; + while segment_ix != 0u { + let segment = segments[segment_ix]; + let y = segment.origin.y - xy.y; + let y0 = clamp(y, 0.0, 1.0); + let y1 = clamp(y + segment.delta.y, 0.0, 1.0); + let dy = y0 - y1; + if dy != 0.0 { + let vec_y_recip = 1.0 / segment.delta.y; + let t0 = (y0 - y) * vec_y_recip; + let t1 = (y1 - y) * vec_y_recip; + let startx = segment.origin.x - xy.x; + let x0 = startx + t0 * segment.delta.x; + let x1 = startx + t1 * segment.delta.x; + let xmin0 = min(x0, x1); + let xmax0 = max(x0, x1); + for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { + let i_f = f32(i); + let xmin = min(xmin0 - i_f, 1.0) - 1.0e-6; + let xmax = xmax0 - i_f; + let b = min(xmax, 1.0); + let c = max(b, 0.0); + let d = max(xmin, 0.0); + let a = (b + 0.5 * (d * d - c * c) - xmin) / (xmax - xmin); + area[i] += a * dy; + } + } + let y_edge = sign(segment.delta.x) * clamp(xy.y - segment.y_edge + 1.0, 0.0, 1.0); + for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { + area[i] += y_edge; + } + segment_ix = segment.next; + } + // nonzero winding rule + for (var i = 0u; i < PIXELS_PER_THREAD; i += 1u) { + area[i] = abs(area[i]); + } + + let bytes = pack4x8unorm(vec4(area[0], area[1], area[2], area[3])); + let out_ix = global_id.y * (config.width_in_tiles * 4u) + global_id.x; + output[out_ix] = bytes; +} diff --git a/piet-wgsl/shader/path_coarse.twgsl b/piet-wgsl/shader/path_coarse.twgsl new file mode 100644 index 0000000..a34d5b1 --- /dev/null +++ b/piet-wgsl/shader/path_coarse.twgsl @@ -0,0 +1,328 @@ +// Copyright 2022 Google LLC +// +// 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. + +{{> pathtag}} + +@group(0) @binding(0) +var path_tags: array; + +@group(0) @binding(1) +var tag_monoids: array; + +// TODO: should probably have single "scene" binding. +@group(0) @binding(2) +var path_data: array; + +{{#if cubics_out}} +@group(0) @binding(3) +var output: array>; +{{else}} +{{> config}} + +struct Tile { + backdrop: atomic, + segments: atomic, +} + +{{> segment}} + +// Should probably be uniform binding +@group(0) @binding(3) +var config: Config; + +@group(0) @binding(4) +var tiles: array; + +@group(0) @binding(5) +var segments: array; +{{/if}} + +fn read_f32_point(ix: u32) -> vec2 { + let x = bitcast(path_data[ix]); + let y = bitcast(path_data[ix + 1u]); + return vec2(x, y); +} + +fn read_i16_point(ix: u32) -> vec2 { + let raw = path_data[ix]; + let x = f32(i32(raw << 16u) >> 16u); + let y = f32(i32(raw) >> 16u); + return vec2(x, y); +} + +{{#unless cubics_out}} +let TILE_WIDTH = 16u; +let TILE_HEIGHT = 16u; + +struct SubdivResult { + val: f32, + a0: f32, + a2: f32, +} + +let D = 0.67; +fn approx_parabola_integral(x: f32) -> f32 { + return x * inverseSqrt(sqrt(1.0 - D + (D * D * D * D + 0.25 * x * x))); +} + +let B = 0.39; +fn approx_parabola_inv_integral(x: f32) -> f32 { + return x * sqrt(1.0 - B + (B * B + 0.5 * x * x)); +} + +fn estimate_subdiv(p0: vec2, p1: vec2, p2: vec2, sqrt_tol: f32) -> SubdivResult { + let d01 = p1 - p0; + let d12 = p2 - p1; + let dd = d01 - d12; + let cross = (p2.x - p0.x) * dd.y - (p2.y - p0.y) * dd.x; + let cross_inv = 1.0 / cross; + let x0 = dot(d01, dd) * cross_inv; + let x2 = dot(d12, dd) * cross_inv; + let scale = abs(cross / (length(dd) * (x2 - x0))); + + let a0 = approx_parabola_integral(x0); + let a2 = approx_parabola_integral(x2); + var val = 0.0; + if scale < 1e9 { + let da = abs(a2 - a0); + let sqrt_scale = sqrt(scale); + if sign(x0) == sign(x2) { + val = sqrt_scale; + } else { + let xmin = sqrt_tol / sqrt_scale; + val = sqrt_tol / approx_parabola_integral(xmin); + } + val *= da; + } + return SubdivResult(val, a0, a2); +} + +fn eval_quad(p0: vec2, p1: vec2, p2: vec2, t: f32) -> vec2 { + let mt = 1.0 - t; + return p0 * (mt * mt) + (p1 * (mt * 2.0) + p2 * t) * t; +} + +fn eval_cubic(p0: vec2, p1: vec2, p2: vec2, p3: vec2, t: f32) -> vec2 { + let mt = 1.0 - t; + return p0 * (mt * mt * mt) + (p1 * (mt * mt * 3.0) + (p2 * (mt * 3.0) + p3 * t) * t) * t; +} + +fn alloc_segment() -> u32 { + // Use 0-index segment (address is sentinel) as counter + // TODO: separate small buffer binding for this? + return atomicAdd(&tiles[4096].segments, 1u) + 1u; +} +{{/unless}} + +let MAX_QUADS = 16u; + +@compute @workgroup_size(256) +fn main( + @builtin(global_invocation_id) global_id: vec3, + @builtin(local_invocation_id) local_id: vec3, +) { + // Obtain exclusive prefix sum of tag monoid + let ix = global_id.x; + let tag_word = path_tags[ix >> 2u]; + let shift = (ix & 3u) * 8u; + var tm = reduce_tag(tag_word & ((1u << shift) - 1u)); + tm = combine_tag_monoid(tag_monoids[ix >> 2u], tm); + var tag_byte = (tag_word >> shift) & 0xffu; + // should be extractBits(tag_word, shift, 8)? + + // Decode path data + let seg_type = tag_byte & PATH_TAG_SEG_TYPE; + if seg_type != 0u { + var p0: vec2; + var p1: vec2; + var p2: vec2; + var p3: vec2; + if (tag_byte & PATH_TAG_F32) != 0u { + p0 = read_f32_point(tm.pathseg_offset); + p1 = read_f32_point(tm.pathseg_offset + 2u); + if seg_type >= PATH_TAG_QUADTO { + p2 = read_f32_point(tm.pathseg_offset + 4u); + if seg_type == PATH_TAG_CUBICTO { + p3 = read_f32_point(tm.pathseg_offset + 6u); + } + } + } else { + p0 = read_i16_point(tm.pathseg_offset); + p1 = read_i16_point(tm.pathseg_offset + 1u); + if seg_type >= PATH_TAG_QUADTO { + p2 = read_i16_point(tm.pathseg_offset + 2u); + if seg_type == PATH_TAG_CUBICTO { + p3 = read_i16_point(tm.pathseg_offset + 3u); + } + } + } + // TODO: transform goes here + // Degree-raise + if seg_type == PATH_TAG_LINETO { + p3 = p1; + p2 = mix(p3, p0, 1.0 / 3.0); + p1 = mix(p0, p3, 1.0 / 3.0); + } else if seg_type == PATH_TAG_QUADTO { + p3 = p2; + p2 = mix(p1, p2, 1.0 / 3.0); + p1 = mix(p1, p0, 1.0 / 3.0); + } +{{#if cubics_out}} + let out_ix = ix * 4u; + output[out_ix] = p0; + output[out_ix + 1u] = p1; + output[out_ix + 2u] = p2; + output[out_ix + 3u] = p3; +{{else}} + let err_v = 3.0 * (p2 - p1) + p0 - p3; + let err = dot(err_v, err_v); + let ACCURACY = 0.25; + let Q_ACCURACY = ACCURACY * 0.1; + let REM_ACCURACY = (ACCURACY - Q_ACCURACY); + let MAX_HYPOT2 = 432.0 * Q_ACCURACY * Q_ACCURACY; + var n_quads = max(u32(ceil(pow(err * (1.0 / MAX_HYPOT2), 1.0 / 6.0))), 1u); + n_quads = min(n_quads, MAX_QUADS); + var keep_params: array; + var val = 0.0; + var qp0 = p0; + let step = 1.0 / f32(n_quads); + for (var i = 0u; i < n_quads; i += 1u) { + let t = f32(i + 1u) * step; + let qp2 = eval_cubic(p0, p1, p2, p3, t); + var qp1 = eval_cubic(p0, p1, p2, p3, t - 0.5 * step); + qp1 = 2.0 * qp1 - 0.5 * (qp0 + qp2); + let params = estimate_subdiv(qp0, qp1, qp2, sqrt(REM_ACCURACY)); + keep_params[i] = params; + val += params.val; + qp0 = qp2; + } + let n = max(u32(ceil(val * (0.5 / sqrt(REM_ACCURACY)))), 1u); + var lp0 = p0; + qp0 = p0; + let v_step = val / f32(n); + var n_out = 1u; + var val_sum = 0.0; + for (var i = 0u; i < n_quads; i += 1u) { + let t = f32(i + 1u) * step; + let qp2 = eval_cubic(p0, p1, p2, p3, t); + var qp1 = eval_cubic(p0, p1, p2, p3, t - 0.5 * step); + qp1 = 2.0 * qp1 - 0.5 * (qp0 + qp2); + let params = keep_params[i]; + let u0 = approx_parabola_inv_integral(params.a0); + let u2 = approx_parabola_inv_integral(params.a2); + let uscale = 1.0 / (u2 - u0); + var val_target = f32(n_out) * v_step; + while n_out == n || val_target < val_sum + params.val { + var lp1: vec2; + if n_out == n { + lp1 = p3; + } else { + let u = (val_target - val_sum) / params.val; + let a = mix(params.a0, params.a2, u); + let au = approx_parabola_inv_integral(a); + let t = (au - u0) * uscale; + lp1 = eval_quad(qp0, qp1, qp2, t); + } + + // Output line segment lp0..lp1 + let xymin = min(lp0, lp1); + let xymax = max(lp0, lp1); + let dp = lp1 - lp0; + let recip_dx = 1.0 / dp.x; + let invslope = select(dp.x / dp.y, 1.0e9, abs(dp.y) < 1.0e-9); + let c = 0.5 * abs(invslope); + let b = invslope; + let SX = 1.0 / f32(TILE_WIDTH); + let SY = 1.0 / f32(TILE_HEIGHT); + let a = (lp0.x - (lp0.y - 0.5 * f32(TILE_HEIGHT)) * b) * SX; + var x0 = i32(floor(xymin.x * SX)); + var x1 = i32(floor(xymax.x * SX) + 1.0); + var y0 = i32(floor(xymin.y * SY)); + var y1 = i32(floor(xymax.y * SY) + 1.0); + x0 = clamp(x0, 0, i32(config.width_in_tiles)); + x1 = clamp(x1, 0, i32(config.width_in_tiles)); + y0 = clamp(y0, 0, i32(config.height_in_tiles)); + y1 = clamp(y1, 0, i32(config.height_in_tiles)); + var xc = a + b * f32(y0); + var xray = i32(floor(lp0.x * SX)); + var last_xray = i32(floor(lp1.x * SX)); + if dp.y < 0.0 { + let tmp = xray; + xray = last_xray; + last_xray = tmp; + } + for (var y = y0; y < y1; y += 1) { + let tile_y0 = f32(y) * f32(TILE_HEIGHT); + let xbackdrop = max(xray + 1, 0); + if xymin.y < tile_y0 && xbackdrop < i32(config.width_in_tiles) { + let backdrop = select(-1, 1, dp.y < 0.0); + let tile_ix = y * i32(config.width_in_tiles) + xbackdrop; + atomicAdd(&tiles[tile_ix].backdrop, backdrop); + } + var next_xray = last_xray; + if y + 1 < y1 { + let tile_y1 = f32(y + 1) * f32(TILE_HEIGHT); + let x_edge = lp0.x + (tile_y1 - lp0.y) * invslope; + next_xray = i32(floor(x_edge * SX)); + } + let min_xray = min(xray, next_xray); + let max_xray = max(xray, next_xray); + var xx0 = min(i32(floor(xc - c)), min_xray); + var xx1 = max(i32(ceil(xc + c)), max_xray + 1); + xx0 = clamp(xx0, x0, x1); + xx1 = clamp(xx1, x0, x1); + var tile_seg: Segment; + for (var x = xx0; x < xx1; x += 1) { + let tile_x0 = f32(x) * f32(TILE_WIDTH); + let tile_ix = y * i32(config.width_in_tiles) + x; + // allocate segment, insert linked list + let seg_ix = alloc_segment(); + let old = atomicExchange(&tiles[tile_ix].segments, seg_ix); + tile_seg.origin = lp0; + tile_seg.delta = dp; + var y_edge = mix(lp0.y, lp1.y, (tile_x0 - lp0.x) * recip_dx); + if xymin.x < tile_x0 { + let p = vec2(tile_x0, y_edge); + if dp.x < 0.0 { + tile_seg.delta = p - lp0; + } else { + tile_seg.origin = p; + tile_seg.delta = lp1 - p; + } + if tile_seg.delta.x == 0.0 { + tile_seg.delta.x = sign(dp.x) * 1e-9; + } + } + if x <= min_xray || max_xray < x { + y_edge = 1e9; + } + tile_seg.y_edge = y_edge; + tile_seg.next = old; + segments[seg_ix] = tile_seg; + } + xc += b; + xray = next_xray; + } + n_out += 1u; + val_target += v_step; + lp0 = lp1; + } + val_sum += params.val; + qp0 = qp2; + } +{{/if}} + } +} diff --git a/piet-wgsl/shader/pathtag.twgsl b/piet-wgsl/shader/pathtag.twgsl new file mode 100644 index 0000000..e4cfda3 --- /dev/null +++ b/piet-wgsl/shader/pathtag.twgsl @@ -0,0 +1,59 @@ +// Copyright 2022 Google LLC +// +// 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. + +struct TagMonoid { + trans_ix: u32, + pathseg_ix: u32, + pathseg_offset: u32, + // Note: piet-gpu has linewidth and path, but not needed here +} + +let PATH_TAG_SEG_TYPE = 3u; +let PATH_TAG_LINETO = 1u; +let PATH_TAG_QUADTO = 2u; +let PATH_TAG_CUBICTO = 3u; +let PATH_TAG_F32 = 8u; +let PATH_TAG_PATH = 0x10u; +let PATH_TAG_TRANSFORM = 0x20u; + +fn tag_monoid_identity() -> TagMonoid { + var c: TagMonoid; + c.trans_ix = 0u; + c.pathseg_ix = 0u; + c.pathseg_offset = 0u; + return c; +} + +fn combine_tag_monoid(a: TagMonoid, b: TagMonoid) -> TagMonoid { + var c: TagMonoid; + c.trans_ix = a.trans_ix + b.trans_ix; + c.pathseg_ix = a.pathseg_ix + b.pathseg_ix; + c.pathseg_offset = a.pathseg_offset + b.pathseg_offset; + return c; +} + +fn reduce_tag(tag_word: u32) -> TagMonoid { + var c: TagMonoid; + let point_count = tag_word & 0x3030303u; + c.pathseg_ix = countOneBits((point_count * 7u) & 0x4040404u); + c.trans_ix = countOneBits(tag_word & (PATH_TAG_TRANSFORM * 0x1010101u)); + let n_points = point_count + ((tag_word >> 2u) & 0x1010101u); + var a = n_points + (n_points & (((tag_word >> 3u) & 0x1010101u) * 15u)); + a += a >> 8u; + a += a >> 16u; + c.pathseg_offset = a & 0xffu; + return c; +} diff --git a/piet-wgsl/shader/pathtag_reduce.twgsl b/piet-wgsl/shader/pathtag_reduce.twgsl new file mode 100644 index 0000000..188f780 --- /dev/null +++ b/piet-wgsl/shader/pathtag_reduce.twgsl @@ -0,0 +1,53 @@ +// Copyright 2022 Google LLC +// +// 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. + +{{> pathtag}} + +// Note: should have a single scene binding, path_tags are a slice +// in that; need a config uniform. +@group(0) @binding(0) +var path_tags: array; + +@group(0) @binding(1) +var reduced: array; + +let LG_WG_SIZE = 8u; +let WG_SIZE = 256u; + +var sh_scratch: array; + +@compute @workgroup_size(256) +fn main( + @builtin(global_invocation_id) global_id: vec3, + @builtin(local_invocation_id) local_id: vec3, +) { + let ix = global_id.x; + let tag_word = path_tags[ix]; + var agg = reduce_tag(tag_word); + sh_scratch[local_id.x] = agg; + for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) { + workgroupBarrier(); + if local_id.x + (1u << i) < WG_SIZE { + let other = sh_scratch[local_id.x + (1u << i)]; + agg = combine_tag_monoid(agg, other); + } + workgroupBarrier(); + sh_scratch[local_id.x] = agg; + } + if local_id.x == 0u { + reduced[ix >> LG_WG_SIZE] = agg; + } +} \ No newline at end of file diff --git a/piet-wgsl/shader/pathtag_scan.twgsl b/piet-wgsl/shader/pathtag_scan.twgsl new file mode 100644 index 0000000..5d5db59 --- /dev/null +++ b/piet-wgsl/shader/pathtag_scan.twgsl @@ -0,0 +1,76 @@ +// Copyright 2022 Google LLC +// +// 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. + +{{> pathtag}} + +@group(0) @binding(0) +var path_tags: array; + +@group(0) @binding(1) +var reduced: array; + +@group(0) @binding(2) +var tag_monoids: array; + +let LG_WG_SIZE = 8u; +let WG_SIZE = 256u; + +var sh_parent: array; +// These could be combined? +var sh_monoid: array; + +@compute @workgroup_size(256) +fn main( + @builtin(global_invocation_id) global_id: vec3, + @builtin(local_invocation_id) local_id: vec3, + @builtin(workgroup_id) wg_id: vec3, +) { + var agg = tag_monoid_identity(); + if (local_id.x < wg_id.x) { + agg = reduced[local_id.x]; + } + sh_parent[local_id.x] = agg; + for (var i = 0u; i < LG_WG_SIZE; i += 1u) { + workgroupBarrier(); + if (local_id.x + (1u << i) < WG_SIZE) { + let other = sh_parent[local_id.x + (1u << i)]; + agg = combine_tag_monoid(agg, other); + } + workgroupBarrier(); + sh_parent[local_id.x] = agg; + } + + let ix = global_id.x; + let tag_word = path_tags[ix]; + agg = reduce_tag(tag_word); + sh_monoid[local_id.x] = agg; + for (var i = 0u; i < LG_WG_SIZE; i += 1u) { + workgroupBarrier(); + if (local_id.x >= 1u << i) { + let other = sh_monoid[local_id.x - (1u << i)]; + agg = combine_tag_monoid(other, agg); + } + workgroupBarrier(); + sh_monoid[local_id.x] = agg; + } + // prefix up to this workgroup + var tm = sh_parent[0]; + if (local_id.x > 0u) { + tm = combine_tag_monoid(tm, sh_monoid[local_id.x - 1u]); + } + // exclusive prefix sum, granularity of 4 tag bytes + tag_monoids[ix] = tm; +} diff --git a/piet-wgsl/shader/segment.twgsl b/piet-wgsl/shader/segment.twgsl new file mode 100644 index 0000000..158fb2e --- /dev/null +++ b/piet-wgsl/shader/segment.twgsl @@ -0,0 +1,22 @@ +// Copyright 2022 Google LLC +// +// 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. + +struct Segment { + origin: vec2, + delta: vec2, + y_edge: f32, + next: u32, +} diff --git a/piet-wgsl/src/engine.rs b/piet-wgsl/src/engine.rs new file mode 100644 index 0000000..2be08db --- /dev/null +++ b/piet-wgsl/src/engine.rs @@ -0,0 +1,356 @@ +// Copyright 2022 Google LLC +// +// 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. + +use std::{ + borrow::Cow, + collections::{hash_map::Entry, HashMap}, + num::NonZeroU64, + sync::atomic::{AtomicU64, Ordering}, +}; + +use futures_intrusive::channel::shared::GenericOneshotReceiver; +use parking_lot::RawMutex; +use wgpu::{ + util::DeviceExt, BindGroup, BindGroupLayout, Buffer, BufferAsyncError, BufferSlice, BufferView, + ComputePipeline, Device, Queue, +}; + +pub type Error = Box; + +#[derive(Clone, Copy)] +pub struct ShaderId(usize); + +#[derive(Clone, Copy, PartialEq, Eq, Hash)] +pub struct Id(NonZeroU64); + +static ID_COUNTER: AtomicU64 = AtomicU64::new(0); + +pub struct Engine { + shaders: Vec, +} + +struct Shader { + pipeline: ComputePipeline, + bind_group_layout: BindGroupLayout, +} + +#[derive(Default)] +pub struct Recording { + commands: Vec, +} + +#[derive(Clone, Copy)] +pub struct BufProxy { + size: u64, + id: Id, +} + +pub enum Command { + Upload(BufProxy, Vec), + // Discussion question: third argument is vec of resources? + // Maybe use tricks to make more ergonomic? + // Alternative: provide bufs & images as separate sequences, like piet-gpu. + Dispatch(ShaderId, (u32, u32, u32), Vec), + Download(BufProxy), + Clear(BufProxy, u64, Option), +} + +#[derive(Default)] +pub struct Downloads { + buf_map: HashMap, +} + +/// 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. + #[allow(unused)] // TODO + Image, + /// A storage image with read only access. + #[allow(unused)] // TODO + ImageRead, + // TODO: Uniform, Sampler, maybe others +} + +#[derive(Default)] +struct BindMap { + buf_map: HashMap, +} + +impl Engine { + pub fn new() -> Engine { + Engine { shaders: vec![] } + } + + /// Add a shader. + /// + /// This function is somewhat limited, it doesn't apply a label, only allows one bind group, + /// doesn't support push constants, and entry point is hardcoded as "main". + /// + /// Maybe should do template instantiation here? But shader compilation pipeline feels maybe + /// a bit separate. + pub fn add_shader( + &mut self, + device: &Device, + wgsl: Cow<'static, str>, + layout: &[BindType], + ) -> Result { + let shader_module = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: None, + source: wgpu::ShaderSource::Wgsl(wgsl), + }); + let entries = layout + .iter() + .enumerate() + .map(|(i, bind_type)| match bind_type { + BindType::Buffer | BindType::BufReadOnly => wgpu::BindGroupLayoutEntry { + binding: i as u32, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Storage { + read_only: *bind_type == BindType::BufReadOnly, + }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + _ => todo!(), + }) + .collect::>(); + let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: None, + entries: &entries, + }); + let compute_pipeline_layout = + device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: None, + bind_group_layouts: &[&bind_group_layout], + push_constant_ranges: &[], + }); + let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: None, + layout: Some(&compute_pipeline_layout), + module: &shader_module, + entry_point: "main", + }); + let shader = Shader { + pipeline, + bind_group_layout, + }; + let id = self.shaders.len(); + self.shaders.push(shader); + Ok(ShaderId(id)) + } + + pub fn run_recording( + &mut self, + device: &Device, + queue: &Queue, + recording: &Recording, + ) -> Result { + let mut bind_map = BindMap::default(); + let mut downloads = Downloads::default(); + + let mut encoder = device.create_command_encoder(&Default::default()); + for command in &recording.commands { + match command { + Command::Upload(buf_proxy, bytes) => { + let buf = device.create_buffer_init(&wgpu::util::BufferInitDescriptor { + label: None, + contents: &bytes, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + }); + bind_map.insert_buf(buf_proxy.id, buf); + } + Command::Dispatch(shader_id, wg_size, bindings) => { + let shader = &self.shaders[shader_id.0]; + let bind_group = + bind_map.create_bind_group(device, &shader.bind_group_layout, bindings)?; + let mut cpass = encoder.begin_compute_pass(&Default::default()); + cpass.set_pipeline(&shader.pipeline); + cpass.set_bind_group(0, &bind_group, &[]); + cpass.dispatch_workgroups(wg_size.0, wg_size.1, wg_size.2); + } + Command::Download(proxy) => { + let src_buf = bind_map.buf_map.get(&proxy.id).ok_or("buffer not in map")?; + let buf = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: proxy.size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + encoder.copy_buffer_to_buffer(src_buf, 0, &buf, 0, proxy.size); + downloads.buf_map.insert(proxy.id, buf); + } + Command::Clear(proxy, offset, size) => { + let buffer = bind_map.get_or_create(*proxy, device)?; + encoder.clear_buffer(buffer, *offset, *size) + } + } + } + queue.submit(Some(encoder.finish())); + Ok(downloads) + } +} + +impl Recording { + pub fn push(&mut self, cmd: Command) { + self.commands.push(cmd); + } + + pub fn upload(&mut self, data: impl Into>) -> BufProxy { + let data = data.into(); + let buf_proxy = BufProxy::new(data.len() as u64); + self.push(Command::Upload(buf_proxy, data)); + buf_proxy + } + + pub fn dispatch( + &mut self, + shader: ShaderId, + wg_size: (u32, u32, u32), + resources: impl Into>, + ) { + self.push(Command::Dispatch(shader, wg_size, resources.into())); + } + + pub fn download(&mut self, buf: BufProxy) { + self.push(Command::Download(buf)); + } + + pub fn clear_all(&mut self, buf: BufProxy) { + self.push(Command::Clear(buf, 0, None)); + } +} + +impl BufProxy { + pub fn new(size: u64) -> Self { + let id = Id::next(); + BufProxy { id, size } + } +} + +impl Id { + pub fn next() -> Id { + let val = ID_COUNTER.fetch_add(1, Ordering::Relaxed); + // could use new_unchecked + Id(NonZeroU64::new(val + 1).unwrap()) + } +} + +impl BindMap { + fn insert_buf(&mut self, id: Id, buf: Buffer) { + self.buf_map.insert(id, buf); + } + + fn create_bind_group( + &mut self, + device: &Device, + layout: &BindGroupLayout, + bindings: &[BufProxy], + ) -> Result { + for proxy in bindings { + if let Entry::Vacant(v) = self.buf_map.entry(proxy.id) { + let buf = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: proxy.size, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + v.insert(buf); + } + } + let entries = bindings + .iter() + .enumerate() + .map(|(i, proxy)| { + let buf = self.buf_map.get(&proxy.id).unwrap(); + Ok(wgpu::BindGroupEntry { + binding: i as u32, + resource: buf.as_entire_binding(), + }) + }) + .collect::, Error>>()?; + let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout, + entries: &entries, + }); + Ok(bind_group) + } + + fn get_or_create(&mut self, proxy: BufProxy, device: &Device) -> Result<&Buffer, Error> { + match self.buf_map.entry(proxy.id) { + Entry::Occupied(occupied) => Ok(occupied.into_mut()), + Entry::Vacant(vacant) => { + let buf = device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: proxy.size, + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + Ok(vacant.insert(buf)) + } + } + } +} + +pub struct DownloadsMapped<'a>( + HashMap< + Id, + ( + BufferSlice<'a>, + GenericOneshotReceiver>, + ), + >, +); + +impl Downloads { + // Discussion: should API change so we get one buffer, rather than mapping all? + pub fn map(&self) -> DownloadsMapped { + let mut map = HashMap::new(); + for (id, buf) in &self.buf_map { + let buf_slice = buf.slice(..); + let (sender, receiver) = futures_intrusive::channel::shared::oneshot_channel(); + buf_slice.map_async(wgpu::MapMode::Read, move |v| sender.send(v).unwrap()); + map.insert(*id, (buf_slice, receiver)); + } + DownloadsMapped(map) + } +} + +impl<'a> DownloadsMapped<'a> { + pub async fn get_mapped(&self, proxy: BufProxy) -> Result { + let (slice, recv) = self.0.get(&proxy.id).ok_or("buffer not in map")?; + if let Some(recv_result) = recv.receive().await { + recv_result?; + } else { + return Err("channel was closed".into()); + } + Ok(slice.get_mapped_range()) + } +} diff --git a/piet-wgsl/src/main.rs b/piet-wgsl/src/main.rs new file mode 100644 index 0000000..95361cd --- /dev/null +++ b/piet-wgsl/src/main.rs @@ -0,0 +1,77 @@ +// Copyright 2022 Google LLC +// +// 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. + +//! A simple application to run a compute shader. + +use std::{fs::File, io::BufWriter}; + +use engine::Engine; + +use render::render; +use test_scene::dump_scene_info; +use wgpu::{Device, Queue}; + +mod engine; +mod render; +mod shaders; +mod template; +mod test_scene; + +async fn run() -> Result<(), Box> { + let instance = wgpu::Instance::new(wgpu::Backends::PRIMARY); + let adapter = instance.request_adapter(&Default::default()).await.unwrap(); + let features = adapter.features(); + let (device, queue) = adapter + .request_device( + &wgpu::DeviceDescriptor { + label: None, + features: features & wgpu::Features::TIMESTAMP_QUERY, + limits: Default::default(), + }, + None, + ) + .await?; + let mut engine = Engine::new(); + do_render(&device, &queue, &mut engine).await?; + + Ok(()) +} + +async fn do_render( + device: &Device, + queue: &Queue, + engine: &mut Engine, +) -> Result<(), Box> { + let shaders = shaders::init_shaders(device, engine)?; + let scene = test_scene::gen_test_scene(); + dump_scene_info(&scene); + let (recording, buf) = render(&scene, &shaders); + let downloads = engine.run_recording(&device, &queue, &recording)?; + let mapped = downloads.map(); + device.poll(wgpu::Maintain::Wait); + let buf = mapped.get_mapped(buf).await?; + + let file = File::create("image.png")?; + let w = BufWriter::new(file); + let encoder = png::Encoder::new(w, 1024, 1024); + let mut writer = encoder.write_header()?; + writer.write_image_data(&buf)?; + Ok(()) +} + +fn main() { + pollster::block_on(run()).unwrap(); +} diff --git a/piet-wgsl/src/render.rs b/piet-wgsl/src/render.rs new file mode 100644 index 0000000..3eec030 --- /dev/null +++ b/piet-wgsl/src/render.rs @@ -0,0 +1,100 @@ +//! Take an encoded scene and create a graph to render it + +use bytemuck::{Pod, Zeroable}; +use piet_scene::Scene; + +use crate::{ + engine::{BufProxy, Recording}, + shaders::{self, Shaders}, +}; + +const TAG_MONOID_SIZE: u64 = 12; + +#[repr(C)] +#[derive(Clone, Copy, Zeroable, Pod)] +struct Config { + width_in_tiles: u32, + height_in_tiles: u32, +} + +#[repr(C)] +#[derive(Clone, Copy, Debug, Zeroable, Pod)] +pub struct PathSegment { + origin: [f32; 2], + delta: [f32; 2], + y_edge: f32, + next: u32, +} + +pub fn render(scene: &Scene, shaders: &Shaders) -> (Recording, BufProxy) { + let mut recording = Recording::default(); + let data = scene.data(); + let n_pathtag = data.tag_stream.len(); + let pathtag_padded = align_up(n_pathtag, 4 * shaders::PATHTAG_REDUCE_WG); + let pathtag_wgs = pathtag_padded / (4 * shaders::PATHTAG_REDUCE_WG as usize); + let mut tag_data: Vec = Vec::with_capacity(pathtag_padded); + tag_data.extend(&data.tag_stream); + tag_data.resize(pathtag_padded, 0); + let pathtag_buf = recording.upload(tag_data); + let reduced_buf = BufProxy::new(pathtag_wgs as u64 * TAG_MONOID_SIZE); + // TODO: really only need pathtag_wgs - 1 + recording.dispatch( + shaders.pathtag_reduce, + (pathtag_wgs as u32, 1, 1), + [pathtag_buf, reduced_buf], + ); + + let tagmonoid_buf = + BufProxy::new(pathtag_wgs as u64 * shaders::PATHTAG_REDUCE_WG as u64 * TAG_MONOID_SIZE); + recording.dispatch( + shaders.pathtag_scan, + (pathtag_wgs as u32, 1, 1), + [pathtag_buf, reduced_buf, tagmonoid_buf], + ); + + let path_coarse_wgs = (data.n_pathseg + shaders::PATH_COARSE_WG - 1) / shaders::PATH_COARSE_WG; + // The clone here is kinda BS, think about reducing copies + // Of course, we'll probably end up concatenating into a single scene binding. + let pathdata_buf = recording.upload(data.pathseg_stream.clone()); + //let cubics_buf = BufProxy::new(data.n_pathseg as u64 * 32); + let config = Config { + width_in_tiles: 64, + height_in_tiles: 64, + }; + let config_buf = recording.upload(bytemuck::bytes_of(&config).to_owned()); + // TODO: more principled size calc + let tiles_buf = BufProxy::new(4097 * 8); + let segments_buf = BufProxy::new(256 * 24); + recording.clear_all(tiles_buf); + recording.dispatch( + shaders.path_coarse, + (path_coarse_wgs, 1, 1), + [ + pathtag_buf, + tagmonoid_buf, + pathdata_buf, + config_buf, + tiles_buf, + segments_buf, + ], + ); + recording.dispatch( + shaders.backdrop, + (config.height_in_tiles, 1, 1), + [config_buf, tiles_buf], + ); + let out_buf_size = config.width_in_tiles * config.height_in_tiles * 256; + let out_buf = BufProxy::new(out_buf_size as u64); + recording.dispatch( + shaders.fine, + (config.width_in_tiles, config.height_in_tiles, 1), + [config_buf, tiles_buf, segments_buf, out_buf], + ); + + recording.download(out_buf); + (recording, out_buf) +} + +pub fn align_up(len: usize, alignment: u32) -> usize { + len + (len.wrapping_neg() & alignment as usize - 1) +} diff --git a/piet-wgsl/src/shaders.rs b/piet-wgsl/src/shaders.rs new file mode 100644 index 0000000..04d6cc4 --- /dev/null +++ b/piet-wgsl/src/shaders.rs @@ -0,0 +1,91 @@ +// Copyright 2022 Google LLC +// +// 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. + +//! Load rendering shaders. + +use serde_json::json; +use wgpu::Device; + +use crate::{ + engine::{BindType, Engine, Error, ShaderId}, + template::ShaderTemplate, +}; + +pub const PATHTAG_REDUCE_WG: u32 = 256; +pub const PATH_COARSE_WG: u32 = 256; + +pub struct Shaders { + pub pathtag_reduce: ShaderId, + pub pathtag_scan: ShaderId, + pub path_coarse: ShaderId, + pub backdrop: ShaderId, + pub fine: ShaderId, +} + +pub fn init_shaders(device: &Device, engine: &mut Engine) -> Result { + let shaders = ShaderTemplate::new(); + let pathtag_reduce = engine.add_shader( + device, + shaders.get_shader("pathtag_reduce", &()).into(), + &[BindType::BufReadOnly, BindType::Buffer], + )?; + let pathtag_scan = engine.add_shader( + device, + shaders.get_shader("pathtag_scan", &()).into(), + &[ + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::Buffer, + ], + )?; + let path_coarse_config = json!({"cubics_out": false}); + let path_coarse = engine.add_shader( + device, + shaders + .get_shader("path_coarse", &path_coarse_config) + .into(), + &[ + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::Buffer, + BindType::Buffer, + ], + )?; + let backdrop = engine.add_shader( + device, + shaders.get_shader("backdrop", &()).into(), + &[BindType::BufReadOnly, BindType::Buffer], + )?; + let fine = engine.add_shader( + device, + shaders.get_shader("fine", &()).into(), + &[ + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::BufReadOnly, + BindType::Buffer, + ], + )?; + Ok(Shaders { + pathtag_reduce, + pathtag_scan, + path_coarse, + backdrop, + fine, + }) +} diff --git a/piet-wgsl/src/template.rs b/piet-wgsl/src/template.rs new file mode 100644 index 0000000..3525298 --- /dev/null +++ b/piet-wgsl/src/template.rs @@ -0,0 +1,37 @@ +// Copyright 2022 Google LLC +// +// 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. + +use handlebars::Handlebars; +use serde::Serialize; + +pub struct ShaderTemplate { + handlebars: Handlebars<'static>, +} + +impl ShaderTemplate { + pub fn new() -> ShaderTemplate { + let mut handlebars = Handlebars::new(); + handlebars + .register_templates_directory("twgsl", concat!(env!("CARGO_MANIFEST_DIR"), "/shader")) + .unwrap(); + handlebars.register_escape_fn(handlebars::no_escape); + ShaderTemplate { handlebars } + } + + pub fn get_shader(&self, shader_name: &str, data: &impl Serialize) -> String { + self.handlebars.render(shader_name, data).unwrap() + } +} diff --git a/piet-wgsl/src/test_scene.rs b/piet-wgsl/src/test_scene.rs new file mode 100644 index 0000000..c844a26 --- /dev/null +++ b/piet-wgsl/src/test_scene.rs @@ -0,0 +1,42 @@ +// Copyright 2022 Google LLC +// +// 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. + +use piet_scene::{Affine, Brush, Color, Fill, PathElement, Point, Scene, SceneBuilder}; + +pub fn gen_test_scene() -> Scene { + let mut scene = Scene::default(); + let mut builder = SceneBuilder::for_scene(&mut scene); + let path = [ + PathElement::MoveTo(Point::new(100.0, 100.0)), + PathElement::LineTo(Point::new(500.0, 120.0)), + PathElement::LineTo(Point::new(300.0, 150.0)), + PathElement::LineTo(Point::new(200.0, 260.0)), + PathElement::LineTo(Point::new(150.0, 210.0)), + PathElement::Close, + ]; + let brush = Brush::Solid(Color::rgb8(0x80, 0x80, 0x80)); + builder.fill(Fill::NonZero, Affine::IDENTITY, &brush, None, &path); + scene +} + +pub fn dump_scene_info(scene: &Scene) { + let data = scene.data(); + println!("tags {:?}", data.tag_stream); + println!( + "pathsegs {:?}", + bytemuck::cast_slice::(&data.pathseg_stream) + ); +}