From 95951e72f1aaa4491d7ab742e77a5e4583373149 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Sun, 18 Aug 2024 14:34:18 +0800 Subject: [PATCH 01/16] initiate webgpu Signed-off-by: GopherJ --- Cargo.lock | 675 ++++++++++++++++++++++++++++++++++++++++++---- miden/Cargo.toml | 2 + prover/Cargo.toml | 8 +- 3 files changed, 635 insertions(+), 50 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index db15dfcc0..68a2dcf78 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -17,6 +17,18 @@ version = "1.0.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "f26201604c87b1e01bd3d98f8d5d9a8fcbb815e8cedb41ffccbeb4bf593a35fe" +[[package]] +name = "ahash" +version = "0.8.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e89da841a80418a9b391ebaea17f5c112ffaaa96f621d2c285b5174da76b9011" +dependencies = [ + "cfg-if", + "once_cell", + "version_check", + "zerocopy", +] + [[package]] name = "aho-corasick" version = "1.1.3" @@ -26,6 +38,21 @@ dependencies = [ "memchr", ] +[[package]] +name = "allocator-api2" +version = "0.2.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5c6cb57a04249c6480766f7f7cef5467412af1490f8d1e243141daddada3264f" + +[[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 = "anes" version = "0.1.6" @@ -98,9 +125,9 @@ checksum = "9d151e35f61089500b617991b791fc8bfd237ae50cd5950803758a179b41e67a" [[package]] name = "arrayvec" -version = "0.7.4" +version = "0.7.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "96d30a06541fbafbc7f82ed10c06164cfbd2c401138f6addd8404629c4b16711" +checksum = "7c02d123df017efcdfbd739ef81735b36c5ba83ec3c59c80a9d7ecc718f92e50" [[package]] name = "ascii-canvas" @@ -111,6 +138,15 @@ dependencies = [ "term", ] +[[package]] +name = "ash" +version = "0.38.0+1.3.281" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f" +dependencies = [ + "libloading", +] + [[package]] name = "assert_cmd" version = "2.0.16" @@ -163,7 +199,16 @@ version = "0.5.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0700ddab506f33b20a03b13996eccd309a48e5ff77d0d95926aa0210fb4e95f1" dependencies = [ - "bit-vec", + "bit-vec 0.6.3", +] + +[[package]] +name = "bit-set" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f0481a0e032742109b1133a095184ee93d88f3dc9e0d28a5d033dc77a073f44f" +dependencies = [ + "bit-vec 0.7.0", ] [[package]] @@ -172,6 +217,12 @@ version = "0.6.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "349f9b6a179ed607305526ca489b34ad0a41aed5f7980fa90eb03160b69598fb" +[[package]] +name = "bit-vec" +version = "0.7.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d2c54ff287cfc0a34f38a6b832ea1bd8e448a330b3e40a50859e6488bee07f22" + [[package]] name = "bitflags" version = "1.3.2" @@ -229,6 +280,12 @@ version = "3.16.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "79296716171880943b8470b5f8d03aa55eb2e645a4874bdbb28adb49162e012c" +[[package]] +name = "bytemuck" +version = "1.17.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6fd4c6dcc3b0aea2f5c0b4b82c2b15fe39ddbc76041a310848f4706edf76bb31" + [[package]] name = "byteorder" version = "1.5.0" @@ -243,12 +300,13 @@ checksum = "37b2a672a2cb129a2e41c10b1224bb368f9f37a2b16b612598138befd7b37eb5" [[package]] name = "cc" -version = "1.1.10" +version = "1.1.13" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e9e8aabfac534be767c909e0690571677d49f41bd8465ae876fe043d52ba5292" +checksum = "72db2f7947ecee9b03b510377e8bb9077afa27176fdbff55c51027e976fdcc48" dependencies = [ "jobserver", "libc", + "shlex", ] [[package]] @@ -257,6 +315,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 = "ciborium" version = "0.2.2" @@ -286,9 +350,9 @@ dependencies = [ [[package]] name = "clap" -version = "4.5.15" +version = "4.5.16" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "11d8838454fda655dafd3accb2b6e2bea645b9e4078abe84a22ceb947235c5cc" +checksum = "ed6719fffa43d0d87e5fd8caeab59be1554fb028cd30edc88fc4369b17971019" dependencies = [ "clap_builder", "clap_derive", @@ -315,7 +379,7 @@ dependencies = [ "heck", "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -333,12 +397,53 @@ dependencies = [ "error-code", ] +[[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 = "colorchoice" version = "1.0.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d3fd119d74b830634cea2a0f58bbd0d54540518a14397557951e79340abc28c0" +[[package]] +name = "com" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7e17887fd17353b65b1b2ef1c526c83e26cd72e74f598a8dc1bee13a48f3d9f6" +dependencies = [ + "com_macros", +] + +[[package]] +name = "com_macros" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d375883580a668c7481ea6631fc1a8863e33cc335bf56bfad8d7e6d4b04b13a5" +dependencies = [ + "com_macros_support", + "proc-macro2", + "syn 1.0.109", +] + +[[package]] +name = "com_macros_support" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ad899a1087a9296d5644792d7cb72b8e34c1bec8e7d4fbc002230169a6e8710c" +dependencies = [ + "proc-macro2", + "quote", + "syn 1.0.109", +] + [[package]] name = "constant_time_eq" version = "0.3.0" @@ -458,6 +563,17 @@ dependencies = [ "typenum", ] +[[package]] +name = "d3d12" +version = "22.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bdbd1f579714e3c809ebd822c81ef148b1ceaeb3d535352afc73fd0c4c6a0017" +dependencies = [ + "bitflags 2.6.0", + "libloading", + "winapi", +] + [[package]] name = "diff" version = "0.1.13" @@ -513,6 +629,15 @@ version = "0.3.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "fea41bba32d969b513997752735605054bc0dfa92b4c56bf1189f2e174be7a10" +[[package]] +name = "document-features" +version = "0.2.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "cb6969eaabd2421f8a2775cfd2471a2b634372b4a25d41e3bd647b79912850a0" +dependencies = [ + "litrs", +] + [[package]] name = "either" version = "1.13.0" @@ -592,6 +717,18 @@ dependencies = [ "num-traits", ] +[[package]] +name = "flume" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "55ac459de2512911e4b674ce33cf20befaba382d05b62b008afc1c8b57cbf181" +dependencies = [ + "futures-core", + "futures-sink", + "nanorand", + "spin", +] + [[package]] name = "fnv" version = "1.0.7" @@ -616,7 +753,7 @@ checksum = "1a5c6c585bc94aaf2c7b51dd4c2ba22680844aba4c687be581871a6f518c5742" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -696,7 +833,7 @@ dependencies = [ "libc", "log", "rustversion", - "windows", + "windows 0.58.0", ] [[package]] @@ -716,8 +853,10 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c4567c8db10ae91089c99af84c68c38da3ec2f087c3f82960bcdbf3656b6f4d7" dependencies = [ "cfg-if", + "js-sys", "libc", "wasi", + "wasm-bindgen", ] [[package]] @@ -726,12 +865,96 @@ version = "0.29.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "40ecd4077b5ae9fd2e9e169b102c6c330d0605168eb0e8bf79952b256dbefffd" +[[package]] +name = "gl_generator" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1a95dfc23a2b4a9a2f5ab41d194f8bfda3cabec42af4e39f08c339eb2a0c124d" +dependencies = [ + "khronos_api", + "log", + "xml-rs", +] + [[package]] name = "glob" version = "0.3.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "d2fabcfbdc87f4758337ca535fb41a6d701b65693ce38287d856d1674551ec9b" +[[package]] +name = "glow" +version = "0.13.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bd348e04c43b32574f2de31c8bb397d96c9fcfa1371bd4ca6d8bdc464ab121b1" +dependencies = [ + "js-sys", + "slotmap", + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "glutin_wgl_sys" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0a4e1951bbd9434a81aa496fe59ccc2235af3820d27b85f9314e279609211e2c" +dependencies = [ + "gl_generator", +] + +[[package]] +name = "gpu-alloc" +version = "0.6.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fbcd2dba93594b227a1f57ee09b8b9da8892c34d55aa332e034a228d0fe6a171" +dependencies = [ + "bitflags 2.6.0", + "gpu-alloc-types", +] + +[[package]] +name = "gpu-alloc-types" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "98ff03b468aa837d70984d55f5d3f846f6ec31fe34bbb97c4f85219caeee1ca4" +dependencies = [ + "bitflags 2.6.0", +] + +[[package]] +name = "gpu-allocator" +version = "0.26.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdd4240fc91d3433d5e5b0fc5b67672d771850dc19bbee03c1381e19322803d7" +dependencies = [ + "log", + "presser", + "thiserror", + "winapi", + "windows 0.52.0", +] + +[[package]] +name = "gpu-descriptor" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c08c1f623a8d0b722b8b99f821eb0ba672a1618f0d3b16ddbee1cedd2dd8557" +dependencies = [ + "bitflags 2.6.0", + "gpu-descriptor-types", + "hashbrown", +] + +[[package]] +name = "gpu-descriptor-types" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdf242682df893b86f33a73828fb09ca4b2d3bb6cc95249707fc684d27484b91" +dependencies = [ + "bitflags 2.6.0", +] + [[package]] name = "half" version = "2.4.1" @@ -747,6 +970,25 @@ name = "hashbrown" version = "0.14.5" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e5274423e17b7c9fc20b6e7e208532f9b19825d82dfd615708b70edd83df41f1" +dependencies = [ + "ahash", + "allocator-api2", +] + +[[package]] +name = "hassle-rs" +version = "0.11.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "af2a7e73e1f34c48da31fb668a907f250794837e08faa144fd24f0b8b741e890" +dependencies = [ + "bitflags 2.6.0", + "com", + "libc", + "libloading", + "thiserror", + "widestring", + "winapi", +] [[package]] name = "heck" @@ -756,9 +998,9 @@ checksum = "2304e00983f87ffb38b55b444b5e3b60a884b5d30c0fca7d82fe33449bbe55ea" [[package]] name = "hermit-abi" -version = "0.3.9" +version = "0.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d231dfb89cfffdbc30e7fc41579ed6066ad03abda9e567ccafae602b97ec5024" +checksum = "fbf6a919d6cf397374f7dfeeea91d974c7c0a7221d0d0f4f20d859d329e53fcc" [[package]] name = "hex" @@ -766,6 +1008,12 @@ version = "0.4.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7f24254aa9a54b5c858eaee2f5bccdb46aaf0e486a595ed5fd8f86ba55232a70" +[[package]] +name = "hexf-parse" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" + [[package]] name = "indenter" version = "0.3.3" @@ -774,9 +1022,9 @@ checksum = "ce23b50ad8242c51a442f3ff322d56b02f08852c77e4c0b4d3fd684abc89c683" [[package]] name = "indexmap" -version = "2.3.0" +version = "2.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "de3fc2e30ba82dd1b3911c8de1ffc143c74a914a14e99514d7637e3099df5ea0" +checksum = "93ead53efc7ea8ed3cfb0c79fc8023fbb782a5432b52830b6518941cebe6505c" dependencies = [ "equivalent", "hashbrown", @@ -784,9 +1032,9 @@ dependencies = [ [[package]] name = "is-terminal" -version = "0.4.12" +version = "0.4.13" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f23ff5ef2b80d608d61efee834934d862cd92461afc0560dedf493e4c033738b" +checksum = "261f68e344040fbd0edea105bef17c66edf46f984ddb1115b775ce31be948f4b" dependencies = [ "hermit-abi", "libc", @@ -829,6 +1077,12 @@ version = "1.0.11" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "49f1f14873335454500d59611f1cf4a4b0f786f9ac11f4312a78e4cf2566695b" +[[package]] +name = "jni-sys" +version = "0.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8eaf4bc02d17cbdd7ff4c7438cafcdf7fb9a4613313ad11b4f8fefe7d3fa0130" + [[package]] name = "jobserver" version = "0.1.32" @@ -856,6 +1110,23 @@ dependencies = [ "cpufeatures", ] +[[package]] +name = "khronos-egl" +version = "6.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76" +dependencies = [ + "libc", + "libloading", + "pkg-config", +] + +[[package]] +name = "khronos_api" +version = "3.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc" + [[package]] name = "lalrpop" version = "0.20.2" @@ -863,7 +1134,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "55cb077ad656299f160924eb2912aa147d7339ea7d69e1b5517326fdcec3c1ca" dependencies = [ "ascii-canvas", - "bit-set", + "bit-set 0.5.3", "ena", "itertools 0.11.0", "lalrpop-util", @@ -891,9 +1162,19 @@ checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe" [[package]] name = "libc" -version = "0.2.155" +version = "0.2.157" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "97b3888a4aecf77e811145cadf6eef5901f4782c53886191b2f693f24761847c" +checksum = "374af5f94e54fa97cf75e945cce8a6b201e88a1a07e688b47dfd2a59c66dbd86" + +[[package]] +name = "libloading" +version = "0.8.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4979f22fdb869068da03c9f7528f8297c6fd2606bc3a4affe42e6a823fdb8da4" +dependencies = [ + "cfg-if", + "windows-targets 0.52.6", +] [[package]] name = "libm" @@ -917,6 +1198,12 @@ version = "0.4.14" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "78b3ae25bc7c8c38cec158d1f2757ee79e9b3740fbc7ccf0e59e4b08d793fa89" +[[package]] +name = "litrs" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b4ce301924b7887e9d637144fdade93f9dfff9b60981d4ac161db09720d39aa5" + [[package]] name = "lock_api" version = "0.4.12" @@ -998,6 +1285,21 @@ dependencies = [ "paste", ] +[[package]] +name = "metal" +version = "0.29.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7ecfd3296f8c56b7c1f6fbac3c71cefa9d78ce009850c45000015f206dc7fa21" +dependencies = [ + "bitflags 2.6.0", + "block", + "core-graphics-types", + "foreign-types", + "log", + "objc", + "paste", +] + [[package]] name = "miden-air" version = "0.11.0" @@ -1080,11 +1382,14 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.2.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ade33603aa2eaf78c6f06fd60f4dfe22b7ae1f5606698e386baf71eb9d246d50" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b0d809bd6c02a38c8075575edc9190c5ac7f0dd9#b0d809bd6c02a38c8075575edc9190c5ac7f0dd9" dependencies = [ - "metal", + "bytemuck", + "elsa", + "flume", + "metal 0.27.0", "once_cell", + "wgpu", "winter-math", ] @@ -1112,7 +1417,7 @@ dependencies = [ "supports-color", "supports-hyperlinks", "supports-unicode", - "syn", + "syn 2.0.75", "terminal_size", "textwrap", "trybuild", @@ -1127,7 +1432,7 @@ checksum = "1cc759f0a2947acae217a2f32f722105cacc57d17d5f93bc16362142943a4edd" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -1212,7 +1517,7 @@ checksum = "0ee4176a0f2e7d29d2a8ee7e60b6deb14ce67a20e94c3e2c7275cdb8804e1862" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -1264,6 +1569,45 @@ dependencies = [ "adler", ] +[[package]] +name = "naga" +version = "22.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8bd5a652b6faf21496f2cfd88fc49989c8db0825d1f6746b1a71a6ede24a63ad" +dependencies = [ + "arrayvec", + "bit-set 0.6.0", + "bitflags 2.6.0", + "cfg_aliases", + "codespan-reporting", + "hexf-parse", + "indexmap", + "log", + "rustc-hash", + "spirv", + "termcolor", + "thiserror", + "unicode-xid", +] + +[[package]] +name = "nanorand" +version = "0.7.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6a51313c5820b0b02bd422f4b44776fbf47961755c74ce64afc73bfad10226c3" +dependencies = [ + "getrandom", +] + +[[package]] +name = "ndk-sys" +version = "0.5.0+25.2.9519653" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8c196769dd60fd4f363e11d948139556a344e79d451aeb2fa2fd040738ef7691" +dependencies = [ + "jni-sys", +] + [[package]] name = "new_debug_unreachable" version = "1.0.6" @@ -1338,7 +1682,7 @@ checksum = "ed3955f1a9c7c0c15e092f9c887db08b1fc683305fdf6eb6684f22555355e202" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -1494,6 +1838,12 @@ version = "0.1.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8b870d8c151b6f2fb93e84a13146138f05d02ed11c7e7c54f8826aaaf7c9f184" +[[package]] +name = "pkg-config" +version = "0.3.30" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d231b230927b5e4ad203db57bbcbee2802f6bce620b1e4a9024a07d94e2907ec" + [[package]] name = "plotters" version = "0.3.6" @@ -1573,6 +1923,12 @@ dependencies = [ "termtree", ] +[[package]] +name = "presser" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa" + [[package]] name = "pretty_assertions" version = "1.4.0" @@ -1592,14 +1948,20 @@ dependencies = [ "unicode-ident", ] +[[package]] +name = "profiling" +version = "1.0.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "43d84d1d7a6ac92673717f9f6d1518374ef257669c24ebc5ac25d5033828be58" + [[package]] name = "proptest" version = "1.5.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b4c2511913b88df1637da85cc8d96ec8e43a3f8bb8ccb71ee1ac240d6f3df58d" dependencies = [ - "bit-set", - "bit-vec", + "bit-set 0.5.3", + "bit-vec 0.6.3", "bitflags 2.6.0", "lazy_static", "num-traits", @@ -1666,6 +2028,18 @@ dependencies = [ "rand_core", ] +[[package]] +name = "range-alloc" +version = "0.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c8a99fddc9f0ba0a85884b8d14e3592853e787d581ca1816c91349b10e4eeab" + +[[package]] +name = "raw-window-handle" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539" + [[package]] name = "rayon" version = "1.10.0" @@ -1750,12 +2124,24 @@ version = "0.8.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7a66a03ae7c801facd77a29370b4faec201768915ac14a721ba36f20bc9c209b" +[[package]] +name = "renderdoc-sys" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832" + [[package]] name = "rustc-demangle" version = "0.1.24" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "719b953e2095829ee67db738b3bfa9fa368c94900df327b3f07fe6e794d2fe1f" +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + [[package]] name = "rustc_version" version = "0.2.3" @@ -1874,29 +2260,29 @@ checksum = "388a1df253eca08550bef6c72392cfe7c30914bf41df5269b68cbd6ff8f570a3" [[package]] name = "serde" -version = "1.0.207" +version = "1.0.208" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5665e14a49a4ea1b91029ba7d3bca9f299e1f7cfa194388ccc20f14743e784f2" +checksum = "cff085d2cb684faa248efb494c39b68e522822ac0de72ccf08109abde717cfb2" dependencies = [ "serde_derive", ] [[package]] name = "serde_derive" -version = "1.0.207" +version = "1.0.208" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "6aea2634c86b0e8ef2cfdc0c340baede54ec27b1e46febd7f80dffb2aa44a00e" +checksum = "24008e81ff7613ed8e5ba0cfaf24e2c2f1e5b8a0495711e44fcd4882fca62bcf" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] name = "serde_json" -version = "1.0.124" +version = "1.0.125" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "66ad62847a56b3dba58cc891acd13884b9c61138d330c0d7b6181713d4fce38d" +checksum = "83c8e735a073ccf5be70aa8066aa984eaf2fa000db6c8d0100ae605b366d31ed" dependencies = [ "itoa", "memchr", @@ -1943,12 +2329,27 @@ dependencies = [ "lazy_static", ] +[[package]] +name = "shlex" +version = "1.3.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64" + [[package]] name = "siphasher" version = "0.3.11" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "38b58827f4464d87d377d175e90bf58eb00fd8716ff0a62f80356b5e61555d0d" +[[package]] +name = "slotmap" +version = "1.0.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dbff4acf519f630b3a3ddcfaea6c06b42174d9a44bc70c620e9ed1649d58b82a" +dependencies = [ + "version_check", +] + [[package]] name = "smallvec" version = "1.13.2" @@ -1966,6 +2367,18 @@ name = "spin" version = "0.9.8" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6980e8d7511241f8acf4aebddbb1ff938df5eebe98691418c4468d0b72a96a67" +dependencies = [ + "lock_api", +] + +[[package]] +name = "spirv" +version = "0.3.0+sdk-1.3.268.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "eda41003dc44290527a59b13432d4a0379379fa074b70174882adfbdfd917844" +dependencies = [ + "bitflags 2.6.0", +] [[package]] name = "stable_deref_trait" @@ -1973,6 +2386,12 @@ version = "1.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a8f112729512f8e442d81f95a8a7ddf2b7c6b8a1a6f509a95864142b30cab2d3" +[[package]] +name = "static_assertions" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" + [[package]] name = "string_cache" version = "0.8.7" @@ -2024,9 +2443,20 @@ checksum = "b7401a30af6cb5818bb64852270bb722533397edcfc7344954a38f420819ece2" [[package]] name = "syn" -version = "2.0.74" +version = "1.0.109" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1fceb41e3d546d0bd83421d3409b1460cc7444cd389341a4c880fe7a042cb3d7" +checksum = "72b64191b275b66ffe2469e8af2c1cfe3bafa67b529ead792a6d0160888b4237" +dependencies = [ + "proc-macro2", + "quote", + "unicode-ident", +] + +[[package]] +name = "syn" +version = "2.0.75" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f6af063034fc1935ede7be0122941bafa9bacb949334d090b77ca98b5817c7d9" dependencies = [ "proc-macro2", "quote", @@ -2100,7 +2530,7 @@ dependencies = [ "cfg-if", "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2111,7 +2541,7 @@ checksum = "5c89e72a01ed4c579669add59014b9a524d609c0c88c6a585ce37485879f6ffb" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", "test-case-core", ] @@ -2143,7 +2573,7 @@ checksum = "a4558b58466b9ad7ca0f102865eccc95938dca1a74a856f2b57b6629050da261" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2228,7 +2658,7 @@ checksum = "34704c8d6ebcbc939824180af020566b01a7c01f80641264eba0999f6c2b6be7" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2431,10 +2861,22 @@ dependencies = [ "once_cell", "proc-macro2", "quote", - "syn", + "syn 2.0.75", "wasm-bindgen-shared", ] +[[package]] +name = "wasm-bindgen-futures" +version = "0.4.43" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "61e9300f63a621e96ed275155c108eb6f843b6a26d053f122ab69724559dc8ed" +dependencies = [ + "cfg-if", + "js-sys", + "wasm-bindgen", + "web-sys", +] + [[package]] name = "wasm-bindgen-macro" version = "0.2.93" @@ -2453,7 +2895,7 @@ checksum = "afc340c74d9005395cf9dd098506f7f44e38f2b4a21c6aaacf9a105ea5e1e836" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", "wasm-bindgen-backend", "wasm-bindgen-shared", ] @@ -2474,6 +2916,118 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "wgpu" +version = "22.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e1d1c4ba43f80542cf63a0a6ed3134629ae73e8ab51e4b765a67f3aa062eb433" +dependencies = [ + "arrayvec", + "cfg_aliases", + "document-features", + "js-sys", + "log", + "naga", + "parking_lot", + "profiling", + "raw-window-handle", + "smallvec", + "static_assertions", + "wasm-bindgen", + "wasm-bindgen-futures", + "web-sys", + "wgpu-core", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-core" +version = "22.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0348c840d1051b8e86c3bcd31206080c5e71e5933dabd79be1ce732b0b2f089a" +dependencies = [ + "arrayvec", + "bit-vec 0.7.0", + "bitflags 2.6.0", + "cfg_aliases", + "document-features", + "indexmap", + "log", + "naga", + "once_cell", + "parking_lot", + "profiling", + "raw-window-handle", + "rustc-hash", + "smallvec", + "thiserror", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-hal" +version = "22.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f6bbf4b4de8b2a83c0401d9e5ae0080a2792055f25859a02bf9be97952bbed4f" +dependencies = [ + "android_system_properties", + "arrayvec", + "ash", + "bit-set 0.6.0", + "bitflags 2.6.0", + "block", + "cfg_aliases", + "core-graphics-types", + "d3d12", + "glow", + "glutin_wgl_sys", + "gpu-alloc", + "gpu-allocator", + "gpu-descriptor", + "hassle-rs", + "js-sys", + "khronos-egl", + "libc", + "libloading", + "log", + "metal 0.29.0", + "naga", + "ndk-sys", + "objc", + "once_cell", + "parking_lot", + "profiling", + "range-alloc", + "raw-window-handle", + "renderdoc-sys", + "rustc-hash", + "smallvec", + "thiserror", + "wasm-bindgen", + "web-sys", + "wgpu-types", + "winapi", +] + +[[package]] +name = "wgpu-types" +version = "22.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bc9d91f0e2c4b51434dfa6db77846f2793149d8e73f800fa2e41f52b8eac3c5d" +dependencies = [ + "bitflags 2.6.0", + "js-sys", + "web-sys", +] + +[[package]] +name = "widestring" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7219d36b6eac893fa81e84ebe06485e7dcbb616177469b142df14f1f4deb1311" + [[package]] name = "winapi" version = "0.3.9" @@ -2505,13 +3059,32 @@ version = "0.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "712e227841d057c1ee1cd2fb22fa7e5a5461ae8e48fa2ca79ec42cfc1931183f" +[[package]] +name = "windows" +version = "0.52.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e48a53791691ab099e5e2ad123536d0fff50652600abaf43bbf952894110d0be" +dependencies = [ + "windows-core 0.52.0", + "windows-targets 0.52.6", +] + [[package]] name = "windows" version = "0.58.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "dd04d41d93c4992d421894c18c8b43496aa748dd4c081bac0dc93eb0489272b6" dependencies = [ - "windows-core", + "windows-core 0.58.0", + "windows-targets 0.52.6", +] + +[[package]] +name = "windows-core" +version = "0.52.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "33ab640c8d7e35bf8ba19b884ba838ceb4fba93a4e8c65a9059d08afcfc683d9" +dependencies = [ "windows-targets 0.52.6", ] @@ -2536,7 +3109,7 @@ checksum = "2bbd5b46c938e506ecbce286b6628a02171d56153ba733b6c741fc627ec9579b" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2547,7 +3120,7 @@ checksum = "053c4c462dc91d3b1504c6fe5a726dd15e216ba718e84a0e46a88fbe5ded3515" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2779,7 +3352,7 @@ checksum = "7ce0f4161cdde50de809b3869c1cb083a09e92e949428ea28f04c0d64045875c" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] [[package]] @@ -2829,6 +3402,12 @@ dependencies = [ "winter-utils", ] +[[package]] +name = "xml-rs" +version = "0.8.21" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "539a77ee7c0de333dcc6da69b177380a0b81e0dacfa4f7344c465a36871ee601" + [[package]] name = "yansi" version = "0.5.1" @@ -2853,5 +3432,5 @@ checksum = "fa4f8080344d4671fb4e831a13ad1e68092748387dfc4f55e356242fae12ce3e" dependencies = [ "proc-macro2", "quote", - "syn", + "syn 2.0.75", ] diff --git a/miden/Cargo.toml b/miden/Cargo.toml index 70d675245..5723ac4ae 100644 --- a/miden/Cargo.toml +++ b/miden/Cargo.toml @@ -24,6 +24,7 @@ required-features = ["executable"] path = "src/lib.rs" bench = false doctest = false +crate-type = ["cdylib", "rlib"] [[bench]] name = "program_compilation" @@ -58,6 +59,7 @@ executable = [ "dep:tracing-subscriber", ] metal = ["prover/metal", "std"] +webgpu = ["prover/webgpu"] std = ["assembly/std", "processor/std", "prover/std", "verifier/std"] [dependencies] diff --git a/prover/Cargo.toml b/prover/Cargo.toml index b2b038b6e..d527dd37c 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -17,7 +17,8 @@ edition.workspace = true concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] default = ["std"] metal = ["dep:miden-gpu", "dep:elsa", "dep:pollster", "concurrent", "std"] -std = ["air/std", "processor/std", "winter-prover/std"] +webgpu = ["dep:miden-gpu", "miden-gpu/webgpu"] +std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std"] [dependencies] air = { package = "miden-air", path = "../air", version = "0.11", default-features = false } @@ -27,5 +28,8 @@ winter-prover = { package = "winter-prover", version = "0.9", default-features = [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] elsa = { version = "1.9", optional = true } -miden-gpu = { version = "0.2", optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b0d809bd6c02a38c8075575edc9190c5ac7f0dd9", default-features = false, optional = true } pollster = { version = "0.3", optional = true } + +[target.'cfg(target_family = "wasm")'.dependencies] +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b0d809bd6c02a38c8075575edc9190c5ac7f0dd9", default-features = false, optional = true } From 873a49ce345c756de8e9ce8d721fd5a7365d187b Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 08:32:05 +0800 Subject: [PATCH 02/16] add webgpu prover Signed-off-by: GopherJ --- Cargo.lock | 57 ++- Cargo.toml | 11 + prover/Cargo.toml | 13 +- prover/src/gpu/mod.rs | 3 + prover/src/gpu/webgpu/mod.rs | 752 +++++++++++++++++++++++++++++++++ prover/src/gpu/webgpu/tests.rs | 227 ++++++++++ prover/src/lib.rs | 27 +- 7 files changed, 1053 insertions(+), 37 deletions(-) create mode 100644 prover/src/gpu/webgpu/mod.rs create mode 100644 prover/src/gpu/webgpu/tests.rs diff --git a/Cargo.lock b/Cargo.lock index 68a2dcf78..19ade88c0 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -163,6 +163,17 @@ dependencies = [ "wait-timeout", ] +[[package]] +name = "async-trait" +version = "0.1.81" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6e0c28dcc82d7c8ead5cb13beb15405b57b8546e93215673ff8ca0349a028107" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.75", +] + [[package]] name = "autocfg" version = "1.3.0" @@ -1354,8 +1365,7 @@ dependencies = [ [[package]] name = "miden-crypto" version = "0.10.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d6fad06fc3af260ed3c4235821daa2132813d993f96d446856036ae97e9606dd" +source = "git+https://github.com/GopherJ/miden-crypto?rev=657b4922f3abe3577d1b0de1f6ae6ad52400cb11#657b4922f3abe3577d1b0de1f6ae6ad52400cb11" dependencies = [ "blake3", "cc", @@ -1382,7 +1392,7 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.2.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b0d809bd6c02a38c8075575edc9190c5ac7f0dd9#b0d809bd6c02a38c8075575edc9190c5ac7f0dd9" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b77b4709d69d42ef663ed5ea41f48f8b165f59f2#b77b4709d69d42ef663ed5ea41f48f8b165f59f2" dependencies = [ "bytemuck", "elsa", @@ -1454,12 +1464,14 @@ dependencies = [ name = "miden-prover" version = "0.11.0" dependencies = [ + "async-trait", "elsa", "miden-air", "miden-gpu", "miden-processor", "pollster", "tracing", + "winter-maybe-async 0.9.0 (registry+https://github.com/rust-lang/crates.io-index)", "winter-prover", ] @@ -3302,8 +3314,7 @@ dependencies = [ [[package]] name = "winter-air" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b72f12b88ebb060b52c0e9aece9bb64a9fc38daf7ba689dd5ce63271b456c883" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "libm", "winter-crypto", @@ -3315,8 +3326,7 @@ dependencies = [ [[package]] name = "winter-crypto" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "00fbb724d2d9fbfd3aa16ea27f5e461d4fe1d74b0c9e0ed1bf79e9e2a955f4d5" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "blake3", "sha3", @@ -3327,8 +3337,7 @@ dependencies = [ [[package]] name = "winter-fri" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3ab6077cf4c23c0411f591f4ba29378e27f26acb8cef3c51cadd93daaf6080b3" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "winter-crypto", "winter-math", @@ -3338,8 +3347,7 @@ dependencies = [ [[package]] name = "winter-math" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "004f85bb051ce986ec0b9a2bd90aaf81b83e3c67464becfdf7db31f14c1019ba" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "winter-utils", ] @@ -3355,26 +3363,35 @@ dependencies = [ "syn 2.0.75", ] +[[package]] +name = "winter-maybe-async" +version = "0.9.0" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.75", +] + [[package]] name = "winter-prover" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f17e3dbae97050f58e01ed4f12906e247841575a0518632e052941a1c37468df" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ + "async-trait", "tracing", "winter-air", "winter-crypto", "winter-fri", "winter-math", - "winter-maybe-async", + "winter-maybe-async 0.9.0 (git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e)", "winter-utils", ] [[package]] name = "winter-rand-utils" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f2b827c901ab0c316d89812858ff451d60855c0a5c7ae734b098c62a28624181" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "rand", "winter-utils", @@ -3382,9 +3399,8 @@ dependencies = [ [[package]] name = "winter-utils" -version = "0.9.1" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "0568612a95bcae3c94fb14da2686f8279ca77723dbdf1e97cf3673798faf6485" +version = "0.9.0" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "rayon", ] @@ -3392,8 +3408,7 @@ dependencies = [ [[package]] name = "winter-verifier" version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "324002ade90f21e85599d51a232a80781efc8cb46f511f8bc89f9c5a4eb9cb65" +source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" dependencies = [ "winter-air", "winter-crypto", diff --git a/Cargo.toml b/Cargo.toml index ac69f241e..d9ae4dd8b 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -32,3 +32,14 @@ inherits = "release" debug = true debug-assertions = true overflow-checks = true + +[patch.crates-io] +winter-prover = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-air = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-utils = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-verifier = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-fri = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-crypto = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-math = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-rand-utils = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +miden-crypto = { git = "https://github.com/GopherJ/miden-crypto", rev = "657b4922f3abe3577d1b0de1f6ae6ad52400cb11" } diff --git a/prover/Cargo.toml b/prover/Cargo.toml index d527dd37c..73b721ec7 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -15,21 +15,24 @@ edition.workspace = true [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["std"] +default = ["webgpu"] metal = ["dep:miden-gpu", "dep:elsa", "dep:pollster", "concurrent", "std"] -webgpu = ["dep:miden-gpu", "miden-gpu/webgpu"] -std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std"] +webgpu = ["dep:miden-gpu", "dep:elsa", "miden-gpu/webgpu", "async"] +async = ["maybe-async/async", "winter-prover/async", "winter-prover/async-trait", "dep:async-trait"] +std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std", "tracing/std"] [dependencies] air = { package = "miden-air", path = "../air", version = "0.11", default-features = false } processor = { package = "miden-processor", path = "../processor", version = "0.11", default-features = false } tracing = { version = "0.1", default-features = false, features = ["attributes"] } winter-prover = { package = "winter-prover", version = "0.9", default-features = false } +maybe-async = { package = "winter-maybe-async", version = "0.9", default-features = false } +async-trait = { version = "0.1", optional = true } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] elsa = { version = "1.9", optional = true } -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b0d809bd6c02a38c8075575edc9190c5ac7f0dd9", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b77b4709d69d42ef663ed5ea41f48f8b165f59f2", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b0d809bd6c02a38c8075575edc9190c5ac7f0dd9", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b77b4709d69d42ef663ed5ea41f48f8b165f59f2", default-features = false, optional = true } diff --git a/prover/src/gpu/mod.rs b/prover/src/gpu/mod.rs index 253db30f7..c57b8b9cc 100644 --- a/prover/src/gpu/mod.rs +++ b/prover/src/gpu/mod.rs @@ -1,2 +1,5 @@ #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] pub mod metal; + +#[cfg(all(feature = "webgpu"))] +pub mod webgpu; diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs new file mode 100644 index 000000000..1189fd3e0 --- /dev/null +++ b/prover/src/gpu/webgpu/mod.rs @@ -0,0 +1,752 @@ +//! This module contains GPU acceleration logic for Browser +//! For now, the logic is limited to GPU accelerating trace and constraint commitments, +//! using the RPO 256 or RPX 256 hash functions. + +#[cfg(feature = "std")] +use std::{boxed::Box, marker::PhantomData, time::Instant, vec::Vec}; + +#[cfg(not(feature = "std"))] +use core::marker::PhantomData; + +#[cfg(not(feature = "std"))] +extern crate alloc; + +#[cfg(not(feature = "std"))] +use alloc::vec; +#[cfg(not(feature = "std"))] +use alloc::vec::Vec; +#[cfg(not(feature = "std"))] +use alloc::boxed::Box; + +use elsa::FrozenVec; +use maybe_async::maybe_async; + +use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; +use miden_gpu::{ + webgpu::{ + build_merkle_tree, get_dispatch_linear, get_wgpu_helper, init_wgpu_helper, RowHasher, + WebGpuHelper, + }, + HashFn, +}; +use processor::{ + crypto::{ElementHasher, Hasher}, + ONE, +}; +use tracing::{event, Level}; +use winter_prover::{ + crypto::{Digest, MerkleTree}, + matrix::{get_evaluation_offsets, ColMatrix, RowMatrix, Segment}, + proof::Queries, + CompositionPoly, CompositionPolyTrace, ConstraintCommitment, ConstraintCompositionCoefficients, + DefaultConstraintEvaluator, EvaluationFrame, Prover, StarkDomain, TraceInfo, TraceLde, + TracePolyTable, +}; + +use crate::{ + crypto::{RandomCoin, Rpo256}, + math::fft, + ExecutionProver, ExecutionTrace, Felt, FieldElement, ProcessorAir, PublicInputs, + WinterProofOptions, +}; + +#[cfg(test)] +mod tests; + +// CONSTANTS +// ================================================================================================ + +// The Rate for RPO and RPX is the same +const RATE: usize = Rpo256::RATE_RANGE.end - Rpo256::RATE_RANGE.start; +const DIGEST_SIZE: usize = Rpo256::DIGEST_RANGE.end - Rpo256::DIGEST_RANGE.start; + +// WebGPU RPO/RPX PROVER +// ================================================================================================ + +/// Wraps an [ExecutionProver] and provides GPU acceleration for building trace commitments. +pub(crate) struct WebGPUExecutionProver +where + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, + R: RandomCoin + Send, +{ + pub execution_prover: ExecutionProver, + pub webgpu_hash_fn: HashFn, + phantom_data: PhantomData, +} + +impl WebGPUExecutionProver +where + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, + R: RandomCoin + Send, +{ + pub fn new(execution_prover: ExecutionProver, hash_fn: HashFn) -> Self { + WebGPUExecutionProver { + execution_prover, + webgpu_hash_fn: hash_fn, + phantom_data: PhantomData, + } + } + + fn build_aligned_segement( + polys: &ColMatrix, + poly_offset: usize, + offsets: &[Felt], + twiddles: &[Felt], + ) -> Segment + where + E: FieldElement, + { + let poly_size = polys.num_rows(); + let domain_size = offsets.len(); + assert!(domain_size.is_power_of_two()); + assert!(domain_size > poly_size); + assert_eq!(poly_size, twiddles.len() * 2); + assert!(poly_offset < polys.num_base_cols()); + + // allocate memory for the segment + // let data = if polys.num_base_cols() - poly_offset >= N { + // // if we will fill the entire segment, we allocate uninitialized memory + // unsafe { page_aligned_uninit_vector(domain_size) } + // } else { + // but if some columns in the segment will remain unfilled, we allocate memory + // initialized to zeros to make sure we don't end up with memory with + // undefined values + let data = alloc::vec![[E::BaseField::ZERO; N]; domain_size]; + // }; + + Segment::new_with_buffer(data, polys, poly_offset, offsets, twiddles) + } + + fn build_aligned_segements( + polys: &ColMatrix, + twiddles: &[Felt], + offsets: &[Felt], + ) -> Vec> + where + E: FieldElement, + { + assert!(N > 0, "batch size N must be greater than zero"); + debug_assert_eq!(polys.num_rows(), twiddles.len() * 2); + debug_assert_eq!(offsets.len() % polys.num_rows(), 0); + + let num_segments = if polys.num_base_cols() % N == 0 { + polys.num_base_cols() / N + } else { + polys.num_base_cols() / N + 1 + }; + + (0..num_segments) + .map(|i| Self::build_aligned_segement(polys, i * N, offsets, twiddles)) + .collect() + } +} + +#[maybe_async] +impl Prover for WebGPUExecutionProver +where + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, + R: RandomCoin + Send, +{ + type BaseField = Felt; + type Air = ProcessorAir; + type Trace = ExecutionTrace; + type HashFn = H; + type RandomCoin = R; + type TraceLde> = WebGPUTraceLde; + type ConstraintEvaluator<'a, E: FieldElement> = + DefaultConstraintEvaluator<'a, ProcessorAir, E>; + + fn get_pub_inputs(&self, trace: &ExecutionTrace) -> PublicInputs { + self.execution_prover.get_pub_inputs(trace) + } + + fn options(&self) -> &WinterProofOptions { + self.execution_prover.options() + } + + async fn new_trace_lde>( + &self, + trace_info: &TraceInfo, + main_trace: &ColMatrix, + domain: &StarkDomain, + ) -> (Self::TraceLde, TracePolyTable) { + WebGPUTraceLde::new(trace_info, main_trace, domain, self.webgpu_hash_fn).await + } + + async fn new_evaluator<'a, E: FieldElement>( + &self, + air: &'a ProcessorAir, + aux_rand_elements: Option>, + composition_coefficients: ConstraintCompositionCoefficients, + ) -> Self::ConstraintEvaluator<'a, E> { + self.execution_prover + .new_evaluator(air, aux_rand_elements, composition_coefficients) + .await + } + + /// Evaluates constraint composition polynomial over the LDE domain and builds a commitment + /// to these evaluations. + /// + /// The evaluation is done by evaluating each composition polynomial column over the LDE + /// domain. + /// + /// The commitment is computed by hashing each row in the evaluation matrix, and then building + /// a Merkle tree from the resulting hashes. + /// + /// The composition polynomial columns are evaluated on the CPU. Afterwards the commitment + /// is computed on the GPU. + /// + /// ```text + /// ───────────────────────────────────────────────────── + /// ┌───┐ ┌───┐ + /// CPU: ... ─┤fft├─┤fft├─┐ ┌─ ... + /// └───┘ └───┘ │ │ + /// ╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴┼╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴┼╴╴╴╴╴╴ + /// │ ┌──────────┐ ┌──────────┐ │ + /// GPU: └─┤ hash ├─┤ hash ├─┘ + /// └──────────┘ └──────────┘ + /// ────┼────────┼────────┼────────┼────────┼────────┼─── + /// t=n t=n+1 t=n+2 t=n+3 t=n+4 t=n+5 + /// ``` + async fn build_constraint_commitment>( + &self, + composition_poly_trace: CompositionPolyTrace, + num_trace_poly_columns: usize, + domain: &StarkDomain, + ) -> (ConstraintCommitment, CompositionPoly) { + // evaluate composition polynomial columns over the LDE domain + #[cfg(feature = "std")] + let now = Instant::now(); + let composition_poly = + CompositionPoly::new(composition_poly_trace, domain, num_trace_poly_columns); + let blowup = domain.trace_to_lde_blowup(); + let offsets = + get_evaluation_offsets::(composition_poly.column_len(), blowup, domain.offset()); + let segments = Self::build_aligned_segements( + composition_poly.data(), + domain.trace_twiddles(), + &offsets, + ); + #[cfg(feature = "std")] + event!( + Level::INFO, + "Evaluated {} composition polynomial columns over LDE domain (2^{} elements) in {} ms", + composition_poly.num_columns(), + offsets.len().ilog2(), + now.elapsed().as_millis() + ); + let helper = get_wgpu_helper().unwrap(); + + // build constraint evaluation commitment + #[cfg(feature = "std")] + let now = Instant::now(); + let lde_domain_size = domain.lde_domain_size(); + let num_base_columns = + composition_poly.num_columns() * ::EXTENSION_DEGREE; + let rpo_requires_padding = num_base_columns % RATE != 0; + let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); + let mut row_hasher = + RowHasher::new(&helper, lde_domain_size, rpo_requires_padding, self.webgpu_hash_fn); + let mut rpo_padded_segment: Vec<[Felt; RATE]>; + for (segment_idx, segment) in segments.iter().enumerate() { + // check if the segment requires padding + if rpo_padded_segment_idx.map_or(false, |pad_idx| pad_idx == segment_idx) { + // duplicate and modify the last segment with Rpo256's padding + // rule ("1" followed by "0"s). Our segments are already + // padded with "0"s we only need to add the "1"s. + let rpo_pad_column = num_base_columns % RATE; + + rpo_padded_segment = segment + .iter() + .map(|x| { + let mut s = x.clone(); + s[rpo_pad_column] = ONE; + s + }) + .collect(); + // For rpx, skip this step + if self.webgpu_hash_fn == HashFn::Rpo256 { + let rpo_pad_column = num_base_columns % RATE; + rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); + } + row_hasher.update(&helper, &rpo_padded_segment); + assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); + break; + } + row_hasher.update(&helper, segment); + } + let row_hashes = row_hasher.finish(&helper).await.unwrap(); + let tree_nodes = + build_merkle_tree(&helper, &row_hashes, self.webgpu_hash_fn).await.unwrap(); + // aggregate segments at the same time as the GPU generates the merkle tree nodes + let composed_evaluations = RowMatrix::::from_segments(segments, num_base_columns); + let nodes = tree_nodes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); + let leaves = row_hashes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); + let commitment = MerkleTree::::from_raw_parts(nodes, leaves).unwrap(); + let constraint_commitment = ConstraintCommitment::new(composed_evaluations, commitment); + #[cfg(feature = "std")] + event!( + Level::INFO, + "Computed constraint evaluation commitment on the GPU (Merkle tree of depth {}) in {} ms", + constraint_commitment.tree_depth(), + now.elapsed().as_millis() + ); + (constraint_commitment, composition_poly) + } +} + +// TRACE LOW DEGREE EXTENSION (WebGPU) +// ================================================================================================ + +/// Contains all segments of the extended execution trace, the commitments to these segments, the +/// LDE blowup factor, and the [TraceInfo]. +/// +/// Segments are stored in two groups: +/// - Main segment: this is the first trace segment generated by the prover. Values in this segment +/// will always be elements in the base field (even when an extension field is used). +/// - Auxiliary segments: a list of 0 or more segments for traces generated after the prover commits +/// to the first trace segment. Currently, at most 1 auxiliary segment is possible. +pub struct WebGPUTraceLde, H: Hasher> { + // low-degree extension of the main segment of the trace + main_segment_lde: RowMatrix, + // commitment to the main segment of the trace + main_segment_tree: MerkleTree, + // low-degree extensions of the auxiliary segments of the trace + aux_segment_lde: Option>, + // commitment to the auxiliary segments of the trace + aux_segment_tree: Option>, + blowup: usize, + trace_info: TraceInfo, + webgpu_hash_fn: HashFn, +} + +impl< + E: FieldElement, + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, + > WebGPUTraceLde +{ + /// Takes the main trace segment columns as input, interpolates them into polynomials in + /// coefficient form, evaluates the polynomials over the LDE domain, commits to the + /// polynomial evaluations, and creates a new [DefaultTraceLde] with the LDE of the main trace + /// segment and the commitment. + /// + /// Returns a tuple containing a [TracePolyTable] with the trace polynomials for the main trace + /// segment and the new [DefaultTraceLde]. + pub async fn new( + trace_info: &TraceInfo, + main_trace: &ColMatrix, + domain: &StarkDomain, + webgpu_hash_fn: HashFn, + ) -> (Self, TracePolyTable) { + init_wgpu_helper().await; + // extend the main execution trace and build a Merkle tree from the extended trace + let (main_segment_lde, main_segment_tree, main_segment_polys) = + build_trace_commitment(main_trace, domain, webgpu_hash_fn).await; + + let trace_poly_table = TracePolyTable::new(main_segment_polys); + let trace_lde = WebGPUTraceLde { + main_segment_lde, + main_segment_tree, + aux_segment_lde: None, + aux_segment_tree: None, + blowup: domain.trace_to_lde_blowup(), + trace_info: trace_info.clone(), + webgpu_hash_fn, + }; + + (trace_lde, trace_poly_table) + } + + // TEST HELPERS + // -------------------------------------------------------------------------------------------- + + /// Returns number of columns in the main segment of the execution trace. + #[allow(unused)] + #[cfg(test)] + pub fn main_segment_width(&self) -> usize { + self.main_segment_lde.num_cols() + } + + /// Returns a reference to [Matrix] representing the main trace segment. + #[allow(unused)] + #[cfg(test)] + pub fn get_main_segment(&self) -> &RowMatrix { + &self.main_segment_lde + } + + /// Returns the entire trace for the column at the specified index. + #[allow(unused)] + #[cfg(test)] + pub fn get_main_segment_column(&self, col_idx: usize) -> Vec { + (0..self.main_segment_lde.num_rows()) + .map(|row_idx| self.main_segment_lde.get(col_idx, row_idx)) + .collect() + } +} + +#[maybe_async] +impl< + E: FieldElement, + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, + > TraceLde for WebGPUTraceLde +{ + type HashFn = H; + + /// Returns the commitment to the low-degree extension of the main trace segment. + fn get_main_trace_commitment(&self) -> D { + let root_hash = self.main_segment_tree.root(); + *root_hash + } + + /// Takes auxiliary trace segment columns as input, interpolates them into polynomials in + /// coefficient form, evaluates the polynomials over the LDE domain, and commits to the + /// polynomial evaluations. + /// + /// Returns a tuple containing the column polynomials in coefficient from and the commitment + /// to the polynomial evaluations over the LDE domain. + /// + /// # Panics + /// + /// This function will panic if any of the following are true: + /// - the number of rows in the provided `aux_trace` does not match the main trace. + /// - this segment would exceed the number of segments specified by the trace layout. + async fn set_aux_trace( + &mut self, + aux_trace: &ColMatrix, + domain: &StarkDomain, + ) -> (ColMatrix, D) { + // extend the auxiliary trace segment and build a Merkle tree from the extended trace + let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = + build_trace_commitment::(aux_trace, domain, self.webgpu_hash_fn).await; + + assert_eq!( + self.main_segment_lde.num_rows(), + aux_segment_lde.num_rows(), + "the number of rows in the auxiliary segment must be the same as in the main segment" + ); + + // save the lde and commitment + self.aux_segment_lde = Some(aux_segment_lde); + let root_hash = *aux_segment_tree.root(); + self.aux_segment_tree = Some(aux_segment_tree); + + (aux_segment_polys, root_hash) + } + + /// Reads current and next rows from the main trace segment into the specified frame. + fn read_main_trace_frame_into(&self, lde_step: usize, frame: &mut EvaluationFrame) { + // at the end of the trace, next state wraps around and we read the first step again + let next_lde_step = (lde_step + self.blowup()) % self.trace_len(); + + // copy main trace segment values into the frame + frame.current_mut().copy_from_slice(self.main_segment_lde.row(lde_step)); + frame.next_mut().copy_from_slice(self.main_segment_lde.row(next_lde_step)); + } + + /// Reads current and next rows from the auxiliary trace segment into the specified frame. + /// + /// # Panics + /// This currently assumes that there is exactly one auxiliary trace segment, and will panic + /// otherwise. + fn read_aux_trace_frame_into(&self, lde_step: usize, frame: &mut EvaluationFrame) { + // at the end of the trace, next state wraps around and we read the first step again + let next_lde_step = (lde_step + self.blowup()) % self.trace_len(); + + // copy auxiliary trace segment values into the frame + if let Some(mat) = self.aux_segment_lde.as_ref() { + frame.current_mut().copy_from_slice(mat.row(lde_step)); + frame.next_mut().copy_from_slice(mat.row(next_lde_step)); + } + } + + /// Returns trace table rows at the specified positions along with Merkle authentication paths + /// from the commitment root to these rows. + fn query(&self, positions: &[usize]) -> Vec { + // build queries for the main trace segment + let mut result = vec![build_segment_queries( + &self.main_segment_lde, + &self.main_segment_tree, + positions, + )]; + + if let (Some(aux_segment_lde), Some(aux_segment_tree)) = + (&self.aux_segment_lde, &self.aux_segment_tree) + { + result.push(build_segment_queries(aux_segment_lde, aux_segment_tree, positions)); + } + + result + } + + /// Returns the number of rows in the execution trace. + fn trace_len(&self) -> usize { + self.main_segment_lde.num_rows() + } + + /// Returns blowup factor which was used to extend the original execution trace into trace LDE. + fn blowup(&self) -> usize { + self.blowup + } + + /// Populates the provided Lagrange kernel frame starting at the current row (as defined by + /// lde_step). + /// Note that unlike EvaluationFrame, the Lagrange kernel frame includes only the Lagrange + /// kernel column (as opposed to all columns). + fn read_lagrange_kernel_frame_into( + &self, + lde_step: usize, + col_idx: usize, + frame: &mut LagrangeKernelEvaluationFrame, + ) { + if let Some(aux_segment) = self.aux_segment_lde.as_ref() { + let frame = frame.frame_mut(); + frame.truncate(0); + + frame.push(aux_segment.get(col_idx, lde_step)); + + let frame_length = self.trace_info.length().ilog2() as usize + 1; + for i in 0..frame_length - 1 { + let shift = self.blowup() * (1 << i); + let next_lde_step = (lde_step + shift) % self.trace_len(); + + frame.push(aux_segment.get(col_idx, next_lde_step)); + } + } + } + + /// Returns the trace info + fn trace_info(&self) -> &TraceInfo { + &self.trace_info + } +} + +/// Computes a low-degree extension (LDE) of the provided execution trace over the specified +/// domain and builds a commitment to the extended trace. +/// +/// The extension is performed by interpolating each column of the execution trace into a +/// polynomial of degree = trace_length - 1, and then evaluating the polynomial over the LDE +/// domain. +/// +/// Trace commitment is computed by hashing each row of the extended execution trace, and then +/// building a Merkle tree from the resulting hashes. +/// +/// Interpolations and evaluations are computed on the CPU while hashes are simultaneously +/// computed on the GPU: +/// +/// ```text +/// ────────────────────────────────────────────────────── +/// ┌───┐ ┌────┐ ┌───┐ ┌────┐ ┌───┐ +/// CPU: ... ──┤fft├─┬─┤ifft├───┤fft├─┬─┤ifft├───┤fft├─┬─ ... +/// └───┘ │ └────┘ └───┘ │ └────┘ └───┘ │ +/// ╴╴╴╴╴╴╴╴╴╴╴╴╴┼╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴┼╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴╴┼╴╴╴╴╴╴ +/// │ ┌──────────┐ │ ┌──────────┐ │ +/// GPU: └─┤ hash │ └─┤ hash │ └─ ... +/// └──────────┘ └──────────┘ +/// ────┼────────┼────────┼────────┼────────┼────────┼──── +/// t=n t=n+1 t=n+2 t=n+3 t=n+4 t=n+5 +/// ``` +async fn build_trace_commitment< + E: FieldElement, + H: Hasher + ElementHasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, +>( + trace: &ColMatrix, + domain: &StarkDomain, + hash_fn: HashFn, +) -> (RowMatrix, MerkleTree, ColMatrix) { + // interpolate the execution trace + #[cfg(feature = "std")] + let now = Instant::now(); + let inv_twiddles = fft::get_inv_twiddles::(trace.num_rows()); + let trace_polys = trace.columns().map(|col| { + let mut poly = col.to_vec(); + fft::interpolate_poly(&mut poly, &inv_twiddles); + poly + }); + + // extend the execution trace and generate hashes on the gpu + let lde_segments = FrozenVec::new(); + let lde_domain_size = domain.lde_domain_size(); + let num_base_columns = trace.num_base_cols(); + let rpo_requires_padding = num_base_columns % RATE != 0; + let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); + let mut row_hasher = RowHasher::new( + &get_wgpu_helper().unwrap(), + num_base_columns, + rpo_requires_padding, + hash_fn, + ); + let mut rpo_padded_segment: Vec<[Felt; RATE]>; + let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); + let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); + for (segment_idx, segment) in &mut lde_segment_iter { + let segment = lde_segments.push_get(Box::new(segment)); + // check if the segment requires padding + if rpo_padded_segment_idx.map_or(false, |pad_idx| pad_idx == segment_idx) { + // duplicate and modify the last segment with Rpo256's padding + // rule ("1" followed by "0"s). Our segments are already + // padded with "0"s we only need to add the "1"s. + let rpo_pad_column = num_base_columns % RATE; + rpo_padded_segment = segment + .iter() + .map(|x| { + let mut s = x.clone(); + s[rpo_pad_column] = ONE; + s + }) + .collect(); + // skip this in case of Rpx + if hash_fn == HashFn::Rpo256 { + let rpo_pad_column = num_base_columns % RATE; + rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); + } + row_hasher.update(&get_wgpu_helper().unwrap(), &rpo_padded_segment); + assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); + break; + } + row_hasher.update(&get_wgpu_helper().unwrap(), segment); + } + let row_hashes = row_hasher.finish(&get_wgpu_helper().unwrap()).await.unwrap(); + let tree_nodes = build_merkle_tree(&get_wgpu_helper().unwrap(), &row_hashes, hash_fn).await.unwrap(); + // aggregate segments at the same time as the GPU generates the merkle tree nodes + let lde_segments = lde_segments.into_vec().into_iter().map(|p| *p).collect(); + let trace_lde = RowMatrix::from_segments(lde_segments, num_base_columns); + let trace_polys = lde_segment_generator.into_polys().unwrap(); + let nodes = tree_nodes.into_iter().map(|dig| D::from(&dig)).collect(); + let leaves = row_hashes.into_iter().map(|dig| D::from(&dig)).collect(); + let trace_tree = MerkleTree::from_raw_parts(nodes, leaves).unwrap(); + #[cfg(feature = "std")] + event!( + Level::INFO, + "Extended (on CPU) and committed (on GPU) to an execution trace of {} columns from 2^{} to 2^{} steps in {} ms", + trace_polys.num_cols(), + trace_polys.num_rows().ilog2(), + trace_lde.num_rows().ilog2(), + now.elapsed().as_millis() + ); + + (trace_lde, trace_tree, trace_polys) +} + +// SEGMENT GENERATOR +// ================================================================================================ + +struct SegmentGenerator<'a, E, I, const N: usize> +where + E: FieldElement, + I: IntoIterator>, +{ + poly_iter: I::IntoIter, + polys: Option>, + poly_offset: usize, + offsets: Vec, + domain: &'a StarkDomain, +} + +impl<'a, E, I, const N: usize> SegmentGenerator<'a, E, I, N> +where + E: FieldElement, + I: IntoIterator>, +{ + fn new(polys: I, domain: &'a StarkDomain) -> Self { + assert!(N > 0, "batch size N must be greater than zero"); + let poly_size = domain.trace_length(); + let lde_blowup = domain.trace_to_lde_blowup(); + let offsets = get_evaluation_offsets::(poly_size, lde_blowup, domain.offset()); + Self { + poly_iter: polys.into_iter(), + polys: None, + poly_offset: 0, + offsets, + domain, + } + } + + /// Returns the matrix of polynomials used to generate segments. + fn into_polys(self) -> Option> { + self.polys + } + + /// Returns a segment generating iterator. + fn gen_segment_iter(&mut self) -> SegmentIterator<'a, '_, E, I, N> { + SegmentIterator(self) + } + + /// Generates the next segment if it exists otherwise returns None. + fn gen_next_segment(&mut self) -> Option> { + // initialize our col matrix + if self.polys.is_none() { + self.polys = Some(ColMatrix::new(vec![self.poly_iter.next()?])); + } + + let offset = self.poly_offset; + let polys = self.polys.as_mut().unwrap(); + while polys.num_base_cols() < offset + N { + if let Some(poly) = self.poly_iter.next() { + polys.merge_column(poly) + } else { + break; + } + } + + // terminate if there are no more segments to create + if polys.num_base_cols() <= offset { + return None; + } + + let domain_size = self.domain.lde_domain_size(); + let mut data = vec![[Felt::ZERO; N]; domain_size]; + if polys.num_base_cols() < offset + N { + // the segment will remain unfilled so we pad it with zeros + data.fill([Felt::ZERO; N]); + } + + let twiddles = self.domain.trace_twiddles(); + let segment = Segment::new_with_buffer(data, &*polys, offset, &self.offsets, twiddles); + self.poly_offset += N; + Some(segment) + } +} + +fn build_segment_queries< + E: FieldElement, + H: Hasher + ElementHasher, +>( + segment_lde: &RowMatrix, + segment_tree: &MerkleTree, + positions: &[usize], +) -> Queries { + // for each position, get the corresponding row from the trace segment LDE and put all these + // rows into a single vector + let trace_states = + positions.iter().map(|&pos| segment_lde.row(pos).to_vec()).collect::>(); + + // build Merkle authentication paths to the leaves specified by positions + let trace_proof = segment_tree + .prove_batch(positions) + .expect("failed to generate a Merkle proof for trace queries"); + + Queries::new(trace_proof, trace_states) +} + +struct SegmentIterator<'a, 'b, E, I, const N: usize>(&'b mut SegmentGenerator<'a, E, I, N>) +where + E: FieldElement, + I: IntoIterator>; + +impl<'a, 'b, E, I, const N: usize> Iterator for SegmentIterator<'a, 'b, E, I, N> +where + E: FieldElement, + I: IntoIterator>, +{ + type Item = Segment; + + fn next(&mut self) -> Option { + self.0.gen_next_segment() + } +} diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs new file mode 100644 index 000000000..f07b44350 --- /dev/null +++ b/prover/src/gpu/webgpu/tests.rs @@ -0,0 +1,227 @@ +use alloc::vec::Vec; + +use air::{ProvingOptions, StarkField}; +use gpu::webgpu::{WebGPUExecutionProver, DIGEST_SIZE, RATE}; +use processor::{ + crypto::{Hasher, RpoDigest, RpoRandomCoin, Rpx256, RpxDigest, RpxRandomCoin}, + math::fft, + StackInputs, StackOutputs, +}; +use winter_prover::{crypto::Digest, math::fields::CubeExtension, CompositionPolyTrace, TraceLde}; + +use crate::*; + +type CubeFelt = CubeExtension; + +fn build_trace_commitment_on_gpu_with_padding_matches_cpu< + R: RandomCoin + Send, + H: ElementHasher + Hasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, +>( + hash_fn: HashFn, +) { + let is_rpx = matches!(hash_fn, HashFn::Rpx256); + + let cpu_prover = create_test_prover::(is_rpx); + let gpu_prover = WebGPUExecutionProver::new(create_test_prover::(is_rpx), hash_fn); + let num_rows = 1 << 8; + let trace_info = get_trace_info(1, num_rows); + let trace = gen_random_trace(num_rows, RATE + 1); + let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); + + let (cpu_trace_lde, cpu_polys) = + cpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + let (gpu_trace_lde, gpu_polys) = + gpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + + assert_eq!( + cpu_trace_lde.get_main_trace_commitment(), + gpu_trace_lde.get_main_trace_commitment() + ); + assert_eq!( + cpu_polys.main_trace_polys().collect::>(), + gpu_polys.main_trace_polys().collect::>() + ); +} + +fn build_trace_commitment_on_gpu_without_padding_matches_cpu< + R: RandomCoin + Send, + H: ElementHasher + Hasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, +>( + hash_fn: HashFn, +) { + let is_rpx = matches!(hash_fn, HashFn::Rpx256); + + let cpu_prover = create_test_prover::(is_rpx); + let gpu_prover = WebGPUExecutionProver::new(create_test_prover::(is_rpx), hash_fn); + let num_rows = 1 << 8; + let trace_info = get_trace_info(1, num_rows); + let trace = gen_random_trace(num_rows, RATE); + let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); + + let (cpu_trace_lde, cpu_polys) = + cpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + let (gpu_trace_lde, gpu_polys) = + gpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + + assert_eq!( + cpu_trace_lde.get_main_trace_commitment(), + gpu_trace_lde.get_main_trace_commitment() + ); + assert_eq!( + cpu_polys.main_trace_polys().collect::>(), + gpu_polys.main_trace_polys().collect::>() + ); +} + +fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< + R: RandomCoin + Send, + H: ElementHasher + Hasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, +>( + hash_fn: HashFn, +) { + let is_rpx = matches!(hash_fn, HashFn::Rpx256); + + let cpu_prover = create_test_prover::(is_rpx); + let gpu_prover = WebGPUExecutionProver::new(create_test_prover::(is_rpx), hash_fn); + let num_rows = 1 << 8; + let ce_blowup_factor = 2; + let values = get_random_values::(num_rows * ce_blowup_factor); + let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); + + let (commitment_cpu, composition_poly_cpu) = cpu_prover.build_constraint_commitment( + CompositionPolyTrace::new(values.clone()), + 2, + &domain, + ); + let (commitment_gpu, composition_poly_gpu) = + gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 2, &domain); + + assert_eq!(commitment_cpu.root(), commitment_gpu.root()); + assert_ne!(0, composition_poly_cpu.data().num_base_cols() % RATE); + assert_eq!(composition_poly_cpu.into_columns(), composition_poly_gpu.into_columns()); +} + +fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< + R: RandomCoin + Send, + H: ElementHasher + Hasher, + D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, +>( + hash_fn: HashFn, +) { + let is_rpx = matches!(hash_fn, HashFn::Rpx256); + + let cpu_prover = create_test_prover::(is_rpx); + let gpu_prover = WebGPUExecutionProver::new(create_test_prover::(is_rpx), hash_fn); + let num_rows = 1 << 8; + let ce_blowup_factor = 8; + let values = get_random_values::(num_rows * ce_blowup_factor); + let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); + + let (commitment_cpu, composition_poly_cpu) = cpu_prover.build_constraint_commitment( + CompositionPolyTrace::new(values.clone()), + 8, + &domain, + ); + let (commitment_gpu, composition_poly_gpu) = + gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 8, &domain); + + assert_eq!(commitment_cpu.root(), commitment_gpu.root()); + assert_eq!(0, composition_poly_cpu.data().num_base_cols() % RATE); + assert_eq!(composition_poly_cpu.into_columns(), composition_poly_gpu.into_columns()); +} + +#[test] +fn rpo_build_trace_commitment_on_gpu_with_padding_matches_cpu() { + build_trace_commitment_on_gpu_with_padding_matches_cpu::( + HashFn::Rpo256, + ); +} + +#[test] +fn rpx_build_trace_commitment_on_gpu_with_padding_matches_cpu() { + build_trace_commitment_on_gpu_with_padding_matches_cpu::( + HashFn::Rpx256, + ); +} + +#[test] +fn rpo_build_trace_commitment_on_gpu_without_padding_matches_cpu() { + build_trace_commitment_on_gpu_without_padding_matches_cpu::( + HashFn::Rpo256, + ); +} + +#[test] +fn rpx_build_trace_commitment_on_gpu_without_padding_matches_cpu() { + build_trace_commitment_on_gpu_without_padding_matches_cpu::( + HashFn::Rpx256, + ); +} + +#[test] +fn rpo_build_constraint_commitment_on_gpu_with_padding_matches_cpu() { + build_constraint_commitment_on_gpu_with_padding_matches_cpu::( + HashFn::Rpo256, + ); +} + +#[test] +fn rpx_build_constraint_commitment_on_gpu_with_padding_matches_cpu() { + build_constraint_commitment_on_gpu_with_padding_matches_cpu::( + HashFn::Rpx256, + ); +} + +#[test] +fn rpo_build_constraint_commitment_on_gpu_without_padding_matches_cpu() { + build_constraint_commitment_on_gpu_without_padding_matches_cpu::< + RpoRandomCoin, + Rpo256, + RpoDigest, + >(HashFn::Rpo256); +} + +#[test] +fn rpx_build_constraint_commitment_on_gpu_without_padding_matches_cpu() { + build_constraint_commitment_on_gpu_without_padding_matches_cpu::< + RpxRandomCoin, + Rpx256, + RpxDigest, + >(HashFn::Rpx256); +} + +fn gen_random_trace(num_rows: usize, num_cols: usize) -> ColMatrix { + ColMatrix::new((0..num_cols as u64).map(|col| vec![Felt::new(col); num_rows]).collect()) +} + +fn get_random_values(num_rows: usize) -> Vec { + (0..num_rows).map(|i| E::from(i as u32)).collect() +} + +fn get_trace_info(num_cols: usize, num_rows: usize) -> TraceInfo { + TraceInfo::new(num_cols, num_rows) +} + +fn create_test_prover< + R: RandomCoin + Send, + H: ElementHasher, +>( + use_rpx: bool, +) -> ExecutionProver { + if use_rpx { + ExecutionProver::new( + ProvingOptions::with_128_bit_security_rpx(), + StackInputs::default(), + StackOutputs::default(), + ) + } else { + ExecutionProver::new( + ProvingOptions::with_128_bit_security(true), + StackInputs::default(), + StackOutputs::default(), + ) + } +} diff --git a/prover/src/lib.rs b/prover/src/lib.rs index b8bee0669..9ea31ec8d 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -1,5 +1,3 @@ -#![no_std] - #[cfg_attr(all(feature = "metal", target_arch = "aarch64", target_os = "macos"), macro_use)] extern crate alloc; @@ -8,8 +6,9 @@ extern crate std; use core::marker::PhantomData; +use maybe_async::maybe_async; + use air::{AuxRandElements, ProcessorAir, PublicInputs}; -#[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] use miden_gpu::HashFn; use processor::{ crypto::{ @@ -51,8 +50,9 @@ pub use winter_prover::Proof; /// /// # Errors /// Returns an error if program execution or STARK proof generation fails for any reason. +#[maybe_async] #[instrument("prove_program", skip_all)] -pub fn prove( +pub async fn prove( program: &Program, stack_inputs: StackInputs, host: H, @@ -86,13 +86,13 @@ where stack_inputs, stack_outputs.clone(), ) - .prove(trace), + .prove(trace).await, HashFunction::Blake3_256 => ExecutionProver::>::new( options, stack_inputs, stack_outputs.clone(), ) - .prove(trace), + .prove(trace).await, HashFunction::Rpo256 => { let prover = ExecutionProver::::new( options, @@ -101,7 +101,9 @@ where ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpo256); - prover.prove(trace) + #[cfg(feature = "webgpu")] + let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, HashFn::Rpo256); + prover.prove(trace).await }, HashFunction::Rpx256 => { let prover = ExecutionProver::::new( @@ -111,7 +113,9 @@ where ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpx256); - prover.prove(trace) + #[cfg(feature = "webgpu")] + let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, HashFn::Rpx256); + prover.prove(trace).await }, } .map_err(ExecutionError::ProverError)?; @@ -174,6 +178,7 @@ where } } +#[maybe_async] impl Prover for ExecutionProver where H: ElementHasher, @@ -207,7 +212,7 @@ where PublicInputs::new(program_info, self.stack_inputs.clone(), self.stack_outputs.clone()) } - fn new_trace_lde>( + async fn new_trace_lde>( &self, trace_info: &TraceInfo, main_trace: &ColMatrix, @@ -216,7 +221,7 @@ where DefaultTraceLde::new(trace_info, main_trace, domain) } - fn new_evaluator<'a, E: FieldElement>( + async fn new_evaluator<'a, E: FieldElement>( &self, air: &'a ProcessorAir, aux_rand_elements: Option>, @@ -225,7 +230,7 @@ where DefaultConstraintEvaluator::new(air, aux_rand_elements, composition_coefficients) } - fn build_aux_trace( + async fn build_aux_trace( &self, trace: &Self::Trace, aux_rand_elements: &AuxRandElements, From c6049548ae67d0c4547fedd6eaaee4cb3d05b4ee Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 15:23:53 +0800 Subject: [PATCH 03/16] fix build Signed-off-by: GopherJ --- Cargo.lock | 19 ++----- Cargo.toml | 1 + prover/Cargo.toml | 4 +- prover/src/gpu/webgpu/mod.rs | 101 ++++++++++++++++++++++----------- prover/src/gpu/webgpu/tests.rs | 24 ++++---- 5 files changed, 87 insertions(+), 62 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 19ade88c0..84ccc59b8 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1173,9 +1173,9 @@ checksum = "bbd2bcb4c963f2ddae06a2efc7e9f3591312473c50c6685e1f298068316e66fe" [[package]] name = "libc" -version = "0.2.157" +version = "0.2.158" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "374af5f94e54fa97cf75e945cce8a6b201e88a1a07e688b47dfd2a59c66dbd86" +checksum = "d8adc4bb1803a324070e64a98ae98f38934d91957a99cfb3a43dcbc01bc56439" [[package]] name = "libloading" @@ -1471,7 +1471,7 @@ dependencies = [ "miden-processor", "pollster", "tracing", - "winter-maybe-async 0.9.0 (registry+https://github.com/rust-lang/crates.io-index)", + "winter-maybe-async", "winter-prover", ] @@ -3352,17 +3352,6 @@ dependencies = [ "winter-utils", ] -[[package]] -name = "winter-maybe-async" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7ce0f4161cdde50de809b3869c1cb083a09e92e949428ea28f04c0d64045875c" -dependencies = [ - "proc-macro2", - "quote", - "syn 2.0.75", -] - [[package]] name = "winter-maybe-async" version = "0.9.0" @@ -3384,7 +3373,7 @@ dependencies = [ "winter-crypto", "winter-fri", "winter-math", - "winter-maybe-async 0.9.0 (git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e)", + "winter-maybe-async", "winter-utils", ] diff --git a/Cargo.toml b/Cargo.toml index d9ae4dd8b..d7e0c24ef 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -41,5 +41,6 @@ winter-verifier = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a0 winter-fri = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } winter-crypto = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } winter-math = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } +winter-maybe-async = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } winter-rand-utils = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } miden-crypto = { git = "https://github.com/GopherJ/miden-crypto", rev = "657b4922f3abe3577d1b0de1f6ae6ad52400cb11" } diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 73b721ec7..4f0c7bf4c 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -15,10 +15,10 @@ edition.workspace = true [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["webgpu"] +default = ["std"] metal = ["dep:miden-gpu", "dep:elsa", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "dep:elsa", "miden-gpu/webgpu", "async"] -async = ["maybe-async/async", "winter-prover/async", "winter-prover/async-trait", "dep:async-trait"] +async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std", "tracing/std"] [dependencies] diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index 1189fd3e0..a9c603592 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -2,26 +2,24 @@ //! For now, the logic is limited to GPU accelerating trace and constraint commitments, //! using the RPO 256 or RPX 256 hash functions. -#[cfg(feature = "std")] -use std::{boxed::Box, marker::PhantomData, time::Instant, vec::Vec}; - #[cfg(not(feature = "std"))] use core::marker::PhantomData; +#[cfg(feature = "std")] +use std::{boxed::Box, marker::PhantomData, time::Instant, vec::Vec}; #[cfg(not(feature = "std"))] extern crate alloc; +#[cfg(not(feature = "std"))] +use alloc::boxed::Box; #[cfg(not(feature = "std"))] use alloc::vec; #[cfg(not(feature = "std"))] use alloc::vec::Vec; -#[cfg(not(feature = "std"))] -use alloc::boxed::Box; +use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use elsa::FrozenVec; use maybe_async::maybe_async; - -use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use miden_gpu::{ webgpu::{ build_merkle_tree, get_dispatch_linear, get_wgpu_helper, init_wgpu_helper, RowHasher, @@ -33,7 +31,7 @@ use processor::{ crypto::{ElementHasher, Hasher}, ONE, }; -use tracing::{event, Level}; +use tracing::{event, info_span, Level}; use winter_prover::{ crypto::{Digest, MerkleTree}, matrix::{get_evaluation_offsets, ColMatrix, RowMatrix, Segment}, @@ -342,10 +340,9 @@ impl< domain: &StarkDomain, webgpu_hash_fn: HashFn, ) -> (Self, TracePolyTable) { - init_wgpu_helper().await; // extend the main execution trace and build a Merkle tree from the extended trace let (main_segment_lde, main_segment_tree, main_segment_polys) = - build_trace_commitment(main_trace, domain, webgpu_hash_fn).await; + build_trace_commitment_sync::(main_trace, domain); let trace_poly_table = TracePolyTable::new(main_segment_polys); let trace_lde = WebGPUTraceLde { @@ -420,22 +417,23 @@ impl< aux_trace: &ColMatrix, domain: &StarkDomain, ) -> (ColMatrix, D) { - // extend the auxiliary trace segment and build a Merkle tree from the extended trace - let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = - build_trace_commitment::(aux_trace, domain, self.webgpu_hash_fn).await; - - assert_eq!( - self.main_segment_lde.num_rows(), - aux_segment_lde.num_rows(), - "the number of rows in the auxiliary segment must be the same as in the main segment" - ); - - // save the lde and commitment - self.aux_segment_lde = Some(aux_segment_lde); - let root_hash = *aux_segment_tree.root(); - self.aux_segment_tree = Some(aux_segment_tree); - - (aux_segment_polys, root_hash) + todo!() + // // extend the auxiliary trace segment and build a Merkle tree from the extended trace + // let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = + // build_trace_commitment::(aux_trace, domain, self.webgpu_hash_fn).await; + // + // assert_eq!( + // self.main_segment_lde.num_rows(), + // aux_segment_lde.num_rows(), + // "the number of rows in the auxiliary segment must be the same as in the main segment" + // ); + // + // // save the lde and commitment + // self.aux_segment_lde = Some(aux_segment_lde); + // let root_hash = *aux_segment_tree.root(); + // self.aux_segment_tree = Some(aux_segment_tree); + // + // (aux_segment_polys, root_hash) } /// Reads current and next rows from the main trace segment into the specified frame. @@ -550,6 +548,45 @@ impl< /// ────┼────────┼────────┼────────┼────────┼────────┼──── /// t=n t=n+1 t=n+2 t=n+3 t=n+4 t=n+5 /// ``` +const DEFAULT_SEGMENT_WIDTH: usize = 8; + +fn build_trace_commitment_sync( + trace: &ColMatrix, + domain: &StarkDomain, +) -> (RowMatrix, MerkleTree, ColMatrix) +where + E: FieldElement, + F: FieldElement, + H: ElementHasher, +{ + // extend the execution trace + let (trace_lde, trace_polys) = { + let span = info_span!( + "extend_execution_trace", + num_cols = trace.num_cols(), + blowup = domain.trace_to_lde_blowup() + ) + .entered(); + let trace_polys = trace.interpolate_columns(); + let trace_lde = + RowMatrix::evaluate_polys_over::(&trace_polys, domain); + drop(span); + + (trace_lde, trace_polys) + }; + assert_eq!(trace_lde.num_cols(), trace.num_cols()); + assert_eq!(trace_polys.num_rows(), trace.num_rows()); + assert_eq!(trace_lde.num_rows(), domain.lde_domain_size()); + + // build trace commitment + let tree_depth = trace_lde.num_rows().ilog2() as usize; + let trace_tree = info_span!("compute_execution_trace_commitment", tree_depth) + .in_scope(|| trace_lde.commit_to_rows()); + assert_eq!(trace_tree.depth(), tree_depth); + + (trace_lde, trace_tree, trace_polys) +} + async fn build_trace_commitment< E: FieldElement, H: Hasher + ElementHasher, @@ -575,12 +612,8 @@ async fn build_trace_commitment< let num_base_columns = trace.num_base_cols(); let rpo_requires_padding = num_base_columns % RATE != 0; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); - let mut row_hasher = RowHasher::new( - &get_wgpu_helper().unwrap(), - num_base_columns, - rpo_requires_padding, - hash_fn, - ); + let mut row_hasher = + RowHasher::new(&get_wgpu_helper().unwrap(), lde_domain_size, rpo_requires_padding, hash_fn); let mut rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); @@ -612,7 +645,9 @@ async fn build_trace_commitment< row_hasher.update(&get_wgpu_helper().unwrap(), segment); } let row_hashes = row_hasher.finish(&get_wgpu_helper().unwrap()).await.unwrap(); - let tree_nodes = build_merkle_tree(&get_wgpu_helper().unwrap(), &row_hashes, hash_fn).await.unwrap(); + let tree_nodes = build_merkle_tree(&get_wgpu_helper().unwrap(), &row_hashes, hash_fn) + .await + .unwrap(); // aggregate segments at the same time as the GPU generates the merkle tree nodes let lde_segments = lde_segments.into_vec().into_iter().map(|p| *p).collect(); let trace_lde = RowMatrix::from_segments(lde_segments, num_base_columns); diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs index f07b44350..10c89318c 100644 --- a/prover/src/gpu/webgpu/tests.rs +++ b/prover/src/gpu/webgpu/tests.rs @@ -13,7 +13,7 @@ use crate::*; type CubeFelt = CubeExtension; -fn build_trace_commitment_on_gpu_with_padding_matches_cpu< +async fn build_trace_commitment_on_gpu_with_padding_matches_cpu< R: RandomCoin + Send, H: ElementHasher + Hasher, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, @@ -30,9 +30,9 @@ fn build_trace_commitment_on_gpu_with_padding_matches_cpu< let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); let (cpu_trace_lde, cpu_polys) = - cpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + cpu_prover.new_trace_lde::(&trace_info, &trace, &domain).await; let (gpu_trace_lde, gpu_polys) = - gpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + gpu_prover.new_trace_lde::(&trace_info, &trace, &domain).await; assert_eq!( cpu_trace_lde.get_main_trace_commitment(), @@ -44,7 +44,7 @@ fn build_trace_commitment_on_gpu_with_padding_matches_cpu< ); } -fn build_trace_commitment_on_gpu_without_padding_matches_cpu< +async fn build_trace_commitment_on_gpu_without_padding_matches_cpu< R: RandomCoin + Send, H: ElementHasher + Hasher, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, @@ -61,9 +61,9 @@ fn build_trace_commitment_on_gpu_without_padding_matches_cpu< let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); let (cpu_trace_lde, cpu_polys) = - cpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + cpu_prover.new_trace_lde::(&trace_info, &trace, &domain).await; let (gpu_trace_lde, gpu_polys) = - gpu_prover.new_trace_lde::(&trace_info, &trace, &domain); + gpu_prover.new_trace_lde::(&trace_info, &trace, &domain).await; assert_eq!( cpu_trace_lde.get_main_trace_commitment(), @@ -75,7 +75,7 @@ fn build_trace_commitment_on_gpu_without_padding_matches_cpu< ); } -fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< +async fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< R: RandomCoin + Send, H: ElementHasher + Hasher, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, @@ -95,16 +95,16 @@ fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< CompositionPolyTrace::new(values.clone()), 2, &domain, - ); + ).await; let (commitment_gpu, composition_poly_gpu) = - gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 2, &domain); + gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 2, &domain).await; assert_eq!(commitment_cpu.root(), commitment_gpu.root()); assert_ne!(0, composition_poly_cpu.data().num_base_cols() % RATE); assert_eq!(composition_poly_cpu.into_columns(), composition_poly_gpu.into_columns()); } -fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< +async fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< R: RandomCoin + Send, H: ElementHasher + Hasher, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, @@ -124,9 +124,9 @@ fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< CompositionPolyTrace::new(values.clone()), 8, &domain, - ); + ).await; let (commitment_gpu, composition_poly_gpu) = - gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 8, &domain); + gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 8, &domain).await; assert_eq!(commitment_cpu.root(), commitment_gpu.root()); assert_eq!(0, composition_poly_cpu.data().num_base_cols() % RATE); From 22ee95ff7b470b9006f978ae338802d334415c1a Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 16:39:07 +0800 Subject: [PATCH 04/16] migrate test to wasm & fix send-sync issue Signed-off-by: GopherJ --- Cargo.lock | 115 +++++++++++++++++++++++++------ prover/Cargo.toml | 22 ++++-- prover/src/gpu/webgpu/tests.rs | 121 ++++++++++++++++++++------------- prover/src/lib.rs | 30 ++++---- 4 files changed, 202 insertions(+), 86 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 84ccc59b8..7c9503073 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -23,7 +23,7 @@ version = "0.8.11" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "e89da841a80418a9b391ebaea17f5c112ffaaa96f621d2c285b5174da76b9011" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "once_cell", "version_check", "zerocopy", @@ -188,7 +188,7 @@ checksum = "5cc23269a4f8976d0a4d2e7109211a419fe30e8d88d677cd60b6bc79c5732e0a" dependencies = [ "addr2line", "cc", - "cfg-if", + "cfg-if 1.0.0", "libc", "miniz_oxide", "object", @@ -255,7 +255,7 @@ dependencies = [ "arrayref", "arrayvec", "cc", - "cfg-if", + "cfg-if 1.0.0", "constant_time_eq", ] @@ -320,6 +320,12 @@ dependencies = [ "shlex", ] +[[package]] +name = "cfg-if" +version = "0.1.10" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4785bdd1c96b2a846b2bd7cc02e86b6b3dbf14e7e53446c4f54c92a361040822" + [[package]] name = "cfg-if" version = "1.0.0" @@ -455,6 +461,16 @@ dependencies = [ "syn 1.0.109", ] +[[package]] +name = "console_error_panic_hook" +version = "0.1.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a06aeb73f470f66dcdbf7223caeebb85984942f22f1adb2a088cf9668146bbbc" +dependencies = [ + "cfg-if 1.0.0", + "wasm-bindgen", +] + [[package]] name = "constant_time_eq" version = "0.3.0" @@ -613,7 +629,7 @@ version = "2.0.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b98cf8ebf19c3d1b223e151f99a4f9f0690dca41414773390fc824184ac833e1" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "dirs-sys-next", ] @@ -840,7 +856,7 @@ version = "0.8.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "979f00864edc7516466d6b3157706e06c032f22715700ddd878228a91d02bc56" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "libc", "log", "rustversion", @@ -863,7 +879,7 @@ version = "0.2.15" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c4567c8db10ae91089c99af84c68c38da3ec2f087c3f82960bcdbf3656b6f4d7" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "js-sys", "libc", "wasi", @@ -972,7 +988,7 @@ version = "2.4.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6dd08c532ae367adf81c312a4580bc67f1d0fe8bc9c460520283f4c0ff277888" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "crunchy", ] @@ -1183,7 +1199,7 @@ version = "0.8.5" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "4979f22fdb869068da03c9f7528f8297c6fd2606bc3a4affe42e6a823fdb8da4" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "windows-targets 0.52.6", ] @@ -1250,7 +1266,7 @@ version = "0.7.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "419e0dc8046cb947daa77eb95ae174acfbddb7673b4151f56d1eed8e93fbfaca" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "generator", "scoped-tls", "tracing", @@ -1281,6 +1297,12 @@ version = "2.7.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "78ca9ab1a0babb1e7d5695e3530886289c18cf2f87ec19a575a0abdce112e3a3" +[[package]] +name = "memory_units" +version = "0.4.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8452105ba047068f40ff7093dd1d9da90898e63dd61736462e9cdda6a90ad3c3" + [[package]] name = "metal" version = "0.27.0" @@ -1392,7 +1414,7 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.2.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b77b4709d69d42ef663ed5ea41f48f8b165f59f2#b77b4709d69d42ef663ed5ea41f48f8b165f59f2" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=9a55000ff8cc113c1f4acb42115c542c1bb8533c#9a55000ff8cc113c1f4acb42115c542c1bb8533c" dependencies = [ "bytemuck", "elsa", @@ -1411,7 +1433,7 @@ checksum = "c532250422d933f15b148fb81e4522a5d649c178ab420d0d596c86228da35570" dependencies = [ "backtrace", "backtrace-ext", - "cfg-if", + "cfg-if 1.0.0", "futures", "indenter", "lazy_static", @@ -1465,12 +1487,17 @@ name = "miden-prover" version = "0.11.0" dependencies = [ "async-trait", + "console_error_panic_hook", "elsa", "miden-air", "miden-gpu", "miden-processor", "pollster", "tracing", + "wasm-bindgen", + "wasm-bindgen-futures", + "wasm-bindgen-test", + "wee_alloc", "winter-maybe-async", "winter-prover", ] @@ -1572,6 +1599,16 @@ dependencies = [ "winter-fri", ] +[[package]] +name = "minicov" +version = "0.3.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5c71e683cd655513b99affab7d317deb690528255a0d5f717f1024093c12b169" +dependencies = [ + "cc", + "walkdir", +] + [[package]] name = "miniz_oxide" version = "0.7.4" @@ -1633,7 +1670,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "2eb04e9c688eff1c89d72b407f168cf79bb9e867a9d3323ed6c01519eb9cc053" dependencies = [ "bitflags 2.6.0", - "cfg-if", + "cfg-if 1.0.0", "libc", ] @@ -1806,7 +1843,7 @@ version = "0.9.10" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "1e401f977ab385c9e4e3ab30627d6f26d00e2c73eef317493c4ec6d468726cf8" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "libc", "redox_syscall", "smallvec", @@ -2210,7 +2247,7 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "02a2d683a4ac90aeef5b1013933f6d977bd37d51ff3f4dad829d4931a7e6be86" dependencies = [ "bitflags 2.6.0", - "cfg-if", + "cfg-if 1.0.0", "clipboard-win", "libc", "log", @@ -2317,7 +2354,7 @@ version = "0.10.8" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "793db75ad2bcafc3ffa7c68b215fee268f537982cd901d132f89c6343f3a3dc8" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "cpufeatures", "digest", ] @@ -2481,7 +2518,7 @@ version = "3.12.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "04cbcdd0c794ebb0d4cf35e88edd2f7d2c4c3e9a5a6dab322839b321c6a87a64" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "fastrand", "once_cell", "rustix", @@ -2539,7 +2576,7 @@ version = "3.3.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "adcb7fd841cd518e279be3d5a3eb0636409487998a4aff22f3de87b81e88384f" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "proc-macro2", "quote", "syn 2.0.75", @@ -2594,7 +2631,7 @@ version = "1.1.8" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8b9ef9bad013ada3808854ceac7b46812a6465ba368859a37e2100283d2d719c" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "once_cell", ] @@ -2857,7 +2894,7 @@ version = "0.2.93" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a82edfc16a6c469f5f44dc7b571814045d60404b55a0ee849f9bcfa2e63dd9b5" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "once_cell", "wasm-bindgen-macro", ] @@ -2883,7 +2920,7 @@ version = "0.4.43" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "61e9300f63a621e96ed275155c108eb6f843b6a26d053f122ab69724559dc8ed" dependencies = [ - "cfg-if", + "cfg-if 1.0.0", "js-sys", "wasm-bindgen", "web-sys", @@ -2918,6 +2955,32 @@ version = "0.2.93" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "c62a0a307cb4a311d3a07867860911ca130c3494e8c2719593806c08bc5d0484" +[[package]] +name = "wasm-bindgen-test" +version = "0.3.43" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68497a05fb21143a08a7d24fc81763384a3072ee43c44e86aad1744d6adef9d9" +dependencies = [ + "console_error_panic_hook", + "js-sys", + "minicov", + "scoped-tls", + "wasm-bindgen", + "wasm-bindgen-futures", + "wasm-bindgen-test-macro", +] + +[[package]] +name = "wasm-bindgen-test-macro" +version = "0.3.43" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4b8220be1fa9e4c889b30fd207d4906657e7e90b12e0e6b0c8b8d8709f5de021" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.75", +] + [[package]] name = "web-sys" version = "0.3.70" @@ -2928,6 +2991,18 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "wee_alloc" +version = "0.4.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dbb3b5a6b2bb17cb6ad44a2e68a43e8d2722c997da10e928665c72ec6c0a0b8e" +dependencies = [ + "cfg-if 0.1.10", + "libc", + "memory_units", + "winapi", +] + [[package]] name = "wgpu" version = "22.1.0" diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 4f0c7bf4c..dc2adf1b7 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -13,11 +13,14 @@ repository.workspace = true rust-version.workspace = true edition.workspace = true +[lib] +crate-type = ["cdylib", "rlib"] + [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["std"] -metal = ["dep:miden-gpu", "dep:elsa", "dep:pollster", "concurrent", "std"] -webgpu = ["dep:miden-gpu", "dep:elsa", "miden-gpu/webgpu", "async"] +default = ["webgpu", "async"] +metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] +webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std", "tracing/std"] @@ -28,11 +31,18 @@ tracing = { version = "0.1", default-features = false, features = ["attributes"] winter-prover = { package = "winter-prover", version = "0.9", default-features = false } maybe-async = { package = "winter-maybe-async", version = "0.9", default-features = false } async-trait = { version = "0.1", optional = true } +elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -elsa = { version = "1.9", optional = true } -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b77b4709d69d42ef663ed5ea41f48f8b165f59f2", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "9a55000ff8cc113c1f4acb42115c542c1bb8533c", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b77b4709d69d42ef663ed5ea41f48f8b165f59f2", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "9a55000ff8cc113c1f4acb42115c542c1bb8533c", default-features = false, optional = true } + +[target.'cfg(target_family = "wasm")'.dev-dependencies] +wasm-bindgen-test = "0.3" +wasm-bindgen-futures = "0.4" +wasm-bindgen = "0.2" +wee_alloc = "0.4" +console_error_panic_hook = "0.1" diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs index 10c89318c..df1f4248c 100644 --- a/prover/src/gpu/webgpu/tests.rs +++ b/prover/src/gpu/webgpu/tests.rs @@ -1,3 +1,5 @@ +#![cfg(target_family = "wasm")] + use alloc::vec::Vec; use air::{ProvingOptions, StarkField}; @@ -11,6 +13,11 @@ use winter_prover::{crypto::Digest, math::fields::CubeExtension, CompositionPoly use crate::*; +use wasm_bindgen::prelude::*; +use wasm_bindgen_test::*; + +wasm_bindgen_test_configure!(run_in_browser); + type CubeFelt = CubeExtension; async fn build_trace_commitment_on_gpu_with_padding_matches_cpu< @@ -91,13 +98,12 @@ async fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< let values = get_random_values::(num_rows * ce_blowup_factor); let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); - let (commitment_cpu, composition_poly_cpu) = cpu_prover.build_constraint_commitment( - CompositionPolyTrace::new(values.clone()), - 2, - &domain, - ).await; - let (commitment_gpu, composition_poly_gpu) = - gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 2, &domain).await; + let (commitment_cpu, composition_poly_cpu) = cpu_prover + .build_constraint_commitment(CompositionPolyTrace::new(values.clone()), 2, &domain) + .await; + let (commitment_gpu, composition_poly_gpu) = gpu_prover + .build_constraint_commitment(CompositionPolyTrace::new(values), 2, &domain) + .await; assert_eq!(commitment_cpu.root(), commitment_gpu.root()); assert_ne!(0, composition_poly_cpu.data().num_base_cols() % RATE); @@ -120,77 +126,96 @@ async fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< let values = get_random_values::(num_rows * ce_blowup_factor); let domain = StarkDomain::from_twiddles(fft::get_twiddles(num_rows), 8, Felt::GENERATOR); - let (commitment_cpu, composition_poly_cpu) = cpu_prover.build_constraint_commitment( - CompositionPolyTrace::new(values.clone()), - 8, - &domain, - ).await; - let (commitment_gpu, composition_poly_gpu) = - gpu_prover.build_constraint_commitment(CompositionPolyTrace::new(values), 8, &domain).await; + let (commitment_cpu, composition_poly_cpu) = cpu_prover + .build_constraint_commitment(CompositionPolyTrace::new(values.clone()), 8, &domain) + .await; + let (commitment_gpu, composition_poly_gpu) = gpu_prover + .build_constraint_commitment(CompositionPolyTrace::new(values), 8, &domain) + .await; assert_eq!(commitment_cpu.root(), commitment_gpu.root()); assert_eq!(0, composition_poly_cpu.data().num_base_cols() % RATE); assert_eq!(composition_poly_cpu.into_columns(), composition_poly_gpu.into_columns()); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpo_build_trace_commitment_on_gpu_with_padding_matches_cpu() { - build_trace_commitment_on_gpu_with_padding_matches_cpu::( - HashFn::Rpo256, - ); + wasm_bindgen_futures::spawn_local(build_trace_commitment_on_gpu_with_padding_matches_cpu::< + RpoRandomCoin, + Rpo256, + RpoDigest, + >(HashFn::Rpo256)); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpx_build_trace_commitment_on_gpu_with_padding_matches_cpu() { - build_trace_commitment_on_gpu_with_padding_matches_cpu::( - HashFn::Rpx256, - ); + wasm_bindgen_futures::spawn_local(build_trace_commitment_on_gpu_with_padding_matches_cpu::< + RpxRandomCoin, + Rpx256, + RpxDigest, + >(HashFn::Rpx256)); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpo_build_trace_commitment_on_gpu_without_padding_matches_cpu() { - build_trace_commitment_on_gpu_without_padding_matches_cpu::( - HashFn::Rpo256, - ); + wasm_bindgen_futures::spawn_local(build_trace_commitment_on_gpu_without_padding_matches_cpu::< + RpoRandomCoin, + Rpo256, + RpoDigest, + >(HashFn::Rpo256)); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpx_build_trace_commitment_on_gpu_without_padding_matches_cpu() { - build_trace_commitment_on_gpu_without_padding_matches_cpu::( - HashFn::Rpx256, - ); + wasm_bindgen_futures::spawn_local(build_trace_commitment_on_gpu_without_padding_matches_cpu::< + RpxRandomCoin, + Rpx256, + RpxDigest, + >(HashFn::Rpx256)); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpo_build_constraint_commitment_on_gpu_with_padding_matches_cpu() { - build_constraint_commitment_on_gpu_with_padding_matches_cpu::( - HashFn::Rpo256, + wasm_bindgen_futures::spawn_local( + build_constraint_commitment_on_gpu_with_padding_matches_cpu::< + RpoRandomCoin, + Rpo256, + RpoDigest, + >(HashFn::Rpo256), ); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpx_build_constraint_commitment_on_gpu_with_padding_matches_cpu() { - build_constraint_commitment_on_gpu_with_padding_matches_cpu::( - HashFn::Rpx256, + wasm_bindgen_futures::spawn_local( + build_constraint_commitment_on_gpu_with_padding_matches_cpu::< + RpxRandomCoin, + Rpx256, + RpxDigest, + >(HashFn::Rpx256), ); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpo_build_constraint_commitment_on_gpu_without_padding_matches_cpu() { - build_constraint_commitment_on_gpu_without_padding_matches_cpu::< - RpoRandomCoin, - Rpo256, - RpoDigest, - >(HashFn::Rpo256); + wasm_bindgen_futures::spawn_local( + build_constraint_commitment_on_gpu_without_padding_matches_cpu::< + RpoRandomCoin, + Rpo256, + RpoDigest, + >(HashFn::Rpo256), + ); } -#[test] +#[wasm_bindgen_test::wasm_bindgen_test] fn rpx_build_constraint_commitment_on_gpu_without_padding_matches_cpu() { - build_constraint_commitment_on_gpu_without_padding_matches_cpu::< - RpxRandomCoin, - Rpx256, - RpxDigest, - >(HashFn::Rpx256); + wasm_bindgen_futures::spawn_local( + build_constraint_commitment_on_gpu_without_padding_matches_cpu::< + RpxRandomCoin, + Rpx256, + RpxDigest, + >(HashFn::Rpx256), + ); } fn gen_random_trace(num_rows: usize, num_cols: usize) -> ColMatrix { diff --git a/prover/src/lib.rs b/prover/src/lib.rs index 9ea31ec8d..ee0077dd7 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -81,18 +81,24 @@ where // generate STARK proof let proof = match hash_fn { - HashFunction::Blake3_192 => ExecutionProver::>::new( - options, - stack_inputs, - stack_outputs.clone(), - ) - .prove(trace).await, - HashFunction::Blake3_256 => ExecutionProver::>::new( - options, - stack_inputs, - stack_outputs.clone(), - ) - .prove(trace).await, + HashFunction::Blake3_192 => { + ExecutionProver::>::new( + options, + stack_inputs, + stack_outputs.clone(), + ) + .prove(trace) + .await + }, + HashFunction::Blake3_256 => { + ExecutionProver::>::new( + options, + stack_inputs, + stack_outputs.clone(), + ) + .prove(trace) + .await + }, HashFunction::Rpo256 => { let prover = ExecutionProver::::new( options, From 6675833ce7726e3df2ab48e7ea791186014ac643 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 16:46:33 +0800 Subject: [PATCH 05/16] print error stack in browser Signed-off-by: GopherJ --- prover/src/gpu/webgpu/mod.rs | 32 ++++++++++++++------------------ prover/src/gpu/webgpu/tests.rs | 4 ++++ 2 files changed, 18 insertions(+), 18 deletions(-) diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index a9c603592..3e27647e0 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -21,17 +21,14 @@ use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use elsa::FrozenVec; use maybe_async::maybe_async; use miden_gpu::{ - webgpu::{ - build_merkle_tree, get_dispatch_linear, get_wgpu_helper, init_wgpu_helper, RowHasher, - WebGpuHelper, - }, + webgpu::{build_merkle_tree, get_wgpu_helper, RowHasher}, HashFn, }; use processor::{ crypto::{ElementHasher, Hasher}, ONE, }; -use tracing::{event, info_span, Level}; +use tracing::info_span; use winter_prover::{ crypto::{Digest, MerkleTree}, matrix::{get_evaluation_offsets, ColMatrix, RowMatrix, Segment}, @@ -247,7 +244,7 @@ where let rpo_requires_padding = num_base_columns % RATE != 0; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = - RowHasher::new(&helper, lde_domain_size, rpo_requires_padding, self.webgpu_hash_fn); + RowHasher::new(helper, lde_domain_size, rpo_requires_padding, self.webgpu_hash_fn); let mut rpo_padded_segment: Vec<[Felt; RATE]>; for (segment_idx, segment) in segments.iter().enumerate() { // check if the segment requires padding @@ -260,7 +257,7 @@ where rpo_padded_segment = segment .iter() .map(|x| { - let mut s = x.clone(); + let mut s = *x; s[rpo_pad_column] = ONE; s }) @@ -270,15 +267,14 @@ where let rpo_pad_column = num_base_columns % RATE; rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); } - row_hasher.update(&helper, &rpo_padded_segment); + row_hasher.update(helper, &rpo_padded_segment); assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); break; } - row_hasher.update(&helper, segment); + row_hasher.update(helper, segment); } - let row_hashes = row_hasher.finish(&helper).await.unwrap(); - let tree_nodes = - build_merkle_tree(&helper, &row_hashes, self.webgpu_hash_fn).await.unwrap(); + let row_hashes = row_hasher.finish(helper).await.unwrap(); + let tree_nodes = build_merkle_tree(helper, &row_hashes, self.webgpu_hash_fn).await.unwrap(); // aggregate segments at the same time as the GPU generates the merkle tree nodes let composed_evaluations = RowMatrix::::from_segments(segments, num_base_columns); let nodes = tree_nodes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); @@ -613,7 +609,7 @@ async fn build_trace_commitment< let rpo_requires_padding = num_base_columns % RATE != 0; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = - RowHasher::new(&get_wgpu_helper().unwrap(), lde_domain_size, rpo_requires_padding, hash_fn); + RowHasher::new(get_wgpu_helper().unwrap(), lde_domain_size, rpo_requires_padding, hash_fn); let mut rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); @@ -628,7 +624,7 @@ async fn build_trace_commitment< rpo_padded_segment = segment .iter() .map(|x| { - let mut s = x.clone(); + let mut s = *x; s[rpo_pad_column] = ONE; s }) @@ -638,14 +634,14 @@ async fn build_trace_commitment< let rpo_pad_column = num_base_columns % RATE; rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); } - row_hasher.update(&get_wgpu_helper().unwrap(), &rpo_padded_segment); + row_hasher.update(get_wgpu_helper().unwrap(), &rpo_padded_segment); assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); break; } - row_hasher.update(&get_wgpu_helper().unwrap(), segment); + row_hasher.update(get_wgpu_helper().unwrap(), segment); } - let row_hashes = row_hasher.finish(&get_wgpu_helper().unwrap()).await.unwrap(); - let tree_nodes = build_merkle_tree(&get_wgpu_helper().unwrap(), &row_hashes, hash_fn) + let row_hashes = row_hasher.finish(get_wgpu_helper().unwrap()).await.unwrap(); + let tree_nodes = build_merkle_tree(get_wgpu_helper().unwrap(), &row_hashes, hash_fn) .await .unwrap(); // aggregate segments at the same time as the GPU generates the merkle tree nodes diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs index df1f4248c..95da5ecd2 100644 --- a/prover/src/gpu/webgpu/tests.rs +++ b/prover/src/gpu/webgpu/tests.rs @@ -27,6 +27,7 @@ async fn build_trace_commitment_on_gpu_with_padding_matches_cpu< >( hash_fn: HashFn, ) { + console_error_panic_hook::set_once(); let is_rpx = matches!(hash_fn, HashFn::Rpx256); let cpu_prover = create_test_prover::(is_rpx); @@ -58,6 +59,7 @@ async fn build_trace_commitment_on_gpu_without_padding_matches_cpu< >( hash_fn: HashFn, ) { + console_error_panic_hook::set_once(); let is_rpx = matches!(hash_fn, HashFn::Rpx256); let cpu_prover = create_test_prover::(is_rpx); @@ -89,6 +91,7 @@ async fn build_constraint_commitment_on_gpu_with_padding_matches_cpu< >( hash_fn: HashFn, ) { + console_error_panic_hook::set_once(); let is_rpx = matches!(hash_fn, HashFn::Rpx256); let cpu_prover = create_test_prover::(is_rpx); @@ -117,6 +120,7 @@ async fn build_constraint_commitment_on_gpu_without_padding_matches_cpu< >( hash_fn: HashFn, ) { + console_error_panic_hook::set_once(); let is_rpx = matches!(hash_fn, HashFn::Rpx256); let cpu_prover = create_test_prover::(is_rpx); From a807f8f9e262276970166678228213aafa7835dc Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 16:48:29 +0800 Subject: [PATCH 06/16] fix build for std Signed-off-by: GopherJ --- prover/Cargo.toml | 2 +- prover/src/gpu/webgpu/mod.rs | 1 + 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/prover/Cargo.toml b/prover/Cargo.toml index dc2adf1b7..96e859ab4 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["webgpu", "async"] +default = ["webgpu", "async", "std"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index 3e27647e0..ae44f45c1 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -37,6 +37,7 @@ use winter_prover::{ DefaultConstraintEvaluator, EvaluationFrame, Prover, StarkDomain, TraceInfo, TraceLde, TracePolyTable, }; +use tracing::{event, Level}; use crate::{ crypto::{RandomCoin, Rpo256}, From cae2b90657430db061c617aab5af629c369ceb41 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 19 Aug 2024 22:25:20 +0800 Subject: [PATCH 07/16] update Signed-off-by: GopherJ --- Cargo.lock | 2 +- prover/Cargo.toml | 6 ++-- prover/src/gpu/webgpu/mod.rs | 64 ++++++++++++++++-------------------- prover/src/lib.rs | 1 + 4 files changed, 34 insertions(+), 39 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 7c9503073..292c0b4d9 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1414,7 +1414,7 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.2.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=9a55000ff8cc113c1f4acb42115c542c1bb8533c#9a55000ff8cc113c1f4acb42115c542c1bb8533c" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70#b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70" dependencies = [ "bytemuck", "elsa", diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 96e859ab4..dbb944ea9 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["webgpu", "async", "std"] +default = ["std"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] @@ -34,11 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "9a55000ff8cc113c1f4acb42115c542c1bb8533c", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "9a55000ff8cc113c1f4acb42115c542c1bb8533c", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index ae44f45c1..f1ef5d4b2 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -13,15 +13,13 @@ extern crate alloc; #[cfg(not(feature = "std"))] use alloc::boxed::Box; #[cfg(not(feature = "std"))] -use alloc::vec; -#[cfg(not(feature = "std"))] -use alloc::vec::Vec; +use alloc::{vec, vec::Vec}; use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use elsa::FrozenVec; use maybe_async::maybe_async; use miden_gpu::{ - webgpu::{build_merkle_tree, get_wgpu_helper, RowHasher}, + webgpu::{build_merkle_tree, get_or_init_wgpu_helper, init_wgpu_helper, RowHasher}, HashFn, }; use processor::{ @@ -29,6 +27,7 @@ use processor::{ ONE, }; use tracing::info_span; +use tracing::{event, Level}; use winter_prover::{ crypto::{Digest, MerkleTree}, matrix::{get_evaluation_offsets, ColMatrix, RowMatrix, Segment}, @@ -37,7 +36,6 @@ use winter_prover::{ DefaultConstraintEvaluator, EvaluationFrame, Prover, StarkDomain, TraceInfo, TraceLde, TracePolyTable, }; -use tracing::{event, Level}; use crate::{ crypto::{RandomCoin, Rpo256}, @@ -49,6 +47,13 @@ use crate::{ #[cfg(test)] mod tests; +#[allow(clippy::uninit_vec)] +pub unsafe fn uninit_vector_real(length: usize) -> Vec { + let mut vector = Vec::with_capacity(length); + vector.set_len(length); + vector +} + // CONSTANTS // ================================================================================================ @@ -102,15 +107,15 @@ where assert!(poly_offset < polys.num_base_cols()); // allocate memory for the segment - // let data = if polys.num_base_cols() - poly_offset >= N { - // // if we will fill the entire segment, we allocate uninitialized memory - // unsafe { page_aligned_uninit_vector(domain_size) } - // } else { - // but if some columns in the segment will remain unfilled, we allocate memory - // initialized to zeros to make sure we don't end up with memory with - // undefined values - let data = alloc::vec![[E::BaseField::ZERO; N]; domain_size]; - // }; + let data = if polys.num_base_cols() - poly_offset >= N { + // if we will fill the entire segment, we allocate uninitialized memory + unsafe { uninit_vector_real(domain_size) } + } else { + // but if some columns in the segment will remain unfilled, we allocate memory + // initialized to zeros to make sure we don't end up with memory with + // undefined values + vec![[E::BaseField::ZERO; N]; domain_size] + }; Segment::new_with_buffer(data, polys, poly_offset, offsets, twiddles) } @@ -234,7 +239,7 @@ where offsets.len().ilog2(), now.elapsed().as_millis() ); - let helper = get_wgpu_helper().unwrap(); + let helper = get_or_init_wgpu_helper().await; // build constraint evaluation commitment #[cfg(feature = "std")] @@ -246,7 +251,7 @@ where let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = RowHasher::new(helper, lde_domain_size, rpo_requires_padding, self.webgpu_hash_fn); - let mut rpo_padded_segment: Vec<[Felt; RATE]>; + let rpo_padded_segment: Vec<[Felt; RATE]>; for (segment_idx, segment) in segments.iter().enumerate() { // check if the segment requires padding if rpo_padded_segment_idx.map_or(false, |pad_idx| pad_idx == segment_idx) { @@ -263,11 +268,6 @@ where s }) .collect(); - // For rpx, skip this step - if self.webgpu_hash_fn == HashFn::Rpo256 { - let rpo_pad_column = num_base_columns % RATE; - rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); - } row_hasher.update(helper, &rpo_padded_segment); assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); break; @@ -337,6 +337,7 @@ impl< domain: &StarkDomain, webgpu_hash_fn: HashFn, ) -> (Self, TracePolyTable) { + init_wgpu_helper().await; // extend the main execution trace and build a Merkle tree from the extended trace let (main_segment_lde, main_segment_tree, main_segment_polys) = build_trace_commitment_sync::(main_trace, domain); @@ -602,6 +603,7 @@ async fn build_trace_commitment< fft::interpolate_poly(&mut poly, &inv_twiddles); poly }); + let helper = get_or_init_wgpu_helper().await; // extend the execution trace and generate hashes on the gpu let lde_segments = FrozenVec::new(); @@ -609,9 +611,8 @@ async fn build_trace_commitment< let num_base_columns = trace.num_base_cols(); let rpo_requires_padding = num_base_columns % RATE != 0; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); - let mut row_hasher = - RowHasher::new(get_wgpu_helper().unwrap(), lde_domain_size, rpo_requires_padding, hash_fn); - let mut rpo_padded_segment: Vec<[Felt; RATE]>; + let mut row_hasher = RowHasher::new(&helper, lde_domain_size, rpo_requires_padding, hash_fn); + let rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); for (segment_idx, segment) in &mut lde_segment_iter { @@ -630,21 +631,14 @@ async fn build_trace_commitment< s }) .collect(); - // skip this in case of Rpx - if hash_fn == HashFn::Rpo256 { - let rpo_pad_column = num_base_columns % RATE; - rpo_padded_segment.iter_mut().for_each(|row| row[rpo_pad_column] = ONE); - } - row_hasher.update(get_wgpu_helper().unwrap(), &rpo_padded_segment); + row_hasher.update(&helper, &rpo_padded_segment); assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); break; } - row_hasher.update(get_wgpu_helper().unwrap(), segment); + row_hasher.update(&helper, segment); } - let row_hashes = row_hasher.finish(get_wgpu_helper().unwrap()).await.unwrap(); - let tree_nodes = build_merkle_tree(get_wgpu_helper().unwrap(), &row_hashes, hash_fn) - .await - .unwrap(); + let row_hashes = row_hasher.finish(&helper).await.unwrap(); + let tree_nodes = build_merkle_tree(&helper, &row_hashes, hash_fn).await.unwrap(); // aggregate segments at the same time as the GPU generates the merkle tree nodes let lde_segments = lde_segments.into_vec().into_iter().map(|p| *p).collect(); let trace_lde = RowMatrix::from_segments(lde_segments, num_base_columns); diff --git a/prover/src/lib.rs b/prover/src/lib.rs index ee0077dd7..69ccc8809 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -9,6 +9,7 @@ use core::marker::PhantomData; use maybe_async::maybe_async; use air::{AuxRandElements, ProcessorAir, PublicInputs}; +#[cfg(any(feature = "metal", feature = "webgpu"))] use miden_gpu::HashFn; use processor::{ crypto::{ From cbdde533982521e49073d19ee7757cb269a47c11 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Tue, 20 Aug 2024 15:21:41 +0800 Subject: [PATCH 08/16] all test passed with build_trace_commitment_sync Signed-off-by: GopherJ --- Cargo.lock | 46 +++++++++++------- Cargo.toml | 12 ----- prover/Cargo.toml | 6 +-- prover/src/gpu/webgpu/mod.rs | 85 +++++++++++++++++++--------------- prover/src/gpu/webgpu/tests.rs | 1 - 5 files changed, 78 insertions(+), 72 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 292c0b4d9..c102703e2 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -248,9 +248,9 @@ checksum = "b048fb63fd8b5923fc5aa7b340d8e156aec7ec02f0c78fa8a6ddc2613f6f71de" [[package]] name = "blake3" -version = "1.5.3" +version = "1.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "e9ec96fe9a81b5e365f9db71fe00edc4fe4ca2cc7dcb7861f0603012a7caa210" +checksum = "d82033247fd8e890df8f740e407ad4d038debb9eb1f40533fffb32e7d17dc6f7" dependencies = [ "arrayref", "arrayvec", @@ -1387,7 +1387,8 @@ dependencies = [ [[package]] name = "miden-crypto" version = "0.10.0" -source = "git+https://github.com/GopherJ/miden-crypto?rev=657b4922f3abe3577d1b0de1f6ae6ad52400cb11#657b4922f3abe3577d1b0de1f6ae6ad52400cb11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d6fad06fc3af260ed3c4235821daa2132813d993f96d446856036ae97e9606dd" dependencies = [ "blake3", "cc", @@ -1414,7 +1415,7 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.2.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70#b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b#ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b" dependencies = [ "bytemuck", "elsa", @@ -2120,9 +2121,9 @@ dependencies = [ [[package]] name = "redox_users" -version = "0.4.5" +version = "0.4.6" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "bd283d9651eeda4b2a83a43c1c91b266c40fd76ecd39a50a8c630ae69dc72891" +checksum = "ba009ff324d1fc1b900bd1fdb31564febe58a8ccc8a6fdbb93b543d33b13ca43" dependencies = [ "getrandom", "libredox", @@ -2815,9 +2816,9 @@ checksum = "0336d538f7abc86d282a4189614dfaa90810dfc2c6f6427eaf88e16311dd225d" [[package]] name = "unicode-xid" -version = "0.2.4" +version = "0.2.5" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f962df74c8c05a667b5ee8bcf162993134c104e96440b663c8daa176dc772d8c" +checksum = "229730647fbc343e3a80e463c1db7f78f3855d3f3739bee0dda773c9a037c90a" [[package]] name = "utf8parse" @@ -3389,7 +3390,8 @@ dependencies = [ [[package]] name = "winter-air" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b72f12b88ebb060b52c0e9aece9bb64a9fc38daf7ba689dd5ce63271b456c883" dependencies = [ "libm", "winter-crypto", @@ -3401,7 +3403,8 @@ dependencies = [ [[package]] name = "winter-crypto" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "00fbb724d2d9fbfd3aa16ea27f5e461d4fe1d74b0c9e0ed1bf79e9e2a955f4d5" dependencies = [ "blake3", "sha3", @@ -3412,7 +3415,8 @@ dependencies = [ [[package]] name = "winter-fri" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3ab6077cf4c23c0411f591f4ba29378e27f26acb8cef3c51cadd93daaf6080b3" dependencies = [ "winter-crypto", "winter-math", @@ -3422,7 +3426,8 @@ dependencies = [ [[package]] name = "winter-math" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "004f85bb051ce986ec0b9a2bd90aaf81b83e3c67464becfdf7db31f14c1019ba" dependencies = [ "winter-utils", ] @@ -3430,7 +3435,8 @@ dependencies = [ [[package]] name = "winter-maybe-async" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7ce0f4161cdde50de809b3869c1cb083a09e92e949428ea28f04c0d64045875c" dependencies = [ "proc-macro2", "quote", @@ -3440,7 +3446,8 @@ dependencies = [ [[package]] name = "winter-prover" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f17e3dbae97050f58e01ed4f12906e247841575a0518632e052941a1c37468df" dependencies = [ "async-trait", "tracing", @@ -3455,7 +3462,8 @@ dependencies = [ [[package]] name = "winter-rand-utils" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "f2b827c901ab0c316d89812858ff451d60855c0a5c7ae734b098c62a28624181" dependencies = [ "rand", "winter-utils", @@ -3463,8 +3471,9 @@ dependencies = [ [[package]] name = "winter-utils" -version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +version = "0.9.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0568612a95bcae3c94fb14da2686f8279ca77723dbdf1e97cf3673798faf6485" dependencies = [ "rayon", ] @@ -3472,7 +3481,8 @@ dependencies = [ [[package]] name = "winter-verifier" version = "0.9.0" -source = "git+https://github.com/GopherJ/winterfell?rev=50a9a088d30ff58e88cf19805f4dfa13b4be356e#50a9a088d30ff58e88cf19805f4dfa13b4be356e" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "324002ade90f21e85599d51a232a80781efc8cb46f511f8bc89f9c5a4eb9cb65" dependencies = [ "winter-air", "winter-crypto", diff --git a/Cargo.toml b/Cargo.toml index d7e0c24ef..ac69f241e 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -32,15 +32,3 @@ inherits = "release" debug = true debug-assertions = true overflow-checks = true - -[patch.crates-io] -winter-prover = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-air = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-utils = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-verifier = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-fri = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-crypto = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-math = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-maybe-async = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -winter-rand-utils = { git = "https://github.com/GopherJ/winterfell", rev = "50a9a088d30ff58e88cf19805f4dfa13b4be356e" } -miden-crypto = { git = "https://github.com/GopherJ/miden-crypto", rev = "657b4922f3abe3577d1b0de1f6ae6ad52400cb11" } diff --git a/prover/Cargo.toml b/prover/Cargo.toml index dbb944ea9..447fdaad0 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["std"] +default = ["webgpu"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] @@ -34,11 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "b8b3a4273dd76a3af2ba8cb7739c9e20823a8f70", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index f1ef5d4b2..25e13030e 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -19,7 +19,7 @@ use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use elsa::FrozenVec; use maybe_async::maybe_async; use miden_gpu::{ - webgpu::{build_merkle_tree, get_or_init_wgpu_helper, init_wgpu_helper, RowHasher}, + webgpu::{build_merkle_tree, get_or_init_wgpu_helper, RowHasher}, HashFn, }; use processor::{ @@ -27,6 +27,7 @@ use processor::{ ONE, }; use tracing::info_span; +#[cfg(feature = "std")] use tracing::{event, Level}; use winter_prover::{ crypto::{Digest, MerkleTree}, @@ -248,9 +249,10 @@ where let num_base_columns = composition_poly.num_columns() * ::EXTENSION_DEGREE; let rpo_requires_padding = num_base_columns % RATE != 0; + let is_rpo = self.webgpu_hash_fn == HashFn::Rpo256; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = - RowHasher::new(helper, lde_domain_size, rpo_requires_padding, self.webgpu_hash_fn); + RowHasher::new(helper, lde_domain_size, num_base_columns, self.webgpu_hash_fn); let rpo_padded_segment: Vec<[Felt; RATE]>; for (segment_idx, segment) in segments.iter().enumerate() { // check if the segment requires padding @@ -260,14 +262,18 @@ where // padded with "0"s we only need to add the "1"s. let rpo_pad_column = num_base_columns % RATE; - rpo_padded_segment = segment - .iter() - .map(|x| { - let mut s = *x; - s[rpo_pad_column] = ONE; - s - }) - .collect(); + rpo_padded_segment = if is_rpo { + segment + .iter() + .map(|x| { + let mut s = *x; + s[rpo_pad_column] = ONE; + s + }) + .collect() + } else { + segment.iter().map(|x| *x).collect() + }; row_hasher.update(helper, &rpo_padded_segment); assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); break; @@ -337,7 +343,6 @@ impl< domain: &StarkDomain, webgpu_hash_fn: HashFn, ) -> (Self, TracePolyTable) { - init_wgpu_helper().await; // extend the main execution trace and build a Merkle tree from the extended trace let (main_segment_lde, main_segment_tree, main_segment_polys) = build_trace_commitment_sync::(main_trace, domain); @@ -410,28 +415,27 @@ impl< /// This function will panic if any of the following are true: /// - the number of rows in the provided `aux_trace` does not match the main trace. /// - this segment would exceed the number of segments specified by the trace layout. - async fn set_aux_trace( + fn set_aux_trace( &mut self, aux_trace: &ColMatrix, domain: &StarkDomain, ) -> (ColMatrix, D) { - todo!() - // // extend the auxiliary trace segment and build a Merkle tree from the extended trace - // let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = - // build_trace_commitment::(aux_trace, domain, self.webgpu_hash_fn).await; - // - // assert_eq!( - // self.main_segment_lde.num_rows(), - // aux_segment_lde.num_rows(), - // "the number of rows in the auxiliary segment must be the same as in the main segment" - // ); - // - // // save the lde and commitment - // self.aux_segment_lde = Some(aux_segment_lde); - // let root_hash = *aux_segment_tree.root(); - // self.aux_segment_tree = Some(aux_segment_tree); - // - // (aux_segment_polys, root_hash) + // extend the auxiliary trace segment and build a Merkle tree from the extended trace + let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = + build_trace_commitment_sync::(aux_trace, domain); + + assert_eq!( + self.main_segment_lde.num_rows(), + aux_segment_lde.num_rows(), + "the number of rows in the auxiliary segment must be the same as in the main segment" + ); + + // save the lde and commitment + self.aux_segment_lde = Some(aux_segment_lde); + let root_hash = *aux_segment_tree.root(); + self.aux_segment_tree = Some(aux_segment_tree); + + (aux_segment_polys, root_hash) } /// Reads current and next rows from the main trace segment into the specified frame. @@ -610,8 +614,9 @@ async fn build_trace_commitment< let lde_domain_size = domain.lde_domain_size(); let num_base_columns = trace.num_base_cols(); let rpo_requires_padding = num_base_columns % RATE != 0; + let is_rpo = hash_fn == HashFn::Rpo256; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); - let mut row_hasher = RowHasher::new(&helper, lde_domain_size, rpo_requires_padding, hash_fn); + let mut row_hasher = RowHasher::new(&helper, lde_domain_size, num_base_columns, hash_fn); let rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); @@ -623,14 +628,18 @@ async fn build_trace_commitment< // rule ("1" followed by "0"s). Our segments are already // padded with "0"s we only need to add the "1"s. let rpo_pad_column = num_base_columns % RATE; - rpo_padded_segment = segment - .iter() - .map(|x| { - let mut s = *x; - s[rpo_pad_column] = ONE; - s - }) - .collect(); + rpo_padded_segment = if is_rpo { + segment + .iter() + .map(|x| { + let mut s = *x; + s[rpo_pad_column] = ONE; + s + }) + .collect() + } else { + segment.iter().map(|x| *x).collect() + }; row_hasher.update(&helper, &rpo_padded_segment); assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); break; diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs index 95da5ecd2..6abdf950d 100644 --- a/prover/src/gpu/webgpu/tests.rs +++ b/prover/src/gpu/webgpu/tests.rs @@ -13,7 +13,6 @@ use winter_prover::{crypto::Digest, math::fields::CubeExtension, CompositionPoly use crate::*; -use wasm_bindgen::prelude::*; use wasm_bindgen_test::*; wasm_bindgen_test_configure!(run_in_browser); From bbd1cf878982de849055b161449f30f5d44c34df Mon Sep 17 00:00:00 2001 From: GopherJ Date: Sat, 24 Aug 2024 10:24:34 +0800 Subject: [PATCH 09/16] update miden-gpu rev Signed-off-by: GopherJ --- Cargo.lock | 47 ++++++++++++++++++----------------------------- prover/Cargo.toml | 6 +++--- 2 files changed, 21 insertions(+), 32 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index c102703e2..ace1ac98b 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1303,21 +1303,6 @@ version = "0.4.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8452105ba047068f40ff7093dd1d9da90898e63dd61736462e9cdda6a90ad3c3" -[[package]] -name = "metal" -version = "0.27.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "c43f73953f8cbe511f021b58f18c3ce1c3d1ae13fe953293e13345bf83217f25" -dependencies = [ - "bitflags 2.6.0", - "block", - "core-graphics-types", - "foreign-types", - "log", - "objc", - "paste", -] - [[package]] name = "metal" version = "0.29.0" @@ -1414,16 +1399,19 @@ dependencies = [ [[package]] name = "miden-gpu" -version = "0.2.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b#ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b" +version = "0.3.0" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=71b573879024d3626e399926b89199b7549c0e13#71b573879024d3626e399926b89199b7549c0e13" dependencies = [ "bytemuck", + "cc", "elsa", "flume", - "metal 0.27.0", + "miden-crypto", "once_cell", + "tracing", "wgpu", "winter-math", + "winterfell", ] [[package]] @@ -1783,16 +1771,6 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" dependencies = [ "malloc_buf", - "objc_exception", -] - -[[package]] -name = "objc_exception" -version = "0.1.2" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "ad970fb455818ad6cba4c122ad012fae53ae8b4795f86378bce65e4f6bab2ca4" -dependencies = [ - "cc", ] [[package]] @@ -3080,7 +3058,7 @@ dependencies = [ "libc", "libloading", "log", - "metal 0.29.0", + "metal", "naga", "ndk-sys", "objc", @@ -3491,6 +3469,17 @@ dependencies = [ "winter-utils", ] +[[package]] +name = "winterfell" +version = "0.9.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "01151ac5fe2d783950743e8a110e0a2f26994f888b4cbe848699142cb3ea1e5b" +dependencies = [ + "winter-air", + "winter-prover", + "winter-verifier", +] + [[package]] name = "xml-rs" version = "0.8.21" diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 447fdaad0..c353ed11c 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["webgpu"] +default = ["std"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] @@ -34,11 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "71b573879024d3626e399926b89199b7549c0e13", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "ee2e51bf1ed54a40212d60ddb55a5b97bfdfae5b", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "71b573879024d3626e399926b89199b7549c0e13", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" From 169d42e4bbe1de7285bb3b08ca78ac35e378750e Mon Sep 17 00:00:00 2001 From: GopherJ Date: Mon, 26 Aug 2024 12:04:03 +0800 Subject: [PATCH 10/16] update miden-gpu rev Signed-off-by: GopherJ --- Cargo.lock | 2 +- prover/Cargo.toml | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 8a24d96ed..6972188b7 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1400,7 +1400,7 @@ dependencies = [ [[package]] name = "miden-gpu" version = "0.3.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=71b573879024d3626e399926b89199b7549c0e13#71b573879024d3626e399926b89199b7549c0e13" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=834c7973893bf6d6a93f4035de4fbf6faadc9c1d#834c7973893bf6d6a93f4035de4fbf6faadc9c1d" dependencies = [ "bytemuck", "cc", diff --git a/prover/Cargo.toml b/prover/Cargo.toml index a51ac75f1..928bbaaf3 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["std"] +default = ["webgpu","std"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] @@ -34,11 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "71b573879024d3626e399926b89199b7549c0e13", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "834c7973893bf6d6a93f4035de4fbf6faadc9c1d", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "71b573879024d3626e399926b89199b7549c0e13", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "834c7973893bf6d6a93f4035de4fbf6faadc9c1d", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" From a116707258e47338fffc3e4c4bbf245ea020b952 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Thu, 10 Oct 2024 11:59:38 +0800 Subject: [PATCH 11/16] use new api Signed-off-by: GopherJ --- prover/Cargo.toml | 4 ++-- prover/src/gpu/webgpu/mod.rs | 18 +++++++++--------- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 928bbaaf3..efbd1558a 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -34,11 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "834c7973893bf6d6a93f4035de4fbf6faadc9c1d", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "389dd1ed47b9b345e978611421afa2eab722e9c4", default-features = false, optional = true } pollster = { version = "0.3", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "834c7973893bf6d6a93f4035de4fbf6faadc9c1d", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "389dd1ed47b9b345e978611421afa2eab722e9c4", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index 25e13030e..dbf97aa84 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -19,7 +19,7 @@ use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; use elsa::FrozenVec; use maybe_async::maybe_async; use miden_gpu::{ - webgpu::{build_merkle_tree, get_or_init_wgpu_helper, RowHasher}, + webgpu::{get_or_init_wgpu_helper, RowHasher}, HashFn, }; use processor::{ @@ -274,14 +274,14 @@ where } else { segment.iter().map(|x| *x).collect() }; - row_hasher.update(helper, &rpo_padded_segment); + row_hasher.update(&rpo_padded_segment); assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); break; } - row_hasher.update(helper, segment); + row_hasher.update(segment); } - let row_hashes = row_hasher.finish(helper).await.unwrap(); - let tree_nodes = build_merkle_tree(helper, &row_hashes, self.webgpu_hash_fn).await.unwrap(); + let row_hashes = row_hasher.finish().await; + let tree_nodes = helper.build_merkle_tree(&row_hashes, self.webgpu_hash_fn).await; // aggregate segments at the same time as the GPU generates the merkle tree nodes let composed_evaluations = RowMatrix::::from_segments(segments, num_base_columns); let nodes = tree_nodes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); @@ -640,14 +640,14 @@ async fn build_trace_commitment< } else { segment.iter().map(|x| *x).collect() }; - row_hasher.update(&helper, &rpo_padded_segment); + row_hasher.update(&rpo_padded_segment); assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); break; } - row_hasher.update(&helper, segment); + row_hasher.update(segment); } - let row_hashes = row_hasher.finish(&helper).await.unwrap(); - let tree_nodes = build_merkle_tree(&helper, &row_hashes, hash_fn).await.unwrap(); + let row_hashes = row_hasher.finish().await; + let tree_nodes = helper.build_merkle_tree(&row_hashes, hash_fn).await; // aggregate segments at the same time as the GPU generates the merkle tree nodes let lde_segments = lde_segments.into_vec().into_iter().map(|p| *p).collect(); let trace_lde = RowMatrix::from_segments(lde_segments, num_base_columns); From 725e97f8f5c5a7617e479cdbdd21334a0cffcd89 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Thu, 10 Oct 2024 12:11:30 +0800 Subject: [PATCH 12/16] fix linux compile --- prover/Cargo.toml | 2 +- prover/src/gpu/mod.rs | 5 ++++- prover/src/gpu/webgpu/mod.rs | 3 --- prover/src/lib.rs | 18 ++++++++++++------ 4 files changed, 17 insertions(+), 11 deletions(-) diff --git a/prover/Cargo.toml b/prover/Cargo.toml index efbd1558a..a5a5f6643 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -18,7 +18,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] -default = ["webgpu","std"] +default = ["std"] metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] diff --git a/prover/src/gpu/mod.rs b/prover/src/gpu/mod.rs index c57b8b9cc..9efe3ee7f 100644 --- a/prover/src/gpu/mod.rs +++ b/prover/src/gpu/mod.rs @@ -1,5 +1,8 @@ #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] pub mod metal; -#[cfg(all(feature = "webgpu"))] +#[cfg(all( + feature = "webgpu", + any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") +))] pub mod webgpu; diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index dbf97aa84..8072ef422 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -7,9 +7,6 @@ use core::marker::PhantomData; #[cfg(feature = "std")] use std::{boxed::Box, marker::PhantomData, time::Instant, vec::Vec}; -#[cfg(not(feature = "std"))] -extern crate alloc; - #[cfg(not(feature = "std"))] use alloc::boxed::Box; #[cfg(not(feature = "std"))] diff --git a/prover/src/lib.rs b/prover/src/lib.rs index 3667a1420..509db836a 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -9,8 +9,6 @@ use core::marker::PhantomData; use maybe_async::maybe_async; use air::{AuxRandElements, ProcessorAir, PublicInputs}; -#[cfg(any(feature = "metal", feature = "webgpu"))] -use miden_gpu::HashFn; use processor::{ crypto::{ Blake3_192, Blake3_256, ElementHasher, RandomCoin, Rpo256, RpoRandomCoin, Rpx256, @@ -108,8 +106,12 @@ where ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpo256); - #[cfg(feature = "webgpu")] - let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, HashFn::Rpo256); + #[cfg(all( + feature = "webgpu", + any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") + ))] + let prover = + gpu::webgpu::WebGPUExecutionProver::new(prover, gpu::webgpu::HashFn::Rpo256); prover.prove(trace).await }, HashFunction::Rpx256 => { @@ -120,8 +122,12 @@ where ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpx256); - #[cfg(feature = "webgpu")] - let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, HashFn::Rpx256); + #[cfg(all( + feature = "webgpu", + any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") + ))] + let prover = + gpu::webgpu::WebGPUExecutionProver::new(prover, gpu::webgpu::HashFn::Rpx256); prover.prove(trace).await }, } From d71bd2e4d852d9ebdfcbcadca2aa0c691aafd0cb Mon Sep 17 00:00:00 2001 From: GopherJ Date: Thu, 10 Oct 2024 12:23:55 +0800 Subject: [PATCH 13/16] fix unresolved hash fn Signed-off-by: GopherJ --- Cargo.lock | 1 + prover/Cargo.toml | 2 +- prover/src/lib.rs | 8 ++++---- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index effb23f57..a6a8e524b 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1412,6 +1412,7 @@ dependencies = [ "cc", "elsa", "flume", + "metal", "miden-crypto", "once_cell", "tracing", diff --git a/prover/Cargo.toml b/prover/Cargo.toml index a5a5f6643..ac5d6c483 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -19,7 +19,7 @@ crate-type = ["cdylib", "rlib"] [features] concurrent = ["processor/concurrent", "std", "winter-prover/concurrent"] default = ["std"] -metal = ["dep:miden-gpu", "dep:pollster", "concurrent", "std"] +metal = ["dep:miden-gpu", "miden-gpu/metal", "dep:pollster", "concurrent", "std"] webgpu = ["dep:miden-gpu", "miden-gpu/webgpu", "async"] async = ["maybe-async/async", "winter-prover/async", "dep:async-trait"] std = ["air/std", "processor/std", "winter-prover/std", "miden-gpu/std", "tracing/std"] diff --git a/prover/src/lib.rs b/prover/src/lib.rs index 509db836a..deb0ccf3f 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -105,13 +105,13 @@ where stack_outputs.clone(), ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] - let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpo256); + let prover = gpu::metal::MetalExecutionProver::new(prover, miden_gpu::HashFn::Rpo256); #[cfg(all( feature = "webgpu", any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") ))] let prover = - gpu::webgpu::WebGPUExecutionProver::new(prover, gpu::webgpu::HashFn::Rpo256); + gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpo256); prover.prove(trace).await }, HashFunction::Rpx256 => { @@ -121,13 +121,13 @@ where stack_outputs.clone(), ); #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] - let prover = gpu::metal::MetalExecutionProver::new(prover, HashFn::Rpx256); + let prover = gpu::metal::MetalExecutionProver::new(prover, miden_gpu::HashFn::Rpx256); #[cfg(all( feature = "webgpu", any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") ))] let prover = - gpu::webgpu::WebGPUExecutionProver::new(prover, gpu::webgpu::HashFn::Rpx256); + gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpx256); prover.prove(trace).await }, } From d8260103382dce65fb8cc5e8a6a0052de43b8bd5 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Thu, 10 Oct 2024 12:29:35 +0800 Subject: [PATCH 14/16] fix tests Signed-off-by: GopherJ --- prover/src/gpu/webgpu/tests.rs | 2 ++ 1 file changed, 2 insertions(+) diff --git a/prover/src/gpu/webgpu/tests.rs b/prover/src/gpu/webgpu/tests.rs index 6abdf950d..3904cbec5 100644 --- a/prover/src/gpu/webgpu/tests.rs +++ b/prover/src/gpu/webgpu/tests.rs @@ -13,6 +13,8 @@ use winter_prover::{crypto::Digest, math::fields::CubeExtension, CompositionPoly use crate::*; +use miden_gpu::HashFn; + use wasm_bindgen_test::*; wasm_bindgen_test_configure!(run_in_browser); From dee76c62801cd8aa2351a07d7f49f10308067893 Mon Sep 17 00:00:00 2001 From: GopherJ Date: Tue, 19 Nov 2024 15:36:43 +0800 Subject: [PATCH 15/16] fix build Signed-off-by: GopherJ --- Cargo.lock | 226 ++++++++++------------------------- prover/Cargo.toml | 5 +- prover/src/gpu/webgpu/mod.rs | 122 +++++++++++-------- prover/src/lib.rs | 9 +- 4 files changed, 145 insertions(+), 217 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 63be06c2e..ce95e1e28 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1324,6 +1324,21 @@ dependencies = [ "paste", ] +[[package]] +name = "metal" +version = "0.30.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "9c3572083504c43e14aec05447f8a3d57cce0f66d7a3c1b9058572eca4d70ab9" +dependencies = [ + "bitflags 2.6.0", + "block", + "core-graphics-types", + "foreign-types", + "log", + "objc", + "paste", +] + [[package]] name = "miden-air" version = "0.11.0" @@ -1332,8 +1347,8 @@ dependencies = [ "miden-core", "miden-thiserror", "proptest", - "winter-air 0.10.2", - "winter-prover 0.10.2", + "winter-air", + "winter-prover", "winter-rand-utils", ] @@ -1362,7 +1377,7 @@ dependencies = [ "lock_api", "loom", "memchr", - "miden-crypto 0.12.0", + "miden-crypto", "miden-formatting", "miden-miette", "miden-thiserror", @@ -1370,28 +1385,9 @@ dependencies = [ "num-traits", "parking_lot", "proptest", - "winter-math 0.10.2", + "winter-math", "winter-rand-utils", - "winter-utils 0.10.2", -] - -[[package]] -name = "miden-crypto" -version = "0.10.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "825fc5d2e4c951f45da54da2d0d23071918d966133f85dbc219694e0b9a1e141" -dependencies = [ - "blake3", - "cc", - "glob", - "num", - "num-complex", - "rand", - "rand_core", - "sha3", - "winter-crypto 0.9.0", - "winter-math 0.9.3", - "winter-utils 0.9.3", + "winter-utils", ] [[package]] @@ -1408,9 +1404,9 @@ dependencies = [ "rand", "rand_core", "sha3", - "winter-crypto 0.10.2", - "winter-math 0.10.2", - "winter-utils 0.10.2", + "winter-crypto", + "winter-math", + "winter-utils", ] [[package]] @@ -1424,19 +1420,19 @@ dependencies = [ [[package]] name = "miden-gpu" -version = "0.3.0" -source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=389dd1ed47b9b345e978611421afa2eab722e9c4#389dd1ed47b9b345e978611421afa2eab722e9c4" +version = "0.4.0" +source = "git+https://github.com/0xPolygonMiden/miden-gpu?rev=471af82af2ba48b0adb2f258c213215461aa8da8#471af82af2ba48b0adb2f258c213215461aa8da8" dependencies = [ "bytemuck", "cc", "elsa", "flume", - "metal", - "miden-crypto 0.10.3", + "metal 0.30.0", + "miden-crypto", "once_cell", "tracing", "wgpu", - "winter-math 0.9.3", + "winter-math", "winterfell", ] @@ -1492,9 +1488,9 @@ dependencies = [ "miden-core", "miden-test-utils", "tracing", - "winter-fri 0.10.2", - "winter-prover 0.10.2", - "winter-utils 0.10.2", + "winter-fri", + "winter-prover", + "winter-utils", ] [[package]] @@ -1513,8 +1509,8 @@ dependencies = [ "wasm-bindgen-futures", "wasm-bindgen-test", "wee_alloc", - "winter-maybe-async 0.10.1", - "winter-prover 0.10.2", + "winter-maybe-async", + "winter-prover", ] [[package]] @@ -1534,8 +1530,8 @@ dependencies = [ "serde_json", "sha2", "sha3", - "winter-air 0.10.2", - "winter-fri 0.10.2", + "winter-air", + "winter-fri", ] [[package]] @@ -1551,7 +1547,7 @@ dependencies = [ "pretty_assertions", "proptest", "test-case", - "winter-prover 0.10.2", + "winter-prover", "winter-rand-utils", ] @@ -1582,7 +1578,7 @@ dependencies = [ "miden-air", "miden-core", "tracing", - "winter-verifier 0.10.2", + "winter-verifier", ] [[package]] @@ -1612,7 +1608,7 @@ dependencies = [ "tracing", "tracing-forest", "tracing-subscriber", - "winter-fri 0.10.2", + "winter-fri", ] [[package]] @@ -3098,7 +3094,7 @@ dependencies = [ "libc", "libloading", "log", - "metal", + "metal 0.29.0", "naga", "ndk-sys", "objc", @@ -3405,19 +3401,6 @@ dependencies = [ "memchr", ] -[[package]] -name = "winter-air" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "b72f12b88ebb060b52c0e9aece9bb64a9fc38daf7ba689dd5ce63271b456c883" -dependencies = [ - "libm", - "winter-crypto 0.9.0", - "winter-fri 0.9.0", - "winter-math 0.9.3", - "winter-utils 0.9.3", -] - [[package]] name = "winter-air" version = "0.10.2" @@ -3425,22 +3408,10 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "74ea405623248442dc08290afcce9a1b38683274e11952e6fcad6f10c22ff81f" dependencies = [ "libm", - "winter-crypto 0.10.2", - "winter-fri 0.10.2", - "winter-math 0.10.2", - "winter-utils 0.10.2", -] - -[[package]] -name = "winter-crypto" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "00fbb724d2d9fbfd3aa16ea27f5e461d4fe1d74b0c9e0ed1bf79e9e2a955f4d5" -dependencies = [ - "blake3", - "sha3", - "winter-math 0.9.3", - "winter-utils 0.9.3", + "winter-crypto", + "winter-fri", + "winter-math", + "winter-utils", ] [[package]] @@ -3451,19 +3422,8 @@ checksum = "3fcae1ada055aa10554910ecffc106cb116a19dba11ac91390ef982f94adb9c5" dependencies = [ "blake3", "sha3", - "winter-math 0.10.2", - "winter-utils 0.10.2", -] - -[[package]] -name = "winter-fri" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "3ab6077cf4c23c0411f591f4ba29378e27f26acb8cef3c51cadd93daaf6080b3" -dependencies = [ - "winter-crypto 0.9.0", - "winter-math 0.9.3", - "winter-utils 0.9.3", + "winter-math", + "winter-utils", ] [[package]] @@ -3472,18 +3432,9 @@ version = "0.10.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ff88657560100f34fb83882a0adf33fb7caee235deb83193d0d251ddb28ed9c9" dependencies = [ - "winter-crypto 0.10.2", - "winter-math 0.10.2", - "winter-utils 0.10.2", -] - -[[package]] -name = "winter-math" -version = "0.9.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "5b0e685b3b872d82e58a86519294a814b7bc7a4d3cd2c93570a7d80c0c5a1aba" -dependencies = [ - "winter-utils 0.9.3", + "winter-crypto", + "winter-math", + "winter-utils", ] [[package]] @@ -3492,18 +3443,7 @@ version = "0.10.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "82479f94efc0b5374a93e2074ba46ef404384fb1ea6e35a847febec53096509b" dependencies = [ - "winter-utils 0.10.2", -] - -[[package]] -name = "winter-maybe-async" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7ce0f4161cdde50de809b3869c1cb083a09e92e949428ea28f04c0d64045875c" -dependencies = [ - "proc-macro2", - "quote", - "syn 2.0.87", + "winter-utils", ] [[package]] @@ -3516,21 +3456,6 @@ dependencies = [ "syn 2.0.87", ] -[[package]] -name = "winter-prover" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "f17e3dbae97050f58e01ed4f12906e247841575a0518632e052941a1c37468df" -dependencies = [ - "tracing", - "winter-air 0.9.0", - "winter-crypto 0.9.0", - "winter-fri 0.9.0", - "winter-math 0.9.3", - "winter-maybe-async 0.9.0", - "winter-utils 0.9.3", -] - [[package]] name = "winter-prover" version = "0.10.2" @@ -3538,12 +3463,12 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "1355271e7e1eef05cbb06fae85f6300321d37eb19d97d7d222a4aa3c16079591" dependencies = [ "tracing", - "winter-air 0.10.2", - "winter-crypto 0.10.2", - "winter-fri 0.10.2", - "winter-math 0.10.2", - "winter-maybe-async 0.10.1", - "winter-utils 0.10.2", + "winter-air", + "winter-crypto", + "winter-fri", + "winter-math", + "winter-maybe-async", + "winter-utils", ] [[package]] @@ -3553,15 +3478,9 @@ source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "4a7616d11fcc26552dada45c803a884ac97c253218835b83a2c63e1c2a988639" dependencies = [ "rand", - "winter-utils 0.10.2", + "winter-utils", ] -[[package]] -name = "winter-utils" -version = "0.9.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "961e81e9388877a25db1c034ba38253de2055f569633ae6a665d857a0556391b" - [[package]] name = "winter-utils" version = "0.10.2" @@ -3571,41 +3490,28 @@ dependencies = [ "rayon", ] -[[package]] -name = "winter-verifier" -version = "0.9.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "324002ade90f21e85599d51a232a80781efc8cb46f511f8bc89f9c5a4eb9cb65" -dependencies = [ - "winter-air 0.9.0", - "winter-crypto 0.9.0", - "winter-fri 0.9.0", - "winter-math 0.9.3", - "winter-utils 0.9.3", -] - [[package]] name = "winter-verifier" version = "0.10.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8003bd5489d5d74036b1d71b2cb8d4bcb54dd05a753bfd2b20527c37cb1f63f0" dependencies = [ - "winter-air 0.10.2", - "winter-crypto 0.10.2", - "winter-fri 0.10.2", - "winter-math 0.10.2", - "winter-utils 0.10.2", + "winter-air", + "winter-crypto", + "winter-fri", + "winter-math", + "winter-utils", ] [[package]] name = "winterfell" -version = "0.9.0" +version = "0.10.2" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "01151ac5fe2d783950743e8a110e0a2f26994f888b4cbe848699142cb3ea1e5b" +checksum = "3f41df9fa738e256cd089f42d0822610488e92d890a6c8169aac60270c5d017c" dependencies = [ - "winter-air 0.9.0", - "winter-prover 0.9.0", - "winter-verifier 0.9.0", + "winter-air", + "winter-prover", + "winter-verifier", ] [[package]] diff --git a/prover/Cargo.toml b/prover/Cargo.toml index 6c29de68e..7f0c2c406 100644 --- a/prover/Cargo.toml +++ b/prover/Cargo.toml @@ -34,12 +34,11 @@ async-trait = { version = "0.1", optional = true } elsa = { version = "1.9" } [target.'cfg(all(target_arch = "aarch64", target_os = "macos"))'.dependencies] -elsa = { version = "1.9", optional = true } -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "389dd1ed47b9b345e978611421afa2eab722e9c4", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "471af82af2ba48b0adb2f258c213215461aa8da8", default-features = false, optional = true } pollster = { version = "0.4", optional = true } [target.'cfg(target_family = "wasm")'.dependencies] -miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "389dd1ed47b9b345e978611421afa2eab722e9c4", default-features = false, optional = true } +miden-gpu = { git = "https://github.com/0xPolygonMiden/miden-gpu", rev = "471af82af2ba48b0adb2f258c213215461aa8da8", default-features = false, optional = true } [target.'cfg(target_family = "wasm")'.dev-dependencies] wasm-bindgen-test = "0.3" diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index 8072ef422..d74a048b6 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -12,11 +12,10 @@ use alloc::boxed::Box; #[cfg(not(feature = "std"))] use alloc::{vec, vec::Vec}; -use air::{AuxRandElements, LagrangeKernelEvaluationFrame}; +use air::{AuxRandElements, LagrangeKernelEvaluationFrame, PartitionOptions}; use elsa::FrozenVec; -use maybe_async::maybe_async; use miden_gpu::{ - webgpu::{get_or_init_wgpu_helper, RowHasher}, + webgpu::{get_or_init_webgpu, RowHasher}, HashFn, }; use processor::{ @@ -26,6 +25,8 @@ use processor::{ use tracing::info_span; #[cfg(feature = "std")] use tracing::{event, Level}; +use winter_maybe_async::{maybe_async, maybe_await}; +use winter_prover::crypto::VectorCommitment; use winter_prover::{ crypto::{Digest, MerkleTree}, matrix::{get_evaluation_offsets, ColMatrix, RowMatrix, Segment}, @@ -76,7 +77,7 @@ where impl WebGPUExecutionProver where - H: Hasher + ElementHasher, + H: Hasher + ElementHasher + Sync, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, R: RandomCoin + Send, { @@ -145,13 +146,14 @@ where #[maybe_async] impl Prover for WebGPUExecutionProver where - H: Hasher + ElementHasher, + H: Hasher + ElementHasher + Sync, D: Digest + for<'a> From<&'a [Felt; DIGEST_SIZE]>, R: RandomCoin + Send, { type BaseField = Felt; type Air = ProcessorAir; type Trace = ExecutionTrace; + type VC = MerkleTree; type HashFn = H; type RandomCoin = R; type TraceLde> = WebGPUTraceLde; @@ -166,24 +168,35 @@ where self.execution_prover.options() } - async fn new_trace_lde>( + #[maybe_async] + fn new_trace_lde>( &self, trace_info: &TraceInfo, main_trace: &ColMatrix, domain: &StarkDomain, + partition_option: PartitionOptions, ) -> (Self::TraceLde, TracePolyTable) { - WebGPUTraceLde::new(trace_info, main_trace, domain, self.webgpu_hash_fn).await + maybe_await!(WebGPUTraceLde::new( + trace_info, + main_trace, + domain, + self.webgpu_hash_fn, + partition_option + )) } - async fn new_evaluator<'a, E: FieldElement>( + #[maybe_async] + fn new_evaluator<'a, E: FieldElement>( &self, air: &'a ProcessorAir, aux_rand_elements: Option>, composition_coefficients: ConstraintCompositionCoefficients, ) -> Self::ConstraintEvaluator<'a, E> { - self.execution_prover - .new_evaluator(air, aux_rand_elements, composition_coefficients) - .await + maybe_await!(self.execution_prover.new_evaluator( + air, + aux_rand_elements, + composition_coefficients + )) } /// Evaluates constraint composition polynomial over the LDE domain and builds a commitment @@ -210,12 +223,16 @@ where /// ────┼────────┼────────┼────────┼────────┼────────┼─── /// t=n t=n+1 t=n+2 t=n+3 t=n+4 t=n+5 /// ``` - async fn build_constraint_commitment>( + #[maybe_async] + fn build_constraint_commitment>( &self, composition_poly_trace: CompositionPolyTrace, num_trace_poly_columns: usize, domain: &StarkDomain, - ) -> (ConstraintCommitment, CompositionPoly) { + ) -> ( + ConstraintCommitment>, + CompositionPoly, + ) { // evaluate composition polynomial columns over the LDE domain #[cfg(feature = "std")] let now = Instant::now(); @@ -237,7 +254,7 @@ where offsets.len().ilog2(), now.elapsed().as_millis() ); - let helper = get_or_init_wgpu_helper().await; + let helper = maybe_await!(get_or_init_webgpu()); // build constraint evaluation commitment #[cfg(feature = "std")] @@ -249,7 +266,7 @@ where let is_rpo = self.webgpu_hash_fn == HashFn::Rpo256; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = - RowHasher::new(helper, lde_domain_size, num_base_columns, self.webgpu_hash_fn); + RowHasher::::new(helper, lde_domain_size, num_base_columns, self.webgpu_hash_fn); let rpo_padded_segment: Vec<[Felt; RATE]>; for (segment_idx, segment) in segments.iter().enumerate() { // check if the segment requires padding @@ -277,19 +294,18 @@ where } row_hasher.update(segment); } - let row_hashes = row_hasher.finish().await; - let tree_nodes = helper.build_merkle_tree(&row_hashes, self.webgpu_hash_fn).await; + let row_hashes = maybe_await!(row_hasher.finish()); + let tree_nodes = + maybe_await!(helper.build_merkle_tree::(&row_hashes, self.webgpu_hash_fn)); // aggregate segments at the same time as the GPU generates the merkle tree nodes let composed_evaluations = RowMatrix::::from_segments(segments, num_base_columns); - let nodes = tree_nodes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); - let leaves = row_hashes.into_iter().map(|dig| H::Digest::from(&dig)).collect(); - let commitment = MerkleTree::::from_raw_parts(nodes, leaves).unwrap(); + let commitment = MerkleTree::::from_raw_parts(tree_nodes, row_hashes).unwrap(); let constraint_commitment = ConstraintCommitment::new(composed_evaluations, commitment); #[cfg(feature = "std")] event!( Level::INFO, "Computed constraint evaluation commitment on the GPU (Merkle tree of depth {}) in {} ms", - constraint_commitment.tree_depth(), + lde_domain_size.ilog2(), now.elapsed().as_millis() ); (constraint_commitment, composition_poly) @@ -319,6 +335,7 @@ pub struct WebGPUTraceLde, H: Hasher> { blowup: usize, trace_info: TraceInfo, webgpu_hash_fn: HashFn, + partition_option: PartitionOptions, } impl< @@ -339,10 +356,11 @@ impl< main_trace: &ColMatrix, domain: &StarkDomain, webgpu_hash_fn: HashFn, + partition_option: PartitionOptions, ) -> (Self, TracePolyTable) { // extend the main execution trace and build a Merkle tree from the extended trace let (main_segment_lde, main_segment_tree, main_segment_polys) = - build_trace_commitment_sync::(main_trace, domain); + build_trace_commitment(main_trace, domain, webgpu_hash_fn).await; let trace_poly_table = TracePolyTable::new(main_segment_polys); let trace_lde = WebGPUTraceLde { @@ -353,6 +371,7 @@ impl< blowup: domain.trace_to_lde_blowup(), trace_info: trace_info.clone(), webgpu_hash_fn, + partition_option, }; (trace_lde, trace_poly_table) @@ -393,6 +412,7 @@ impl< > TraceLde for WebGPUTraceLde { type HashFn = H; + type VC = MerkleTree; /// Returns the commitment to the low-degree extension of the main trace segment. fn get_main_trace_commitment(&self) -> D { @@ -419,7 +439,11 @@ impl< ) -> (ColMatrix, D) { // extend the auxiliary trace segment and build a Merkle tree from the extended trace let (aux_segment_lde, aux_segment_tree, aux_segment_polys) = - build_trace_commitment_sync::(aux_trace, domain); + build_trace_commitment_sync::( + aux_trace, + domain, + self.partition_option.partition_size::(aux_trace.num_cols()), + ); assert_eq!( self.main_segment_lde.num_rows(), @@ -549,14 +573,16 @@ impl< /// ``` const DEFAULT_SEGMENT_WIDTH: usize = 8; -fn build_trace_commitment_sync( +fn build_trace_commitment_sync( trace: &ColMatrix, domain: &StarkDomain, -) -> (RowMatrix, MerkleTree, ColMatrix) + partition_size: usize, +) -> (RowMatrix, V, ColMatrix) where E: FieldElement, F: FieldElement, H: ElementHasher, + V: VectorCommitment, { // extend the execution trace let (trace_lde, trace_polys) = { @@ -578,12 +604,12 @@ where assert_eq!(trace_lde.num_rows(), domain.lde_domain_size()); // build trace commitment - let tree_depth = trace_lde.num_rows().ilog2() as usize; - let trace_tree = info_span!("compute_execution_trace_commitment", tree_depth) - .in_scope(|| trace_lde.commit_to_rows()); - assert_eq!(trace_tree.depth(), tree_depth); + let commitment_domain_size = trace_lde.num_rows(); + let trace_vector_com = info_span!("compute_execution_trace_commitment", commitment_domain_size) + .in_scope(|| trace_lde.commit_to_rows::(partition_size)); + assert_eq!(trace_vector_com.domain_len(), commitment_domain_size); - (trace_lde, trace_tree, trace_polys) + (trace_lde, trace_vector_com, trace_polys) } async fn build_trace_commitment< @@ -604,7 +630,7 @@ async fn build_trace_commitment< fft::interpolate_poly(&mut poly, &inv_twiddles); poly }); - let helper = get_or_init_wgpu_helper().await; + let helper = get_or_init_webgpu().await; // extend the execution trace and generate hashes on the gpu let lde_segments = FrozenVec::new(); @@ -613,7 +639,7 @@ async fn build_trace_commitment< let rpo_requires_padding = num_base_columns % RATE != 0; let is_rpo = hash_fn == HashFn::Rpo256; let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); - let mut row_hasher = RowHasher::new(&helper, lde_domain_size, num_base_columns, hash_fn); + let mut row_hasher = RowHasher::::new(&helper, lde_domain_size, num_base_columns, hash_fn); let rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); @@ -644,14 +670,12 @@ async fn build_trace_commitment< row_hasher.update(segment); } let row_hashes = row_hasher.finish().await; - let tree_nodes = helper.build_merkle_tree(&row_hashes, hash_fn).await; + let tree_nodes = helper.build_merkle_tree::(&row_hashes, hash_fn).await; // aggregate segments at the same time as the GPU generates the merkle tree nodes let lde_segments = lde_segments.into_vec().into_iter().map(|p| *p).collect(); let trace_lde = RowMatrix::from_segments(lde_segments, num_base_columns); let trace_polys = lde_segment_generator.into_polys().unwrap(); - let nodes = tree_nodes.into_iter().map(|dig| D::from(&dig)).collect(); - let leaves = row_hashes.into_iter().map(|dig| D::from(&dig)).collect(); - let trace_tree = MerkleTree::from_raw_parts(nodes, leaves).unwrap(); + let trace_tree = MerkleTree::from_raw_parts(tree_nodes, row_hashes).unwrap(); #[cfg(feature = "std")] event!( Level::INFO, @@ -745,25 +769,27 @@ where } } -fn build_segment_queries< - E: FieldElement, - H: Hasher + ElementHasher, ->( +fn build_segment_queries( segment_lde: &RowMatrix, - segment_tree: &MerkleTree, + segment_vector_com: &V, positions: &[usize], -) -> Queries { +) -> Queries +where + E: FieldElement, + H: ElementHasher, + V: VectorCommitment, +{ // for each position, get the corresponding row from the trace segment LDE and put all these // rows into a single vector let trace_states = positions.iter().map(|&pos| segment_lde.row(pos).to_vec()).collect::>(); - // build Merkle authentication paths to the leaves specified by positions - let trace_proof = segment_tree - .prove_batch(positions) - .expect("failed to generate a Merkle proof for trace queries"); + // build a batch opening proof to the leaves specified by positions + let trace_proof = segment_vector_com + .open_many(positions) + .expect("failed to generate a batch opening proof for trace queries"); - Queries::new(trace_proof, trace_states) + Queries::new::(trace_proof.1, trace_states) } struct SegmentIterator<'a, 'b, E, I, const N: usize>(&'b mut SegmentGenerator<'a, E, I, N>) @@ -771,7 +797,7 @@ where E: FieldElement, I: IntoIterator>; -impl<'a, 'b, E, I, const N: usize> Iterator for SegmentIterator<'a, 'b, E, I, N> +impl Iterator for SegmentIterator<'_, '_, E, I, N> where E: FieldElement, I: IntoIterator>, diff --git a/prover/src/lib.rs b/prover/src/lib.rs index c723eee9d..6fcb8fa44 100644 --- a/prover/src/lib.rs +++ b/prover/src/lib.rs @@ -6,7 +6,7 @@ extern crate std; use core::marker::PhantomData; -use air::{AuxRandElements, ProcessorAir, PartitionOptions, PublicInputs}; +use air::{AuxRandElements, PartitionOptions, ProcessorAir, PublicInputs}; #[cfg(all(feature = "metal", target_arch = "aarch64", target_os = "macos"))] use miden_gpu::HashFn; use processor::{ @@ -51,7 +51,6 @@ pub use winter_prover::{crypto::MerkleTree as MerkleTreeVC, Proof}; /// /// # Errors /// Returns an error if program execution or STARK proof generation fails for any reason. -#[maybe_async] #[instrument("prove_program", skip_all)] #[maybe_async] pub fn prove( @@ -111,8 +110,7 @@ where feature = "webgpu", any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") ))] - let prover = - gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpo256); + let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpo256); maybe_await!(prover.prove(trace)) }, HashFunction::Rpx256 => { @@ -127,8 +125,7 @@ where feature = "webgpu", any(all(target_arch = "aarch64", target_os = "macos"), target_family = "wasm") ))] - let prover = - gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpx256); + let prover = gpu::webgpu::WebGPUExecutionProver::new(prover, miden_gpu::HashFn::Rpx256); maybe_await!(prover.prove(trace)) }, } From fe6397bc12fa65729db2d3a2d3c8984a9f2ea6ab Mon Sep 17 00:00:00 2001 From: GopherJ Date: Tue, 19 Nov 2024 15:52:18 +0800 Subject: [PATCH 16/16] migrate to next release of winterfell Signed-off-by: GopherJ --- prover/src/gpu/webgpu/mod.rs | 66 ++++-------------------------------- 1 file changed, 6 insertions(+), 60 deletions(-) diff --git a/prover/src/gpu/webgpu/mod.rs b/prover/src/gpu/webgpu/mod.rs index d74a048b6..4e9e34b25 100644 --- a/prover/src/gpu/webgpu/mod.rs +++ b/prover/src/gpu/webgpu/mod.rs @@ -89,7 +89,7 @@ where } } - fn build_aligned_segement( + fn build_aligned_segment( polys: &ColMatrix, poly_offset: usize, offsets: &[Felt], @@ -119,7 +119,7 @@ where Segment::new_with_buffer(data, polys, poly_offset, offsets, twiddles) } - fn build_aligned_segements( + fn build_aligned_segments( polys: &ColMatrix, twiddles: &[Felt], offsets: &[Felt], @@ -138,7 +138,7 @@ where }; (0..num_segments) - .map(|i| Self::build_aligned_segement(polys, i * N, offsets, twiddles)) + .map(|i| Self::build_aligned_segment(polys, i * N, offsets, twiddles)) .collect() } } @@ -241,7 +241,7 @@ where let blowup = domain.trace_to_lde_blowup(); let offsets = get_evaluation_offsets::(composition_poly.column_len(), blowup, domain.offset()); - let segments = Self::build_aligned_segements( + let segments = Self::build_aligned_segments( composition_poly.data(), domain.trace_twiddles(), &offsets, @@ -262,36 +262,9 @@ where let lde_domain_size = domain.lde_domain_size(); let num_base_columns = composition_poly.num_columns() * ::EXTENSION_DEGREE; - let rpo_requires_padding = num_base_columns % RATE != 0; - let is_rpo = self.webgpu_hash_fn == HashFn::Rpo256; - let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = RowHasher::::new(helper, lde_domain_size, num_base_columns, self.webgpu_hash_fn); - let rpo_padded_segment: Vec<[Felt; RATE]>; - for (segment_idx, segment) in segments.iter().enumerate() { - // check if the segment requires padding - if rpo_padded_segment_idx.map_or(false, |pad_idx| pad_idx == segment_idx) { - // duplicate and modify the last segment with Rpo256's padding - // rule ("1" followed by "0"s). Our segments are already - // padded with "0"s we only need to add the "1"s. - let rpo_pad_column = num_base_columns % RATE; - - rpo_padded_segment = if is_rpo { - segment - .iter() - .map(|x| { - let mut s = *x; - s[rpo_pad_column] = ONE; - s - }) - .collect() - } else { - segment.iter().map(|x| *x).collect() - }; - row_hasher.update(&rpo_padded_segment); - assert_eq!(segments.len() - 1, segment_idx, "padded segment should be the last"); - break; - } + for segment in segments.iter() { row_hasher.update(segment); } let row_hashes = maybe_await!(row_hasher.finish()); @@ -636,37 +609,10 @@ async fn build_trace_commitment< let lde_segments = FrozenVec::new(); let lde_domain_size = domain.lde_domain_size(); let num_base_columns = trace.num_base_cols(); - let rpo_requires_padding = num_base_columns % RATE != 0; - let is_rpo = hash_fn == HashFn::Rpo256; - let rpo_padded_segment_idx = rpo_requires_padding.then_some(num_base_columns / RATE); let mut row_hasher = RowHasher::::new(&helper, lde_domain_size, num_base_columns, hash_fn); - let rpo_padded_segment: Vec<[Felt; RATE]>; let mut lde_segment_generator = SegmentGenerator::new(trace_polys, domain); - let mut lde_segment_iter = lde_segment_generator.gen_segment_iter().enumerate(); - for (segment_idx, segment) in &mut lde_segment_iter { + for segment in lde_segment_generator.gen_segment_iter() { let segment = lde_segments.push_get(Box::new(segment)); - // check if the segment requires padding - if rpo_padded_segment_idx.map_or(false, |pad_idx| pad_idx == segment_idx) { - // duplicate and modify the last segment with Rpo256's padding - // rule ("1" followed by "0"s). Our segments are already - // padded with "0"s we only need to add the "1"s. - let rpo_pad_column = num_base_columns % RATE; - rpo_padded_segment = if is_rpo { - segment - .iter() - .map(|x| { - let mut s = *x; - s[rpo_pad_column] = ONE; - s - }) - .collect() - } else { - segment.iter().map(|x| *x).collect() - }; - row_hasher.update(&rpo_padded_segment); - assert!(lde_segment_iter.next().is_none(), "padded segment should be the last"); - break; - } row_hasher.update(segment); } let row_hashes = row_hasher.finish().await;