diff --git a/Cargo.lock b/Cargo.lock index e15b793..68b4682 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -349,6 +349,15 @@ dependencies = [ "either", ] +[[package]] +name = "itertools" +version = "0.15.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8b4baf93f58d4425749ca49a51c50ebab072c5df6994d08fed93541c331481dc" +dependencies = [ + "either", +] + [[package]] name = "itoa" version = "1.0.18" @@ -377,6 +386,12 @@ version = "0.2.183" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b5b646652bf6661599e1da8901b3b9522896f01e736bad5f723fe7a3a27f899d" +[[package]] +name = "libm" +version = "0.2.16" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b6d2cec3eae94f9f509c767b45932f1ada8350c4bdb85af2fcab4a3c14807981" + [[package]] name = "linux-raw-sys" version = "0.12.1" @@ -442,7 +457,6 @@ dependencies = [ "p3-field", "p3-fri", "p3-goldilocks", - "p3-interpolation", "p3-keccak", "p3-matrix", "p3-maybe-rayon", @@ -467,9 +481,9 @@ dependencies = [ [[package]] name = "num-bigint" -version = "0.4.6" +version = "0.5.1" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "a5e44f723f1133c9deac646763579fdb3ac745e418f2a7af9cd0c431da1f20b9" +checksum = "93e7820bc0a80a0238e650327316f929ba18d5be054b647490a3a6a339f3e7c0" dependencies = [ "num-integer", "num-traits", @@ -507,20 +521,20 @@ checksum = "d6790f58c7ff633d8771f42965289203411a5e5c68388703c06e14f24770b41e" [[package]] name = "p3-air" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-field", "p3-matrix", + "serde", "tracing", ] [[package]] name = "p3-baby-bear" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "p3-challenger", "p3-field", "p3-mds", "p3-monty-31", @@ -532,8 +546,8 @@ dependencies = [ [[package]] name = "p3-blake3" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "blake3", "p3-symmetric", @@ -542,12 +556,11 @@ dependencies = [ [[package]] name = "p3-challenger" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-field", "p3-maybe-rayon", - "p3-monty-31", "p3-symmetric", "p3-util", "tracing", @@ -555,25 +568,24 @@ dependencies = [ [[package]] name = "p3-commit" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", - "p3-challenger", + "itertools 0.15.0", "p3-dft", "p3-field", - "p3-interpolation", "p3-matrix", + "p3-multilinear-util", "p3-util", "serde", ] [[package]] name = "p3-dft" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "p3-field", "p3-matrix", "p3-maybe-rayon", @@ -584,10 +596,10 @@ dependencies = [ [[package]] name = "p3-field" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "num-bigint", "p3-maybe-rayon", "p3-util", @@ -599,17 +611,17 @@ dependencies = [ [[package]] name = "p3-fri" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "p3-challenger", "p3-commit", "p3-dft", "p3-field", - "p3-interpolation", "p3-matrix", "p3-maybe-rayon", + "p3-security", "p3-util", "rand", "serde", @@ -620,11 +632,10 @@ dependencies = [ [[package]] name = "p3-goldilocks" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "num-bigint", - "p3-challenger", "p3-dft", "p3-field", "p3-mds", @@ -635,23 +646,13 @@ dependencies = [ "paste", "rand", "serde", -] - -[[package]] -name = "p3-interpolation" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" -dependencies = [ - "p3-field", - "p3-matrix", - "p3-maybe-rayon", - "p3-util", + "spin", ] [[package]] name = "p3-keccak" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-symmetric", "p3-util", @@ -660,10 +661,10 @@ dependencies = [ [[package]] name = "p3-matrix" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "p3-field", "p3-maybe-rayon", "p3-util", @@ -674,16 +675,16 @@ dependencies = [ [[package]] name = "p3-maybe-rayon" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "rayon", ] [[package]] name = "p3-mds" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-dft", "p3-field", @@ -694,10 +695,10 @@ dependencies = [ [[package]] name = "p3-merkle-tree" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "p3-commit", "p3-field", "p3-matrix", @@ -706,16 +707,17 @@ dependencies = [ "p3-util", "rand", "serde", + "spin", "thiserror", "tracing", ] [[package]] name = "p3-monty-31" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "num-bigint", "p3-dft", "p3-field", @@ -733,20 +735,36 @@ dependencies = [ "tracing", ] +[[package]] +name = "p3-multilinear-util" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" +dependencies = [ + "itertools 0.15.0", + "p3-field", + "p3-matrix", + "p3-maybe-rayon", + "p3-util", + "rand", + "serde", + "tracing", +] + [[package]] name = "p3-poseidon1" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-field", + "p3-mds", "p3-symmetric", "rand", ] [[package]] name = "p3-poseidon2" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ "p3-field", "p3-mds", @@ -755,12 +773,24 @@ dependencies = [ "rand", ] +[[package]] +name = "p3-security" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" +dependencies = [ + "libm", + "p3-air", + "p3-field", + "p3-util", + "serde", +] + [[package]] name = "p3-symmetric" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ - "itertools 0.14.0", + "itertools 0.15.0", "p3-field", "p3-util", "serde", @@ -768,11 +798,11 @@ dependencies = [ [[package]] name = "p3-util" -version = "0.5.1" -source = "git+https://github.com/Plonky3/Plonky3?rev=e9d75614dd6816f9b5dbb4413c69be63536efd64#e9d75614dd6816f9b5dbb4413c69be63536efd64" +version = "0.6.0" +source = "git+https://github.com/Plonky3/Plonky3?rev=3152b14a89067c83775a8076cc262ffc48a1fd7c#3152b14a89067c83775a8076cc262ffc48a1fd7c" dependencies = [ + "p3-maybe-rayon", "serde", - "transpose", ] [[package]] @@ -881,9 +911,9 @@ checksum = "0c8d0fd677905edcbeedbf2edb6494d676f0e98d54d5cf9bda0b061cb8fb8aba" [[package]] name = "rayon" -version = "1.11.0" +version = "1.12.0" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "368f01d005bf8fd9b1206fb6fa653e6c4a81ceb1466406b81792d87c5677a58f" +checksum = "fb39b166781f92d482534ef4b4b1b2568f42613b53e5b6c160e24cfbfa30926d" dependencies = [ "either", "rayon-core", @@ -1043,19 +1073,13 @@ checksum = "8ed6a63f02c8539c91a8685a86f4099661ba3da017932f6ebbea6de3f0fa7c90" [[package]] name = "spin" -version = "0.10.0" +version = "0.12.3" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "d5fe4ccb98d9c292d56fec89a5e07da7fc4cf0dc11e156b41793132775d3e591" +checksum = "0134f9043ed38b087ac4f7d4af44c79e2c9e5094421fe3164f435ce585953b10" dependencies = [ "lock_api", ] -[[package]] -name = "strength_reduce" -version = "0.2.4" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "fe895eb47f22e2ddd4dabc02bce419d2e643c8e3b585c78158b349195bc24d82" - [[package]] name = "syn" version = "2.0.117" @@ -1198,16 +1222,6 @@ dependencies = [ "tracing-subscriber", ] -[[package]] -name = "transpose" -version = "0.2.3" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "1ad61aed86bc3faea4300c7aee358b4c6d0c8d6ccc36524c96e4c92ccf26e77e" -dependencies = [ - "num-integer", - "strength_reduce", -] - [[package]] name = "unicode-ident" version = "1.0.24" diff --git a/Cargo.toml b/Cargo.toml index a4373f3..52315db 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -4,7 +4,7 @@ version = "0.1.0" edition = "2024" authors = ["Argument Engineering "] license = "MIT OR Apache-2.0" -rust-version = "1.91" +rust-version = "1.98" # See more keys and their definitions at https://doc.rust-lang.org/cargo/reference/manifest.html @@ -16,25 +16,24 @@ tracing = "0.1" itertools = { version = "0.14", optional = true } serde = { version = "1", features = ["derive"] } bincode = { version = "2", features = ["serde"] } -p3-air = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-goldilocks = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-challenger = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-commit = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-dft = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-field = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-fri = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-keccak = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-interpolation = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64", optional = true } -p3-blake3 = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-matrix = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-maybe-rayon = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-merkle-tree = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-symmetric = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } -p3-util = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } +p3-air = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-goldilocks = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-challenger = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-commit = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-dft = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-field = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-fri = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-keccak = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-blake3 = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-matrix = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-maybe-rayon = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-merkle-tree = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-symmetric = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } +p3-util = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } [dev-dependencies] criterion = "0.5" -p3-baby-bear = { git = "https://github.com/Plonky3/Plonky3", rev = "e9d75614dd6816f9b5dbb4413c69be63536efd64" } +p3-baby-bear = { git = "https://github.com/Plonky3/Plonky3", rev = "3152b14a89067c83775a8076cc262ffc48a1fd7c" } rand = "0.10" tracing-subscriber = "0.3" tracing-texray = { git = "https://github.com/argumentcomputer/tracing-texray", rev = "465bbca0bea4721e58419c11cabd8cce21757822" } @@ -48,7 +47,7 @@ parallel = ["p3-maybe-rayon/parallel"] # Use the first-party CUDA Goldilocks DFT/LDE backend. Enabling this feature # requires a CUDA toolkit at build time and an NVIDIA GPU at runtime; the # default CPU build never invokes nvcc or links the CUDA runtime. -cuda = ["dep:itertools", "dep:p3-interpolation"] +cuda = ["dep:itertools"] # Similar to `release`, but preserves debug info [profile.dev-ci] @@ -74,7 +73,6 @@ explicit_into_iter_loop = "warn" fallible_impl_from = "warn" filter_map_next = "warn" flat_map_option = "warn" -from_iter_instead_of_collect = "warn" implicit_clone = "warn" inefficient_to_string = "warn" large_stack_arrays = "warn" diff --git a/cuda/kernels.cu b/cuda/kernels.cu index 3c3f13d..a9d6680 100644 --- a/cuda/kernels.cu +++ b/cuda/kernels.cu @@ -1624,6 +1624,19 @@ cudaError_t launch_blake3_digest_pairs(uint8_t* digests, return cudaGetLastError(); } +__global__ void gather_resident_lde_group( + uint64_t* output, const uint64_t* const* columns, const size_t* strides, + size_t row_start, size_t rows, size_t width) { + const size_t count = rows * width; + const size_t grid_stride = static_cast(blockDim.x) * gridDim.x; + for (size_t index = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + index < count; index += grid_stride) { + const size_t column = index % width; + const size_t row = row_start + index / width; + output[index] = columns[column][row * strides[column]]; + } +} + cudaError_t hash_resident_lde_group(uint8_t* digests, const void* const* handles, size_t handle_count, size_t height) { @@ -1642,35 +1655,65 @@ cudaError_t hash_resident_lde_group(uint8_t* digests, return cudaErrorInvalidValue; } - constexpr size_t ROW_STAGING_BYTES = size_t(256) << 20; + constexpr size_t ROW_STAGING_BYTES = size_t(32) << 20; const size_t row_bytes = total_width * sizeof(uint64_t); const size_t rows_per_chunk = (ROW_STAGING_BYTES / row_bytes) > 0 ? (ROW_STAGING_BYTES / row_bytes) : 1; + const uint64_t** host_columns = + new (std::nothrow) const uint64_t*[total_width]; + size_t* host_strides = new (std::nothrow) size_t[total_width]; + if (host_columns == nullptr || host_strides == nullptr) { + delete[] host_columns; + delete[] host_strides; + return cudaErrorMemoryAllocation; + } + size_t column_offset = 0; + for (size_t index = 0; index < handle_count; ++index) { + const ResidentLde* lde = static_cast(handles[index]); + if (lde == nullptr || lde->height != height) { + continue; + } + for (size_t column = 0; column < lde->width; ++column) { + host_columns[column_offset] = lde->values + column; + host_strides[column_offset++] = lde->width; + } + } + DeviceBuffer device_columns; + DeviceBuffer device_strides; + cudaError_t status = device_columns.allocate(total_width); + if (status == cudaSuccess) { + status = device_strides.allocate(total_width); + } + if (status == cudaSuccess) { + status = cudaMemcpy(device_columns.get(), host_columns, + total_width * sizeof(uint64_t*), + cudaMemcpyHostToDevice); + } + if (status == cudaSuccess) { + status = cudaMemcpy(device_strides.get(), host_strides, + total_width * sizeof(size_t), + cudaMemcpyHostToDevice); + } + delete[] host_columns; + delete[] host_strides; DeviceBuffer combined_rows; - cudaError_t status = combined_rows.allocate( - (height < rows_per_chunk ? height : rows_per_chunk) * total_width); + if (status == cudaSuccess) { + status = combined_rows.allocate( + (height < rows_per_chunk ? height : rows_per_chunk) * total_width); + } for (size_t row_start = 0; status == cudaSuccess && row_start < height; row_start += rows_per_chunk) { const size_t rows = (height - row_start < rows_per_chunk) ? height - row_start : rows_per_chunk; - size_t column_offset = 0; - for (size_t index = 0; status == cudaSuccess && index < handle_count; - ++index) { - const ResidentLde* lde = - static_cast(handles[index]); - if (lde == nullptr || lde->height != height) { - continue; - } - status = cudaMemcpy2DAsync( - combined_rows.get() + column_offset, - total_width * sizeof(uint64_t), - lde->values + row_start * lde->width, - lde->width * sizeof(uint64_t), - lde->width * sizeof(uint64_t), rows, cudaMemcpyDeviceToDevice, - cudaStreamPerThread); - column_offset += lde->width; - } + const size_t count = rows * total_width; + gather_resident_lde_group<<>>( + combined_rows.get(), + reinterpret_cast(device_columns.get()), + reinterpret_cast(device_strides.get()), row_start, + rows, total_width); + status = cudaGetLastError(); if (status == cudaSuccess) { status = launch_blake3_rows( digests + row_start * 32, @@ -2453,13 +2496,13 @@ extern "C" int multi_stark_cuda_fri_workspace_destroy(int device_id,void* handle cudaError_t status=cudaSetDevice(device_id);if(status==cudaSuccess)delete static_cast(handle);return static_cast(status); } -extern "C" int multi_stark_cuda_reduced_to_lde(int device_id,void** output,const void* reduced){ +extern "C" int multi_stark_cuda_reduced_into_lde(int device_id,void** output,void* reduced){ if(!output||!reduced)return static_cast(cudaErrorInvalidValue);*output=nullptr; - auto* r=static_cast(reduced);cudaError_t status=cudaSetDevice(device_id); + auto* r=static_cast(reduced);cudaError_t status=cudaSetDevice(device_id); ResidentLde* l=nullptr;if(status==cudaSuccess)status=create_resident_lde(&l);if(status==cudaSuccess){l->height=r->height;l->width=2;} - if(status==cudaSuccess)status=cudaMalloc(reinterpret_cast(&l->values),r->height*sizeof(Ext2)); - if(status==cudaSuccess)status=cudaMemcpy(l->values,r->values,r->height*sizeof(Ext2),cudaMemcpyDeviceToDevice); - if(status!=cudaSuccess){destroy_resident_lde(l);return static_cast(status);}*output=l;return static_cast(cudaSuccess); + if(status!=cudaSuccess){destroy_resident_lde(l);return static_cast(status);} + l->values=reinterpret_cast(r->values);r->values=nullptr;*output=l; + return static_cast(cudaSuccess); } extern "C" int multi_stark_cuda_fri_fold_resident(int device_id,void** output,const void* input, diff --git a/cuda/smoke.sh b/cuda/smoke.sh index 4089dac..0ad061a 100755 --- a/cuda/smoke.sh +++ b/cuda/smoke.sh @@ -27,7 +27,7 @@ cargo test --release --locked --features parallel,cuda -- --test-threads=1 compat_dir="$(mktemp -d)" trap 'rm -rf "$compat_dir"' EXIT cargo run --release --locked --example proof_compatibility -- "$compat_dir/cpu.proof" -echo "c3c8ff942efc36fe508d6dcea3450cf4fa4d92a84f8d16bc651430595f4d244a $compat_dir/cpu.proof" \ +echo "25564a01d1d352b1ec2de56b019b641d24acc81083133e79274a86829b2a5dd5 $compat_dir/cpu.proof" \ | sha256sum --check --status cargo run --release --locked --features parallel,cuda \ --example proof_compatibility -- "$compat_dir/cuda.proof" diff --git a/examples/proof_compatibility.rs b/examples/proof_compatibility.rs index eea52cd..a8a33b6 100644 --- a/examples/proof_compatibility.rs +++ b/examples/proof_compatibility.rs @@ -66,7 +66,7 @@ fn main() { let (system, key) = System::new(config, [LookupAir::new(WidePythagoreanAir, vec![])]); let height = 1 << LOG_HEIGHT; let mut values = Val::zero_vec(height * WIDTH); - for row in values.chunks_exact_mut(WIDTH) { + for row in values.as_chunks_mut::().0 { row[0] = Val::from_u8(3); row[1] = Val::from_u8(4); row[2] = Val::from_u8(5); diff --git a/flake.nix b/flake.nix index 8ed206b..b4848e1 100644 --- a/flake.nix +++ b/flake.nix @@ -40,7 +40,7 @@ # Pins the Rust toolchain rustToolchain = fenix.packages.${system}.fromToolchainFile { file = ./rust-toolchain.toml; - sha256 = "sha256-SDu4snEWjuZU475PERvu+iO50Mi39KVjqCeJeNvpguU="; + sha256 = "sha256-P30Tm3O7vQAE725YtDCDHGjNrSsfZO4us11UwJGZSJo="; }; craneLib = (crane.mkLib pkgs).overrideToolchain rustToolchain; diff --git a/rust-toolchain.toml b/rust-toolchain.toml index dccca02..23de054 100644 --- a/rust-toolchain.toml +++ b/rust-toolchain.toml @@ -1,4 +1,4 @@ [toolchain] # The default profile includes rustc, rust-std, cargo, rust-docs, rustfmt and clippy. profile = "default" -channel = "1.91" +channel = "1.98" diff --git a/src/cuda/mmcs.rs b/src/cuda/mmcs.rs index df45dbe..8f19966 100644 --- a/src/cuda/mmcs.rs +++ b/src/cuda/mmcs.rs @@ -93,41 +93,6 @@ pub trait CudaCommitMmcs: Mmcs { ) -> (Self::Commitment, Self::ProverData>); } -pub trait CudaBatchOpenMmcs: Mmcs { - fn open_batches>( - &self, - indices: &[usize], - prover_data: &Self::ProverData, - ) -> Vec> - where - Self: Sized; -} - -impl CudaBatchOpenMmcs for CudaMmcs { - fn open_batches>( - &self, - indices: &[usize], - prover_data: &Self::ProverData, - ) -> Vec> { - match prover_data { - CudaMmcsData::Cpu(_) => indices - .iter() - .map(|&index| self.open_batch(index, prover_data)) - .collect(), - CudaMmcsData::Cuda { resident, tree, .. } => { - let rows = mixed_lde_open_rows(resident, indices); - let paths = tree.open_siblings_batch(indices); - rows.into_iter() - .zip(paths) - .map(|(opened_values, opening_proof)| { - BatchOpening::new(opened_values, opening_proof) - }) - .collect() - } - } - } -} - impl CudaCommitMmcs for CudaMmcs { fn cuda_device_id(&self) -> i32 { self.device_id @@ -218,6 +183,7 @@ impl Mmcs for CudaMmcs { type ProverData = CudaMmcsData; type Commitment = MerkleCap; type Proof = Vec<[u8; 32]>; + type MultiProof = >::MultiProof; type Error = MerkleTreeError; fn commit>( @@ -332,4 +298,35 @@ impl Mmcs for CudaMmcs { BatchOpeningRef::new(batch_opening.opened_values, batch_opening.opening_proof), ) } + + fn open_multi_batch>( + &self, + indices: &[usize], + prover_data: &Self::ProverData, + ) -> (Vec>>, Self::MultiProof) { + match prover_data { + CudaMmcsData::Cpu(data) => self.cpu.open_multi_batch(indices, data), + CudaMmcsData::Cuda { resident, tree, .. } => { + let opened_values = if indices.is_empty() { + Vec::new() + } else { + mixed_lde_open_rows(resident, indices) + }; + let opening_proof = tree.open_pruned_siblings(indices); + (opened_values, opening_proof) + } + } + } + + fn verify_multi_batch + PartialEq>( + &self, + commit: &Self::Commitment, + dimensions: &[Dimensions], + indices: &[usize], + opened_values: &[Vec], + proof: &Self::MultiProof, + ) -> Result<(), Self::Error> { + self.cpu + .verify_multi_batch(commit, dimensions, indices, opened_values, proof) + } } diff --git a/src/cuda/mod.rs b/src/cuda/mod.rs index d7c24a2..017a365 100644 --- a/src/cuda/mod.rs +++ b/src/cuda/mod.rs @@ -23,6 +23,7 @@ use p3_goldilocks::Goldilocks; use p3_matrix::Matrix; use p3_matrix::bitrev::{BitReversalPerm, BitReversedMatrixView}; use p3_matrix::dense::RowMajorMatrix; +use p3_merkle_tree::PrunedMerklePaths; use p3_util::log2_strict_usize; const _: () = assert!(size_of::() == size_of::()); @@ -1104,10 +1105,10 @@ impl CudaReducedOpening { check_cuda(status, "reduced opening copy"); out } - pub(crate) fn to_lde(&self) -> CudaLde { + pub(crate) fn into_lde(self) -> CudaLde { let mut handle = core::ptr::null_mut(); let status = unsafe { - multi_stark_cuda_reduced_to_lde(self.device_id, &mut handle, self.handle.as_ptr()) + multi_stark_cuda_reduced_into_lde(self.device_id, &mut handle, self.handle.as_ptr()) }; check_cuda(status, "reduced opening to resident FRI codeword"); CudaLde { @@ -1902,6 +1903,9 @@ impl CudaMixedMerkleTree { "mixed Merkle opening index out of bounds" ); let levels = self.row_count.trailing_zeros() as usize; + if levels == 0 { + return vec![Vec::new(); indices.len()]; + } let device_indices: Vec = indices .iter() .map(|&index| u64::try_from(index).expect("Merkle index exceeds u64")) @@ -1922,6 +1926,57 @@ impl CudaMixedMerkleTree { .map(<[[u8; 32]]>::to_vec) .collect() } + + /// Opens several leaves and prunes their overlapping binary authentication paths. + /// + /// The CUDA kernel returns one full path per requested index. Plonky3's multiproof + /// format keeps only boundary digests, ordered by tree level and then parent index. + #[must_use] + pub(crate) fn open_pruned_siblings(&self, indices: &[usize]) -> PrunedMerklePaths { + if indices.is_empty() { + return PrunedMerklePaths { + sibling_hashes: Vec::new(), + }; + } + + let paths = self.open_siblings_batch(indices); + // Each frontier entry is (node index at this level, source query slot). + // Sorting and deduplicating matches Plonky3's treatment of repeated queries. + let mut frontier: Vec<_> = indices + .iter() + .copied() + .enumerate() + .map(|(slot, index)| (index, slot)) + .collect(); + frontier.sort_unstable_by_key(|&(index, _)| index); + frontier.dedup_by_key(|entry| entry.0); + + let mut parents = Vec::with_capacity(frontier.len()); + let mut sibling_hashes = Vec::new(); + for (level, _) in paths[0].iter().enumerate() { + parents.clear(); + let mut at = 0; + while at < frontier.len() { + let parent = frontier[at].0 >> 1; + let lead_slot = frontier[at].1; + let group_start = at; + at += 1; + while at < frontier.len() && frontier[at].0 >> 1 == parent { + at += 1; + } + + let group_len = at - group_start; + debug_assert!(group_len <= 2, "binary frontier group exceeds two children"); + if group_len == 1 { + sibling_hashes.push(paths[lead_slot][level]); + } + parents.push((parent, lead_slot)); + } + core::mem::swap(&mut frontier, &mut parents); + } + + PrunedMerklePaths { sibling_hashes } + } } impl Drop for CudaMixedMerkleTree { @@ -2167,10 +2222,10 @@ unsafe extern "C" { ext_w: u64, ) -> i32; fn multi_stark_cuda_fri_workspace_destroy(device_id: i32, handle: *mut c_void) -> i32; - fn multi_stark_cuda_reduced_to_lde( + fn multi_stark_cuda_reduced_into_lde( device_id: i32, output: *mut *mut c_void, - reduced: *const c_void, + reduced: *mut c_void, ) -> i32; fn multi_stark_cuda_fri_fold_resident( device_id: i32, @@ -2489,7 +2544,7 @@ mod tests { use p3_field::{ BasedVectorSpace, batch_multiplicative_inverse, extension::BinomialExtensionField, }; - use p3_interpolation::interpolate_coset_with_precomputation; + use p3_matrix::interpolation::{Interpolate, compute_adjusted_weights}; use p3_util::reverse_slice_index_bits; type Ext = BinomialExtensionField; let height = 256; @@ -2516,13 +2571,11 @@ mod tests { .unwrap(); let inv = batch_multiplicative_inverse(&coset.iter().map(|&x| point - x).collect::>()); - let expected = interpolate_coset_with_precomputation( - &bitrev.split_rows(height).0, - Goldilocks::GENERATOR, - point, - &coset[..height], - &inv[..height], - ); + let adjusted = compute_adjusted_weights(point, &inv[..height]); + let expected = bitrev + .split_rows(height) + .0 + .interpolate_coset_with_precomputation(Goldilocks::GENERATOR, point, &adjusted); let inv2: Vec<[Goldilocks; 2]> = inv[..height] .iter() .map(|x| x.as_basis_coefficients_slice().try_into().unwrap()) @@ -2701,7 +2754,9 @@ mod tests { .collect(); while layer.len() > 1 { layer = layer - .chunks_exact(2) + .as_chunks::<2>() + .0 + .iter() .map(|children| { Blake3.hash_iter(children[0].iter().chain(&children[1]).copied()) }) @@ -2748,13 +2803,17 @@ mod tests { ]; let mut layer: Vec<[u8; 32]> = level_8 - .chunks_exact(16) + .as_chunks::<16>() + .0 + .iter() .map(|row| Blake3.hash_iter(row.iter().copied())) .collect(); let mut digest_layers = vec![layer.clone()]; while layer.len() > 1 { layer = layer - .chunks_exact(2) + .as_chunks::<2>() + .0 + .iter() .map(|children| Blake3.hash_iter(children.iter().flatten().copied())) .collect(); let injected = match layer.len() { diff --git a/src/cuda/pcs.rs b/src/cuda/pcs.rs index 99f3c90..73a5265 100644 --- a/src/cuda/pcs.rs +++ b/src/cuda/pcs.rs @@ -14,11 +14,10 @@ //! If we changed our domain construction (e.g., using multiple cosets), we would need to carefully reconsider these assumptions. //! //! The CPU PCS control flow in this module is derived from Plonky3's -//! `two_adic_pcs.rs` at revision e9d75614dd6816f9b5dbb4413c69be63536efd64 +//! `two_adic_pcs.rs` at revision 3152b14a89067c83775a8076cc262ffc48a1fd7c //! (MIT/Apache-2.0). CUDA-resident commitments, openings, and FRI are maintained //! here so their transcript order can be reviewed directly against that source. -use core::fmt::Debug; use core::iter; use core::marker::PhantomData; use core::mem::size_of; @@ -28,26 +27,23 @@ use std::vec::Vec; use itertools::{Itertools, izip}; use p3_challenger::{CanObserve, FieldChallenger, GrindingChallenger}; -use p3_commit::{ - BatchOpening, BuildPeriodicLdeTableFast, ExtensionMmcs, Mmcs, OpenedValues, Pcs, - PeriodicLdeTable, -}; +use p3_commit::{ExtensionMmcs, Mmcs, OpenedValues, Pcs, PeriodicLdeTable}; use p3_dft::{Radix2DFTSmallBatch, TwoAdicSubgroupDft}; use p3_field::coset::TwoAdicMultiplicativeCoset; use p3_field::{ BasedVectorSpace, ExtensionField, PackedFieldExtension, PrimeCharacteristicRing, PrimeField64, TwoAdicField, batch_multiplicative_inverse, dot_product, }; -use p3_interpolation::interpolate_coset_with_precomputation; use p3_matrix::Matrix; use p3_matrix::bitrev::{BitReversedMatrixView, BitReversibleMatrix}; use p3_matrix::dense::{RowMajorMatrix, RowMajorMatrixCow}; +use p3_matrix::interpolation::{Interpolate, compute_adjusted_weights}; use p3_maybe_rayon::prelude::*; use p3_util::linear_map::LinearMap; -use p3_util::{log2_strict_usize, reverse_bits_len, reverse_slice_index_bits}; +use p3_util::{log2_strict_usize, reverse_slice_index_bits}; use tracing::{debug_span, instrument}; -use super::mmcs::{CudaBatchOpenMmcs, CudaCommitMmcs}; +use super::mmcs::CudaCommitMmcs; use super::{CudaFriWorkspace, CudaLde, CudaMixedMerkleTree, CudaReducedOpening}; use p3_goldilocks::Goldilocks; use p3_symmetric::MerkleCap; @@ -64,8 +60,9 @@ pub trait CudaPcsDft: TwoAdicSubgroupDft { } use p3_fri::{ - CommitPhaseProofStep, FriFoldingStrategy, FriParameters, FriProof, QueryProof, - build_periodic_lde_table_two_adic, compute_log_arity_for_round, prover, + BatchMultiOpening, CommitPhaseMultiStep, FriParameters, FriProof, TwoAdicFriFolding, + TwoAdicFriFoldingForMmcs, build_periodic_lde_table_two_adic, compute_log_arity_for_round, + prover, verifier::{self, FriError}, }; @@ -85,7 +82,7 @@ trait CudaFriMmcs: Mmcs { &self, round: &CudaFriRound, rows: &[usize], - ) -> Vec<(Vec, Self::Proof)>; + ) -> (Vec>, Self::MultiProof); } impl CudaFriMmcs @@ -115,27 +112,26 @@ where &self, round: &CudaFriRound, rows: &[usize], - ) -> Vec<(Vec, Self::Proof)> { + ) -> (Vec>, Self::MultiProof) { let codeword_rows = rows .iter() .flat_map(|&row| (0..round.arity).map(move |column| row * round.arity + column)) .collect_vec(); let opened = round.codeword.rows(&codeword_rows); - let paths = round.tree.open_siblings_batch(rows); - opened + let opened_rows = opened .chunks_exact(round.arity) - .zip(paths) - .map(|(query_rows, opening_proof)| { - let values = query_rows + .map(|query_rows| { + query_rows .iter() .map(|pair| { Challenge::from_basis_coefficients_slice(pair) .expect("quadratic extension row") }) - .collect(); - (values, opening_proof) + .collect() }) - .collect() + .collect(); + let opening_proof = round.tree.open_pruned_siblings(rows); + (opened_rows, opening_proof) } } @@ -153,40 +149,6 @@ pub struct CudaTwoAdicFriPcs { _phantom: PhantomData, } -struct CudaFriFolding(PhantomData<(InputProof, InputError)>); -type CudaFriFoldingForMmcs = CudaFriFolding>, >::Error>; -impl FriFoldingStrategy - for CudaFriFolding -where - Val: TwoAdicField + PrimeField64, - Challenge: ExtensionField, -{ - type InputProof = InputProof; - type InputError = InputError; - fn extra_query_index_bits(&self) -> usize { - 0 - } - fn fold_row( - &self, - index: usize, - log_height: usize, - log_arity: usize, - beta: Challenge, - evals: impl Iterator, - ) -> Challenge { - TwoAdicFriFolding::(PhantomData) - .fold_row(index, log_height, log_arity, beta, evals) - } - fn fold_matrix>( - &self, - beta: Challenge, - log_arity: usize, - m: M, - ) -> Vec { - TwoAdicFriFolding::(PhantomData).fold_matrix(beta, log_arity, m) - } -} - fn prove_fri_cuda_resident( params: &FriParameters, mut inputs: Vec, @@ -199,20 +161,38 @@ fn prove_fri_cuda_resident( >], input_mmcs: &InputMmcs, ext_w: Goldilocks, -) -> FriProof>> +) -> FriProof>> where Val: TwoAdicField + PrimeField64, Challenge: ExtensionField, - InputMmcs: Mmcs + CudaBatchOpenMmcs, + InputMmcs: Mmcs, FriMmcs: CudaFriMmcs, Challenger: FieldChallenger + GrindingChallenger + CanObserve, { assert!(!inputs.is_empty()); + assert!( + params.num_queries > 0, + "num_queries must be at least 1 for FRI soundness" + ); + assert!( + params.max_log_arity > 0, + "max_log_arity must be at least 1 to guarantee folding progress" + ); assert!( inputs .windows(2) - .all(|pair| pair[0].height() > pair[1].height()) + .all(|pair| pair[0].height() > pair[1].height()), + "inputs are not sorted in strictly descending order of height" ); + assert_eq!( + log_global_max_height, + log2_strict_usize(inputs[0].height()), + "log_global_max_height must match the largest input height" + ); + let log_min_height = log2_strict_usize(inputs.last().unwrap().height()); + if params.log_final_poly_len > 0 { + assert!(log_min_height > params.log_final_poly_len + params.log_blowup); + } let to_pair = |value: Challenge| { let coefficients = value.as_basis_coefficients_slice(); [ @@ -228,7 +208,7 @@ where .expect("quadratic extension element") }; - let mut codeword = inputs.remove(0).to_lde(); + let mut codeword = inputs.remove(0).into_lde(); let mut commits = Vec::new(); let mut rounds = Vec::new(); let mut log_arities = Vec::new(); @@ -286,10 +266,12 @@ where } assert!(inputs.is_empty()); - let mut final_values = codeword - .to_row_major_matrix() + let final_matrix = codeword.to_row_major_matrix(); + let mut final_values = final_matrix .values - .chunks_exact(2) + .as_chunks::<2>() + .0 + .iter() .take(params.final_poly_len()) .map(|pair| from_pair([pair[0], pair[1]])) .collect_vec(); @@ -301,12 +283,10 @@ where challenger.observe(Val::from_usize(log_arity)); } let query_pow_witness = challenger.grind(params.query_proof_of_work_bits); - let query_indices = iter::repeat_with(|| { - challenger.sample_bits(log2_strict_usize(rounds[0].codeword.height())) - }) - .take(params.num_queries) - .collect_vec(); - let input_batches = prover_data_with_opening_points + let query_indices = iter::repeat_with(|| challenger.sample_bits(log_global_max_height)) + .take(params.num_queries) + .collect_vec(); + let input_openings = prover_data_with_opening_points .iter() .map(|(data, _)| { let log_height = log2_strict_usize(input_mmcs.get_max_height(data)); @@ -314,11 +294,15 @@ where .iter() .map(|&index| index >> (log_global_max_height - log_height)) .collect_vec(); - input_mmcs.open_batches(&indices, data) + let (opened_values, opening_proof) = input_mmcs.open_multi_batch(&indices, data); + BatchMultiOpening { + opened_values, + opening_proof, + } }) .collect_vec(); let mut current_indices = query_indices; - let commit_batches = rounds + let commit_phase_openings = rounds .iter() .zip(&log_arities) .map(|(round, &log_arity)| { @@ -331,45 +315,34 @@ where .iter() .map(|&index| index >> log_arity) .collect_vec(); - let openings = params.mmcs.open_cuda_fri_batch(round, &group_indices); + let (opened_rows, opening_proof) = + params.mmcs.open_cuda_fri_batch(round, &group_indices); current_indices = group_indices; - positions + let sibling_values = positions .into_iter() - .zip(openings) - .map(|(index_in_group, (opened, opening_proof))| { - let sibling_values = opened + .zip(opened_rows) + .map(|(index_in_group, opened)| { + assert_eq!(opened.len(), arity, "FRI opening has the wrong arity"); + opened .into_iter() .enumerate() .filter_map(|(column, value)| (column != index_in_group).then_some(value)) - .collect(); - CommitPhaseProofStep { - log_arity: u8::try_from(log_arity).expect("FRI arity exceeds u8"), - sibling_values, - opening_proof, - } + .collect() }) - .collect_vec() + .collect_vec(); + CommitPhaseMultiStep { + log_arity: u8::try_from(log_arity).expect("FRI arity exceeds u8"), + sibling_values, + opening_proof, + } }) .collect_vec(); - let mut input_iters = input_batches.into_iter().map(Vec::into_iter).collect_vec(); - let mut commit_iters = commit_batches.into_iter().map(Vec::into_iter).collect_vec(); - let query_proofs = (0..params.num_queries) - .map(|_| QueryProof { - input_proof: input_iters - .iter_mut() - .map(|openings| openings.next().expect("complete input opening batch")) - .collect(), - commit_phase_openings: commit_iters - .iter_mut() - .map(|openings| openings.next().expect("complete FRI opening batch")) - .collect(), - }) - .collect(); FriProof { commit_phase_commits: commits, commit_pow_witnesses, - query_proofs, + input_openings, + commit_phase_openings, final_poly, query_pow_witness, } @@ -411,180 +384,12 @@ pub type CommitmentWithOpeningPoints = ( )>, ); -pub struct TwoAdicFriFolding(pub PhantomData<(InputProof, InputError)>); - -pub type TwoAdicFriFoldingForMmcs = - TwoAdicFriFolding>, >::Error>; - -impl> - FriFoldingStrategy for TwoAdicFriFolding -{ - type InputProof = InputProof; - type InputError = InputError; - - fn extra_query_index_bits(&self) -> usize { - 0 - } - - fn fold_row( - &self, - index: usize, - log_height: usize, - log_arity: usize, - beta: EF, - evals: impl Iterator, - ) -> EF { - let arity = 1 << log_arity; - let evals: Vec<_> = evals.collect(); - assert_eq!(evals.len(), arity, "Expected {} evaluations", arity); - - // Compute the evaluation points in the subgroup - let subgroup_start = F::two_adic_generator(log_height + log_arity) - .exp_u64(reverse_bits_len(index, log_height) as u64); - let mut xs: Vec = F::two_adic_generator(log_arity) - .shifted_powers(subgroup_start) - .take(arity) - .collect(); - reverse_slice_index_bits(&mut xs); - - // Lagrange interpolation at beta - lagrange_interpolate_at(&xs, &evals, beta) - } - - #[instrument(skip_all)] - fn fold_matrix>(&self, beta: EF, log_arity: usize, m: M) -> Vec { - if log_arity == 1 { - // Optimized path for arity 2 - // We use the fact that - // p_e(x^2) = (p(x) + p(-x)) / 2 - // p_o(x^2) = (p(x) - p(-x)) / (2 x) - // that is, - // p_e(g^(2i)) = (p(g^i) + p(g^(n/2 + i))) / 2 - // p_o(g^(2i)) = (p(g^i) - p(g^(n/2 + i))) / (2 g^i) - // so - // result(g^(2i)) = p_e(g^(2i)) + beta p_o(g^(2i)) - // - // As p_e, p_o will be in the extension field we want to find ways to avoid extension multiplications. - // We should only need a single one (namely multiplication by beta). - let g_inv = F::two_adic_generator(log2_strict_usize(m.height()) + 1).inverse(); - - // As beta is in the extension field, we want to avoid multiplying by it - // for as long as possible. Here we precompute the powers `g_inv^i / 2` in the base field. - let mut halve_inv_powers = g_inv.shifted_powers(F::ONE.halve()).collect_n(m.height()); - reverse_slice_index_bits(&mut halve_inv_powers); - - m.par_rows() - .zip(halve_inv_powers) - .map(|(mut row, halve_inv_power)| { - let (lo, hi) = row.next_tuple().unwrap(); - (lo + hi).halve() + (lo - hi) * beta * halve_inv_power - }) - .collect() - } else { - // Decompose arity-2^k fold into k sequential arity-2 folds. - // This way, an arity-2^k fold with a single challenge beta is equivalent to - // k arity-2 folds with challenges beta, beta^2, beta^4, ..., beta^{2^{k-1}}. - // - // For arity 4 with evaluation points {s, -s, si, -si}: - // Step 1 (beta): fold pairs → g(s^2), g(-s^2) where g = f_e + beta*f_o - // Step 2 (beta^2): fold pair → g(beta^2) = f(beta) - - let mut data = m.to_row_major_matrix().values; - - let initial_height = data.len() / 2; - let g_inv = F::two_adic_generator(log2_strict_usize(initial_height) + 1).inverse(); - let mut halve_inv_powers = g_inv - .shifted_powers(F::ONE.halve()) - .collect_n(initial_height); - reverse_slice_index_bits(&mut halve_inv_powers); - - let two = F::ONE + F::ONE; - let mut current_beta = beta; - let mut next_data = EF::zero_vec(initial_height); - - for step in 0..log_arity { - let current_len = data.len(); - let height = current_len / 2; - // Since j << 1 is always >= j, we never overwrite data we haven't read yet. - if step > 0 { - for j in 0..height { - halve_inv_powers[j] = two * halve_inv_powers[j << 1].square(); - } - } - next_data[..height] - .par_iter_mut() - .zip(data.par_chunks_exact(2)) - .zip(&halve_inv_powers[..height]) - .for_each(|((out, chunk), &halve_inv_power)| { - // chunk is guaranteed to be size 2 by par_chunks_exact - let lo = chunk[0]; - let hi = chunk[1]; - - *out = (lo + hi).halve() + (lo - hi) * current_beta * halve_inv_power; - }); - current_beta = current_beta.square(); - - // Swap buffers conceptually (just truncate data and copy back, or ping-pong). - data.truncate(height); - data.copy_from_slice(&next_data[..height]); - } - - data - } - } -} - -/// Lagrange interpolation: given points (xs[i], ys[i]), evaluate at z. -/// -/// Uses the barycentric formula for efficiency when xs are roots of unity. -fn lagrange_interpolate_at>( - xs: &[F], - ys: &[EF], - z: EF, -) -> EF { - debug_assert_eq!(xs.len(), ys.len()); - let n = xs.len(); - - if n == 0 { - return EF::ZERO; - } - - // If z equals one of the interpolation points, return early. - for i in 0..n { - if (z - xs[i]).is_zero() { - return ys[i]; - } - } - - let log_n = log2_strict_usize(n); - - // All xs lie in a coset of the 2^log_n roots of unity. - let coset_power = xs[0].exp_power_of_2(log_n); - let weight_scale = (F::from_usize(n) * coset_power).inverse(); - - // Compute (z - x_i)^{-1} as a batch inversion - let diffs: Vec<_> = xs.iter().map(|&x| z - x).collect(); - let diff_invs = batch_multiplicative_inverse(&diffs); - - // Compute L(z) = prod_i (z - x_i) - let l_z = diffs.iter().copied().product::(); - - // Barycentric formula: sum_i (w_i * y_i / (z - x_i)) - // where w_i = 1 / prod_{j != i} (x_i - x_j) = x_i * weight_scale. - let mut result = EF::ZERO; - for ((&x, &y), &diff_inv) in xs.iter().zip(ys).zip(diff_invs.iter()) { - let weight = x * weight_scale; - result += y * weight * diff_inv; - } - result * l_z -} - impl Pcs for CudaTwoAdicFriPcs where Val: TwoAdicField + PrimeField64, Dft: TwoAdicSubgroupDft + CudaPcsDft + Sync, - InputMmcs: Mmcs + CudaCommitMmcs + CudaBatchOpenMmcs, + InputMmcs: Mmcs + CudaCommitMmcs, FriMmcs: Mmcs + CudaFriMmcs, Challenge: ExtensionField, Challenger: @@ -594,7 +399,7 @@ where type Commitment = InputMmcs::Commitment; type ProverData = InputMmcs::ProverData>; type EvaluationsOnDomain<'a> = BitReversedMatrixView>; - type Proof = FriProof>>; + type Proof = FriProof>>; type Error = FriError; const ZK: bool = false; @@ -606,6 +411,10 @@ where TwoAdicMultiplicativeCoset::new(Val::ONE, log2_strict_usize(degree)).unwrap() } + fn log_max_lde_height(&self) -> usize { + Val::TWO_ADICITY + } + /// Commit to a collection of evaluation matrices. /// /// Each element of `evaluations` contains a coset `shift * H` and a matrix `mat` with `mat.height() = |H|`. @@ -737,6 +546,14 @@ where } fn commit_ldes(&self, ldes: Vec>) -> (Self::Commitment, Self::ProverData) { + let min_height = 1 << self.fri.log_blowup; + for lde in &ldes { + assert!( + lde.height() >= min_height, + "committed LDE height {} is smaller than the blowup factor {min_height}", + lde.height() + ); + } self.mmcs.commit_cuda_storage(ldes) } @@ -774,6 +591,7 @@ where let result = self .dft .coset_dft_batch(coeffs, domain.shift()) + .bit_reverse_rows() .to_row_major_matrix(); let result_width = result.width(); @@ -982,8 +800,8 @@ where let alpha: Challenge = challenger.sample_algebra_element(); let alpha_powers: Vec<_> = alpha.powers().take(global_max_width).collect(); let alpha_pairs: Vec<_> = alpha_powers.iter().copied().map(to_pair).collect(); - let mut num_reduced = [0usize; 32]; - let mut reduced: [Option; 32] = core::array::from_fn(|_| None); + let mut num_reduced = [0usize; 33]; + let mut reduced: [Option; 33] = core::array::from_fn(|_| None); let mut reduction_tasks = Vec::new(); for ((ldes, points), openings_round) in rounds.iter().zip(all_opened_values.iter()) { for ((lde, ps), openings) in @@ -1067,6 +885,13 @@ where // for that point, and precompute 1/(z - X) for the largest subgroup (in bitrev order). let inv_denoms = compute_inverse_denominators(&mats_and_points, &coset); + // Convert the inverse denominators into the adjusted barycentric weights expected by + // the matrix interpolation API. Reuse them across every matrix opened at each point. + let adjusted_weights: LinearMap> = inv_denoms + .iter() + .map(|(point, denoms)| (*point, compute_adjusted_weights(*point, denoms))) + .collect(); + // Evaluate coset representations and write openings to the challenger let all_opened_values = mats_and_points .iter() @@ -1086,7 +911,6 @@ where // `subgroup` and `mat` are both in bit-reversed order, so we can truncate. let (low_coset, _) = mat.split_rows(h); - let coset_h = &coset[..h]; points_for_mat .iter() @@ -1100,16 +924,11 @@ where "compute opened values with Lagrange interpolation" ) .in_scope(|| { - // Get the relevant inverse denominators for this point and use these to - // interpolate to get the evaluation of each polynomial in the matrix - // at the desired point. - let inv_denoms = &inv_denoms.get(&point).unwrap()[..h]; - interpolate_coset_with_precomputation( - &low_coset, + let adjusted = &adjusted_weights.get(&point).unwrap()[..h]; + low_coset.interpolate_coset_with_precomputation( Val::GENERATOR, point, - coset_h, - inv_denoms, + adjusted, ) }); @@ -1157,13 +976,13 @@ where // num_reduced records the number of (function, opening point) pairs for each `log_height`. // TODO: This should really be `[0; Val::TWO_ADICITY]` but that runs into issues with generics. - let mut num_reduced = [0; 32]; + let mut num_reduced = [0; 33]; // For each `log_height` from 2^1 -> 2^32, reduced_openings will contain either `None` // if there are no matrices of that height, or `Some(vec)` where `vec` is equal to // a weighted sum of `(f(zeta) - f(x))/(zeta - x)` over all `f`'s of that height and // for each `f`, all opening points `zeta`. The sum is weighted by powers of the challenge alpha. - let mut reduced_openings: [_; 32] = core::array::from_fn(|_| None); + let mut reduced_openings: [_; 33] = core::array::from_fn(|_| None); for ((mats, points), openings_for_round) in mats_and_points.iter().zip(all_opened_values.iter()) @@ -1227,7 +1046,7 @@ where // low degree functions. let fri_input = reduced_openings.into_iter().rev().flatten().collect_vec(); - let folding: CudaFriFoldingForMmcs = CudaFriFolding(PhantomData); + let folding: TwoAdicFriFoldingForMmcs = TwoAdicFriFolding(PhantomData); // Produce the FRI proof. let fri_proof = prover::prove_fri( @@ -1275,32 +1094,19 @@ where Ok(()) } -} - -impl BuildPeriodicLdeTableFast - for CudaTwoAdicFriPcs -where - Val: TwoAdicField, - Dft: TwoAdicSubgroupDft + CudaPcsDft + Default, -{ - type PeriodicDomain = TwoAdicMultiplicativeCoset; - fn maybe_build_periodic_lde_table_fast( + fn build_periodic_lde_table( &self, - periodic_cols: &[Vec>], - trace_domain: Self::PeriodicDomain, - quotient_domain: Self::PeriodicDomain, - ) -> Option>> - where - p3_commit::Val: Clone, - { - let periodic_cols_val: &[Vec] = unsafe { core::mem::transmute(periodic_cols) }; - let table = build_periodic_lde_table_two_adic::( - periodic_cols_val, + periodic_cols: &[Vec], + trace_domain: Self::Domain, + quotient_domain: Self::Domain, + ) -> PeriodicLdeTable { + build_periodic_lde_table_two_adic::( + &self.dft, + periodic_cols, &trace_domain, "ient_domain, - ); - Some(table) + ) } } diff --git a/src/p3_adapter.rs b/src/p3_adapter.rs index f29c348..686c798 100644 --- a/src/p3_adapter.rs +++ b/src/p3_adapter.rs @@ -258,6 +258,9 @@ impl AirBuilder for P3AirBuilder { // Public inputs are reserved for the lookup argument in this system, so // the builder exposes none (the default `public_values()` is empty). type PublicVar = P3Var; + // Periodic columns are rejected by `circuit_inputs_from_air`, but the + // Plonky3 builder trait still requires a concrete variable type. + type PeriodicVar = P3Var; fn main(&self) -> Self::MainWindow { self.main.clone() @@ -275,6 +278,10 @@ impl AirBuilder for P3AirBuilder { P3Expr(Expr::IsLastRow) } + fn is_transition(&self) -> Self::Expr { + P3Expr(Expr::IsTransition) + } + /// # Panics /// Panics if `size` is not `2`; only two-row windows are supported. fn is_transition_window(&self, size: usize) -> Self::Expr { diff --git a/src/test_circuits/blake3.rs b/src/test_circuits/blake3.rs index 2cdcd52..3b5bdd1 100644 --- a/src/test_circuits/blake3.rs +++ b/src/test_circuits/blake3.rs @@ -117,8 +117,8 @@ mod tests { fn words_from_little_endian_bytes(bytes: &[u8], words: &mut [u32]) { debug_assert_eq!(bytes.len(), 4 * words.len()); - for (four_bytes, word) in bytes.chunks_exact(4).zip(words) { - *word = u32::from_le_bytes(four_bytes.try_into().unwrap()); + for (four_bytes, word) in bytes.as_chunks::<4>().0.iter().zip(words) { + *word = u32::from_le_bytes(*four_bytes); } } diff --git a/src/types.rs b/src/types.rs index 3128af0..d580468 100644 --- a/src/types.rs +++ b/src/types.rs @@ -24,8 +24,6 @@ use p3_fri::TwoAdicFriPcs; use p3_goldilocks::Goldilocks; #[cfg(feature = "cuda")] use p3_matrix::dense::RowMajorMatrix; -#[cfg(feature = "cuda")] -use p3_maybe_rayon::prelude::*; use p3_merkle_tree::MerkleTreeMmcs; use p3_symmetric::{CompressionFunctionFromHasher, SerializingHasher}; @@ -258,14 +256,16 @@ impl StarkGenericConfig for GoldilocksBlake3Config { fri.commit_pow_witnesses.iter_mut().for_each(canonical_base); canonical_base(&mut fri.query_pow_witness); fri.final_poly.iter_mut().for_each(canonical_ext); - for query in &mut fri.query_proofs { - for opening in &mut query.input_proof { - for row in &mut opening.opened_values { + for opening in &mut fri.input_openings { + for query in &mut opening.opened_values { + for row in query { row.iter_mut().for_each(canonical_base); } } - for step in &mut query.commit_phase_openings { - step.sibling_values.iter_mut().for_each(canonical_ext); + } + for step in &mut fri.commit_phase_openings { + for siblings in &mut step.sibling_values { + siblings.iter_mut().for_each(canonical_ext); } } } @@ -328,7 +328,9 @@ impl StarkGenericConfig for GoldilocksBlake3Config { ); Some( flat.values - .chunks_exact(2) + .as_chunks::<2>() + .0 + .iter() .map(|coords| { ExtVal::from_basis_coefficients_slice(coords) .expect("CUDA quotient has two coordinates") @@ -345,7 +347,7 @@ impl StarkGenericConfig for GoldilocksBlake3Config { ) -> Option<(crate::config::Com, crate::config::PcsData)> { use crate::cuda::mmcs::CudaCommitMmcs; let ldes: Option> = inputs - .par_iter() + .iter() .map(|input| { let main = input.stage_1.0.resident(input.stage_1.1)?; let stage2 = input.stage_2.0.resident(input.stage_2.1)?; diff --git a/src/verifier.rs b/src/verifier.rs index 51a7cff..a99af01 100644 --- a/src/verifier.rs +++ b/src/verifier.rs @@ -837,14 +837,14 @@ mod tests { assert_eq!(proof_bytes, second_proof_bytes); assert_eq!( proof_bytes.len(), - 77_637, + 22_017, "proof encoding changed; update only with an intentional protocol review" ); assert_eq!( Blake3.hash_slice(&proof_bytes), [ - 132, 122, 135, 163, 73, 111, 225, 81, 221, 201, 107, 28, 30, 21, 49, 58, 253, 13, - 161, 49, 19, 184, 213, 239, 107, 152, 43, 42, 67, 255, 151, 193, + 17, 55, 200, 132, 9, 170, 249, 195, 35, 120, 241, 110, 48, 5, 20, 38, 250, 196, 78, + 102, 166, 45, 101, 2, 221, 109, 185, 39, 56, 102, 245, 182, ], "proof bytes changed; update only with an intentional protocol review" ); @@ -896,6 +896,20 @@ mod tests { assert!(result.is_err()); } + #[test] + fn test_tampered_opening_proof_rejected() { + let (system, mut proof) = small_system_and_proof(); + // The context-building helper deliberately does not consume the PCS + // proof, but the public verifier must pass it to FRI verification. + proof.opening_proof.final_poly[0] += ExtVal::ONE; + let no_claims: &[&[Val]] = &[]; + let result = system.verify_multiple_claims(no_claims, &proof); + assert!(matches!( + result, + Err(VerificationError::InvalidOpeningArgument(_)) + )); + } + #[test] fn test_tampered_accumulator_rejected() { let (system, mut proof) = small_system_and_proof();