alex/vello
1
0
Fork 0
mirror of https://github.com/italicsjenga/vello.git synced 2025-01-10 04:31:30 +11:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #193 from linebender/wgsl

Browse source 
Initial commit of piet-wgsl
This commit is contained in:
Raph Levien 2022-10-26 13:29:31 -07:00 committed by GitHub
parent 4edea5fbc9 911dee23cd
commit b42679c675
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
19 changed files with 2019 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
.gitattributes vendored
View file

@ -1,3 +1,3 @@
**/shader/* linguist-language=glsl
**/shader/gen/* linguist-generated
piet-wgpu/shader/* linguist-language=wgsl

563
Cargo.lock generated
View file

@ -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"

6
Cargo.toml
View file

@ -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" }

20
piet-wgsl/Cargo.toml Normal file
View file

@ -0,0 +1,20 @@
[package]
name = "piet-wgsl"
version = "0.1.0"
authors = ["Raph Levien <raph@google.com>"]
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" }

11
piet-wgsl/README.md Normal file
View file

@ -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.

61
piet-wgsl/shader/backdrop.twgsl Normal file
View file

@ -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<storage> config: Config;
@group(0) @binding(1)
var<storage, read_write> tiles: array<Tile>;
let WG_SIZE = 64u;
var<workgroup> sh_backdrop: array<i32, WG_SIZE>;
// 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<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
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;
}
}

20
piet-wgsl/shader/config.twgsl Normal file
View file

@ -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,
}

96
piet-wgsl/shader/fine.twgsl Normal file
View file

@ -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<storage> config: Config;
@group(0) @binding(1)
var<storage> tiles: array<Tile>;
@group(0) @binding(2)
var<storage> segments: array<Segment>;
// This will become a texture, but keeping things simple for now
@group(0) @binding(3)
var<storage, read_write> output: array<u32>;
let PIXELS_PER_THREAD = 4u;
@compute @workgroup_size(4, 16)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
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<f32, PIXELS_PER_THREAD>;
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<f32>(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;
}

328
piet-wgsl/shader/path_coarse.twgsl Normal file
View file

@ -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<storage> path_tags: array<u32>;
@group(0) @binding(1)
var<storage> tag_monoids: array<TagMonoid>;
// TODO: should probably have single "scene" binding.
@group(0) @binding(2)
var<storage> path_data: array<u32>;
{{#if cubics_out}}
@group(0) @binding(3)
var<storage, read_write> output: array<vec2<f32>>;
{{else}}
{{> config}}
struct Tile {
backdrop: atomic<i32>,
segments: atomic<u32>,
}
{{> segment}}
// Should probably be uniform binding
@group(0) @binding(3)
var<storage> config: Config;
@group(0) @binding(4)
var<storage, read_write> tiles: array<Tile>;
@group(0) @binding(5)
var<storage, read_write> segments: array<Segment>;
{{/if}}
fn read_f32_point(ix: u32) -> vec2<f32> {
let x = bitcast<f32>(path_data[ix]);
let y = bitcast<f32>(path_data[ix + 1u]);
return vec2<f32>(x, y);
}
fn read_i16_point(ix: u32) -> vec2<f32> {
let raw = path_data[ix];
let x = f32(i32(raw << 16u) >> 16u);
let y = f32(i32(raw) >> 16u);
return vec2<f32>(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<f32>, p1: vec2<f32>, p2: vec2<f32>, 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<f32>, p1: vec2<f32>, p2: vec2<f32>, t: f32) -> vec2<f32> {
let mt = 1.0 - t;
return p0 * (mt * mt) + (p1 * (mt * 2.0) + p2 * t) * t;
}
fn eval_cubic(p0: vec2<f32>, p1: vec2<f32>, p2: vec2<f32>, p3: vec2<f32>, t: f32) -> vec2<f32> {
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<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
) {
// 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<f32>;
var p1: vec2<f32>;
var p2: vec2<f32>;
var p3: vec2<f32>;
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<SubdivResult, MAX_QUADS>;
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<f32>;
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}}
}
}

59
piet-wgsl/shader/pathtag.twgsl Normal file
View file

@ -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;
}

53
piet-wgsl/shader/pathtag_reduce.twgsl Normal file
View file

@ -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<storage> path_tags: array<u32>;
@group(0) @binding(1)
var<storage, read_write> reduced: array<TagMonoid>;
let LG_WG_SIZE = 8u;
let WG_SIZE = 256u;
var<workgroup> sh_scratch: array<TagMonoid, WG_SIZE>;
@compute @workgroup_size(256)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
) {
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;
}
}

76
piet-wgsl/shader/pathtag_scan.twgsl Normal file
View file

@ -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<storage> path_tags: array<u32>;
@group(0) @binding(1)
var<storage> reduced: array<TagMonoid>;
@group(0) @binding(2)
var<storage, read_write> tag_monoids: array<TagMonoid>;
let LG_WG_SIZE = 8u;
let WG_SIZE = 256u;
var<workgroup> sh_parent: array<TagMonoid, WG_SIZE>;
// These could be combined?
var<workgroup> sh_monoid: array<TagMonoid, WG_SIZE>;
@compute @workgroup_size(256)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
@builtin(workgroup_id) wg_id: vec3<u32>,
) {
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;
}

22
piet-wgsl/shader/segment.twgsl Normal file
View file

@ -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<f32>,
delta: vec2<f32>,
y_edge: f32,
next: u32,
}

356
piet-wgsl/src/engine.rs Normal file
View file

@ -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<dyn std::error::Error>;
#[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<Shader>,
}
struct Shader {
pipeline: ComputePipeline,
bind_group_layout: BindGroupLayout,
}
#[derive(Default)]
pub struct Recording {
commands: Vec<Command>,
}
#[derive(Clone, Copy)]
pub struct BufProxy {
size: u64,
id: Id,
}
pub enum Command {
Upload(BufProxy, Vec<u8>),
// 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<BufProxy>),
Download(BufProxy),
Clear(BufProxy, u64, Option<NonZeroU64>),
}
#[derive(Default)]
pub struct Downloads {
buf_map: HashMap<Id, Buffer>,
}
/// 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<Id, Buffer>,
}
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<ShaderId, Error> {
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::<Vec<_>>();
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<Downloads, Error> {
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<Vec<u8>>) -> 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<Vec<BufProxy>>,
) {
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<BindGroup, Error> {
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::<Result<Vec<_>, 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<RawMutex, Result<(), BufferAsyncError>>,
),
>,
);
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<BufferView, Error> {
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())
}
}

77
piet-wgsl/src/main.rs Normal file
View file

@ -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<dyn std::error::Error>> {
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<dyn std::error::Error>> {
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();
}

100
piet-wgsl/src/render.rs Normal file
View file

@ -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<u8> = 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)
}

91
piet-wgsl/src/shaders.rs Normal file
View file

@ -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<Shaders, Error> {
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,
})
}

37
piet-wgsl/src/template.rs Normal file
View file

@ -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()
}
}

42
piet-wgsl/src/test_scene.rs Normal file
View file

@ -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::<u8, f32>(&data.pathseg_stream)
);
}