From b6c4963d4cf6edd4b010e944e8471238bcf253f2 Mon Sep 17 00:00:00 2001 From: Raph Levien Date: Mon, 24 Oct 2022 14:53:12 -0700 Subject: [PATCH] Initial commit of piet-wgsl Starting an experimental port to WGSL shader language, using wgpu to run the examples. As of this commit, it's quite hacky and takes some shortcuts, but does render paths to a grayscale texture. --- Cargo.lock | 562 ++++++++++++++++++++++++++ Cargo.toml | 4 + piet-wgsl/.gitignore | 3 + piet-wgsl/Cargo.toml | 20 + piet-wgsl/README.md | 11 + piet-wgsl/shader/backdrop.twgsl | 63 +++ piet-wgsl/shader/config.twgsl | 20 + piet-wgsl/shader/fine.twgsl | 96 +++++ piet-wgsl/shader/path_coarse.twgsl | 328 +++++++++++++++ piet-wgsl/shader/pathtag.twgsl | 59 +++ piet-wgsl/shader/pathtag_reduce.twgsl | 53 +++ piet-wgsl/shader/pathtag_scan.twgsl | 76 ++++ piet-wgsl/shader/segment.twgsl | 22 + piet-wgsl/src/engine.rs | 356 ++++++++++++++++ piet-wgsl/src/main.rs | 78 ++++ piet-wgsl/src/render.rs | 99 +++++ piet-wgsl/src/shaders.rs | 91 +++++ piet-wgsl/src/template.rs | 37 ++ piet-wgsl/src/test_scene.rs | 39 ++ 19 files changed, 2017 insertions(+) create mode 100644 piet-wgsl/.gitignore create mode 100644 piet-wgsl/Cargo.toml create mode 100644 piet-wgsl/README.md create mode 100644 piet-wgsl/shader/backdrop.twgsl create mode 100644 piet-wgsl/shader/config.twgsl create mode 100644 piet-wgsl/shader/fine.twgsl create mode 100644 piet-wgsl/shader/path_coarse.twgsl create mode 100644 piet-wgsl/shader/pathtag.twgsl create mode 100644 piet-wgsl/shader/pathtag_reduce.twgsl create mode 100644 piet-wgsl/shader/pathtag_scan.twgsl create mode 100644 piet-wgsl/shader/segment.twgsl create mode 100644 piet-wgsl/src/engine.rs create mode 100644 piet-wgsl/src/main.rs create mode 100644 piet-wgsl/src/render.rs create mode 100644 piet-wgsl/src/shaders.rs create mode 100644 piet-wgsl/src/template.rs create mode 100644 piet-wgsl/src/test_scene.rs diff --git a/Cargo.lock b/Cargo.lock index acc4ca4..2de38e1 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,24 @@ dependencies = [ "pinot", ] +[[package]] +name = "naga" +version = "0.10.0" +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 +1046,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 +1136,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 +1273,23 @@ dependencies = [ "smallvec", ] +[[package]] +name = "piet-twgsl" +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 +1314,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 +1355,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 +1441,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 +1496,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 +1522,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 +1537,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 +1601,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 +1646,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 +1758,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 +1788,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 +1806,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 +1854,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 +1978,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..1d4503d 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -8,5 +8,9 @@ members = [ "piet-gpu-hal", "piet-gpu-types", "piet-scene", + "piet-wgsl", "tests", ] + +[patch.crates-io] +naga = { path = "../../dl/naga" } diff --git a/piet-wgsl/.gitignore b/piet-wgsl/.gitignore new file mode 100644 index 0000000..28f0e82 --- /dev/null +++ b/piet-wgsl/.gitignore @@ -0,0 +1,3 @@ +target +.ninja_deps +.ninja_log diff --git a/piet-wgsl/Cargo.toml b/piet-wgsl/Cargo.toml new file mode 100644 index 0000000..9e0150f --- /dev/null +++ b/piet-wgsl/Cargo.toml @@ -0,0 +1,20 @@ +[package] +name = "piet-twgsl" +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..c85aa3d --- /dev/null +++ b/piet-wgsl/shader/backdrop.twgsl @@ -0,0 +1,63 @@ +// 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, +} + +struct Config { + width_in_tiles: u32, + height_in_tiles: u32, +} + +@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; + 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..0bf18fa --- /dev/null +++ b/piet-wgsl/src/main.rs @@ -0,0 +1,78 @@ +// 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::{io::BufWriter, fs::File}; + +use engine::Engine; + +use bytemuck; +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..0672c49 --- /dev/null +++ b/piet-wgsl/src/render.rs @@ -0,0 +1,99 @@ +//! 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.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..fadafb8 --- /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", "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..3ea20f4 --- /dev/null +++ b/piet-wgsl/src/test_scene.rs @@ -0,0 +1,39 @@ +// 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)); +}