From 9d470ac7a5feb1014918750a27979f5647c94ef2 Mon Sep 17 00:00:00 2001 From: Arthur Paulino Date: Mon, 31 Aug 2026 21:05:29 +0000 Subject: [PATCH 1/4] Bump Plonky3 dependency --- Cargo.lock | 188 +++++++++++++++++--------------- Cargo.toml | 36 +++--- examples/proof_compatibility.rs | 2 +- rust-toolchain.toml | 2 +- src/cuda/mod.rs | 14 +-- src/cuda/pcs.rs | 21 ++-- src/p3_adapter.rs | 7 ++ src/test_circuits/blake3.rs | 4 +- src/types.rs | 12 +- src/verifier.rs | 20 +++- 10 files changed, 170 insertions(+), 136 deletions(-) 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/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/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/mod.rs b/src/cuda/mod.rs index d7c24a2..86e2f89 100644 --- a/src/cuda/mod.rs +++ b/src/cuda/mod.rs @@ -2489,7 +2489,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 +2516,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()) diff --git a/src/cuda/pcs.rs b/src/cuda/pcs.rs index 99f3c90..d1f4049 100644 --- a/src/cuda/pcs.rs +++ b/src/cuda/pcs.rs @@ -38,10 +38,10 @@ 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}; @@ -1067,6 +1067,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 +1093,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 +1106,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, ) }); 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..01121e9 100644 --- a/src/types.rs +++ b/src/types.rs @@ -258,14 +258,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); } } } 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(); From 33b4454aa63b88427b3d3cc09980fe86ddeb7e5f Mon Sep 17 00:00:00 2001 From: Arthur Paulino Date: Tue, 1 Sep 2026 09:25:23 +0000 Subject: [PATCH 2/4] Update the CUDA prover for Plonky3 multiproofs Adapt the CUDA PCS to Plonky3's pruned multiproof protocol while preserving the recursive verifier's compact per-query circuit. Add a verified advice re-encoder that expands native pruned authentication frontiers only at the recursion boundary, leaving native proof serialization and verification unchanged. Bring the CUDA PCS in line with upstream domain semantics by enforcing the minimum committed LDE height, restoring row bit reversal for domain evaluations, and covering Goldilocks log-height 32. Transfer the initial reduced FRI opening into its resident LDE without an extra allocation or device copy. Replace per-column cudaMemcpy2D submissions during Merkle construction with one coalesced gather kernel and cap staging at 32 MiB. Serialize concurrent quotient scheduling to prevent transient recursive-prover OOMs while retaining parallel GPU kernel execution. Update the compatibility digest and keep the CPU-only build independent of CUDA. The CUDA smoke suite passes 50 tests and produces byte-identical 17,213-byte CPU/GPU proofs. On the RTX PRO 6000 q50 Vector.extract_append workload, inner plus outer STARK proving is 65.87s on CPU and 8.81s on CUDA (7.48x), with identical proof sizes. --- cuda/kernels.cu | 93 ++++-- cuda/smoke.sh | 2 +- src/advice.rs | 795 +++++++++++++++++++++++++++++++++++++++++++++++ src/cuda/mmcs.rs | 67 ++-- src/cuda/mod.rs | 75 ++++- src/cuda/pcs.rs | 391 ++++++----------------- src/lib.rs | 1 + src/prover.rs | 2 +- src/types.rs | 8 +- src/verifier.rs | 80 ++++- 10 files changed, 1142 insertions(+), 372 deletions(-) create mode 100644 src/advice.rs 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/src/advice.rs b/src/advice.rs new file mode 100644 index 0000000..d9d9dbf --- /dev/null +++ b/src/advice.rs @@ -0,0 +1,795 @@ +//! Re-encoding of a proof's FRI opening transport for per-query verifiers. +//! +//! Plonky3 v0.6.0 ships FRI query openings as *pruned Merkle multiproofs*: +//! per commitment, one flat list of boundary sibling digests shared by all +//! queries, verified by an amortized bottom-up walk +//! (`MerkleTreeMmcs::verify_batch_pruned`). The in-circuit recursive +//! verifier instead consumes one full authentication path per query — the +//! legacy per-query layout — because its per-query control flow is a far +//! smaller circuit than the amortized walk's sort/merge bookkeeping. +//! +//! [`proof_to_advice_bytes`] converts a verified [`Proof`] into that +//! per-query **advice encoding**: the outer proof fields unchanged, the FRI +//! opening transport expanded from pruned multiproofs to per-query paths. +//! +//! # Soundness: encoding freedom +//! +//! The advice bytes are untrusted prover input to the recursive verifier — +//! never digest-bound, never observed into the transcript. What the +//! transcript binds are the *commitments* (read from the advice and +//! observed), and every expanded sibling digest is authenticated against +//! them by the per-query Merkle checks. Pruning is transport compression: +//! a pruned proof and its expansion authenticate identical opened values +//! against identical commitments, so a valid advice encoding exists iff a +//! valid pruned proof exists for the same statement. Per-query +//! verification is at least as strong as the amortized walk — duplicate +//! queries that disagree on opened values would need a hash collision to +//! both authenticate. +//! +//! # How the expansion recovers interior digests +//! +//! No tree data is available (the input is a proof, not prover data), and +//! the walk that recomputes interior digests lives inside p3. Rather than +//! reimplement it, the expansion *runs* it: `verify_fri` is invoked with a +//! [`MerkleTreeMmcs`] whose compression function records every +//! `(inputs → output)` call. The recorded map is functional (a collision +//! would break Blake3), so each query's full path is read back by walking +//! the tree top-down from its cap entry: at every level the map yields the +//! two children — one continues the path, the other is the sibling the +//! legacy wire format carries. Matrix-injection levels (a shorter matrix's +//! row hash compressed into the running digest) consume no wire sibling, +//! mirroring the per-path verifier's schedule. Query indices come from the +//! same instrumented run, via a challenger wrapper that records +//! `sample_bits` results. + +use std::collections::HashMap; +use std::marker::PhantomData; +use std::sync::{Arc, Mutex}; + +use bincode::serde::encode_to_vec; +use p3_blake3::Blake3; +use p3_challenger::{CanObserve, CanSample, CanSampleBits, FieldChallenger, GrindingChallenger}; +use p3_commit::{ExtensionMmcs, Mmcs, OpenedValuesForRound}; +use p3_field::{BasedVectorSpace, Field}; +use p3_fri::verifier::verify_fri; +use p3_fri::{BatchMultiOpening, CommitPhaseMultiStep, FriProof}; +use p3_fri::{FriParameters as InnerFriParameters, TwoAdicFriFolding}; +use p3_matrix::Dimensions; +use p3_merkle_tree::MerkleTreeMmcs; +use p3_symmetric::{ + CompressionFunctionFromHasher, CryptographicHasher, PseudoCompressionFunction, + SerializingHasher, +}; +use p3_util::log2_strict_usize; +use serde::{Deserialize, Serialize}; + +use crate::prover::{Commitments, Proof}; +use crate::system::System; +use crate::types::{ + Challenger, Commitment, CommitmentParameters, ExtVal, FriParameters, GoldilocksBlake3Config, + Val, +}; + +type Digest = [u8; 32]; +type Blake3Compress = CompressionFunctionFromHasher; +type RecMmcs = MerkleTreeMmcs, RecordingCompress, 2, 32>; +type RecExtMmcs = ExtensionMmcs; + +/// Why an advice re-encoding could not be produced. Every variant except +/// [`AdviceError::Encode`] indicates an invalid proof or a bug: expansion +/// only runs after native verification succeeds. +#[derive(Debug)] +pub enum AdviceError { + /// Native verification of the proof failed; nothing was expanded. + Verification(String), + /// The instrumented PCS run failed (unreachable after native + /// verification passed, absent a bug). + Recording(String), + /// A digest needed for path read-back was never computed by the + /// instrumented run. + MissingDigest, + /// A cap index fell outside the commitment. + CapIndexOutOfBounds, + /// The recorded `sample_bits` log did not contain the query indices. + MissingQueryIndices, + /// Serialization of the legacy layout failed. + Encode(bincode::error::EncodeError), +} + +// --------------------------------------------------------------------------- +// Instrumentation +// --------------------------------------------------------------------------- + +/// A compression function that forwards to Blake3 and records every call. +/// The map is keyed by output: Blake3 collisions aside, each digest has a +/// unique preimage pair, so read-back is unambiguous. +#[derive(Clone)] +struct RecordingCompress { + inner: Blake3Compress, + log: Arc>>, +} + +impl PseudoCompressionFunction for RecordingCompress { + fn compress(&self, input: [Digest; 2]) -> Digest { + let output = self.inner.compress(input); + self.log.lock().unwrap().insert(output, input); + output + } +} + +/// A challenger that forwards everything and records `sample_bits` calls. +/// FRI query indices are the last `num_queries` recorded samples: they are +/// drawn after the query proof-of-work check, and nothing samples bits +/// after them. +#[derive(Clone)] +struct RecordingChallenger { + inner: C, + samples: Arc>>, +} + +impl> CanObserve for RecordingChallenger { + fn observe(&mut self, value: T) { + self.inner.observe(value); + } +} + +impl> CanSample for RecordingChallenger { + fn sample(&mut self) -> T { + self.inner.sample() + } +} + +impl> CanSampleBits for RecordingChallenger { + fn sample_bits(&mut self, bits: usize) -> usize { + let value = self.inner.sample_bits(bits); + self.samples.lock().unwrap().push(value); + value + } +} + +impl> FieldChallenger for RecordingChallenger {} + +impl GrindingChallenger for RecordingChallenger { + type Witness = C::Witness; + + fn grind(&mut self, bits: usize) -> Self::Witness { + self.inner.grind(bits) + } + + // Delegate so the inner challenger's own `sample_bits` runs unrecorded: + // proof-of-work samples never pollute the query-index log. + fn check_witness(&mut self, bits: usize, witness: Self::Witness) -> bool { + self.inner.check_witness(bits, witness) + } +} + +// --------------------------------------------------------------------------- +// Legacy wire layout (the advice encoding) +// --------------------------------------------------------------------------- + +/// One round's input opening for one query: the opened rows and a full +/// Merkle authentication path. +#[derive(Serialize, Deserialize)] +pub struct AdviceBatchOpening { + pub opened_values: Vec>, + pub opening_proof: Vec, +} + +/// One FRI folding step for one query: the folding arity, the `arity - 1` +/// sibling evaluations, and a full path into that round's commitment. +#[derive(Serialize, Deserialize)] +pub struct AdviceCommitPhaseStep { + pub log_arity: u8, + pub sibling_values: Vec, + pub opening_proof: Vec, +} + +/// One query's complete opening data: input openings per round, then one +/// folding step per commit phase. +#[derive(Serialize, Deserialize)] +pub struct AdviceQueryProof { + pub input_proof: Vec, + pub commit_phase_openings: Vec, +} + +/// The FRI proof in per-query transport. +#[derive(Serialize, Deserialize)] +pub struct AdviceFriProof { + pub commit_phase_commits: Vec, + pub commit_pow_witnesses: Vec, + pub query_proofs: Vec, + pub final_poly: Vec, + pub query_pow_witness: Val, +} + +/// The full advice proof: every field of [`Proof`] unchanged except the +/// opening transport. +#[derive(Serialize, Deserialize)] +pub struct AdviceProof { + pub active: Vec, + pub commitments: Commitments, + pub intermediate_accumulators: Vec, + pub log_degrees: Vec, + pub opening_proof: AdviceFriProof, + pub quotient_opened_values: OpenedValuesForRound, + pub preprocessed_opened_values: Option>, + pub stage_1_opened_values: OpenedValuesForRound, + pub stage_2_opened_values: OpenedValuesForRound, +} + +// --------------------------------------------------------------------------- +// Path read-back +// --------------------------------------------------------------------------- + +/// One step of the per-path verifier's ascent, derived from the matrix +/// dimensions exactly as `verify_batch` derives it. +enum LevelOp { + /// An arity-`step` compression; the path consumes `step - 1` wire + /// siblings here. + Fold { step: usize }, + /// A shorter matrix's row hash is compressed into the running digest; + /// nothing on the wire. + Inject, +} + +/// The ascent schedule (leaf to cap) for a tree of the given dimensions, +/// mirroring `verify_batch`'s traversal: one `Fold` per arity-schedule +/// step, an `Inject` after any step whose folded height picks up shorter +/// matrices. +fn ascent_schedule( + schedule: &[usize], + dimensions: &[Dimensions], + max_height: usize, +) -> Vec { + let mut heights: Vec = dimensions.iter().map(|d| d.height).collect(); + heights.sort_unstable_by(|a, b| b.cmp(a)); + let leaf_npt = max_height.next_power_of_two(); + // Heights hashed into the leaf layer are consumed up front. + let mut next = heights + .iter() + .position(|h| h.next_power_of_two() != leaf_npt) + .unwrap_or(heights.len()); + + let mut ops = Vec::new(); + let mut curr_height_padded = max_height.next_multiple_of(2); + for &step in schedule { + ops.push(LevelOp::Fold { step }); + let logical_next = curr_height_padded / step; + curr_height_padded = logical_next.next_multiple_of(2); + let logical_next_npt = logical_next.next_power_of_two(); + if next < heights.len() && heights[next].next_power_of_two() == logical_next_npt { + ops.push(LevelOp::Inject); + while next < heights.len() && heights[next].next_power_of_two() == logical_next_npt { + next += 1; + } + } + } + ops +} + +/// Reads one query's full authentication path out of the recorded +/// compression map by walking its tree top-down from the cap entry, +/// emitting siblings in the bottom-up order the legacy wire carries. +/// +/// `leaf_check`, when supplied, is the expected leaf digest (the hash of +/// the query's opened rows); the walk must land exactly there. +fn expand_path( + map: &HashMap, + commitment: &Commitment, + schedule: &[usize], + dimensions: &[Dimensions], + max_height: usize, + index: usize, + leaf_check: Option, +) -> Result, AdviceError> { + let ops = ascent_schedule(schedule, dimensions, max_height); + + // Per-fold node index at each level, bottom-up. + let mut level_indices = Vec::new(); + let mut idx = index; + for op in &ops { + if let LevelOp::Fold { step } = op { + level_indices.push(idx); + idx /= step; + } + } + let cap_index = idx; + if cap_index >= commitment.num_roots() { + return Err(AdviceError::CapIndexOutOfBounds); + } + + let mut digest: Digest = commitment[cap_index]; + let mut fold_level = level_indices.len(); + // Wire order is bottom-up; the walk is top-down, so collect reversed. + let mut siblings_rev: Vec = Vec::new(); + for op in ops.iter().rev() { + let inputs = map.get(&digest).ok_or(AdviceError::MissingDigest)?; + match op { + LevelOp::Inject => { + // inject_inputs = [running digest, row hash]; descend left. + digest = inputs[0]; + } + LevelOp::Fold { step } => { + fold_level -= 1; + let pos_in_group = level_indices[fold_level] % step; + // The recording mmcs is binary (N = 2), so `step` is 2. + for k in (0..*step).rev() { + if k != pos_in_group { + siblings_rev.push(inputs[k]); + } + } + digest = inputs[pos_in_group]; + } + } + } + if let Some(expected) = leaf_check + && digest != expected + { + return Err(AdviceError::MissingDigest); + } + siblings_rev.reverse(); + Ok(siblings_rev) +} + +// --------------------------------------------------------------------------- +// Re-encoding +// --------------------------------------------------------------------------- + +/// Converts a natively-verified proof into the per-query advice encoding. +/// +/// `commitment_parameters` and `fri_parameters` must be the ones the +/// system's config was built from (the config does not expose them back). +/// Verifies the proof natively first and refuses to expand on failure. +pub fn proof_to_advice_bytes( + system: &System, + commitment_parameters: CommitmentParameters, + fri_parameters: FriParameters, + claims: &[&[Val]], + proof: &Proof, +) -> Result, AdviceError> { + system + .verify_multiple_claims(claims, proof) + .map_err(|e| AdviceError::Verification(format!("{e:?}")))?; + let ctx = system + .pcs_verification_context(claims, proof) + .map_err(|e| AdviceError::Verification(format!("{e:?}")))?; + + // Instrumented components sharing one compression log. + let log: Arc>> = Arc::new(Mutex::new(HashMap::new())); + let rec_compress = RecordingCompress { + inner: Blake3Compress::new(Blake3), + log: Arc::clone(&log), + }; + let rec_mmcs = RecMmcs::new( + SerializingHasher::new(Blake3), + rec_compress, + commitment_parameters.cap_height, + ); + let rec_params = InnerFriParameters { + log_blowup: commitment_parameters.log_blowup, + log_final_poly_len: fri_parameters.log_final_poly_len, + max_log_arity: fri_parameters.max_log_arity, + num_queries: fri_parameters.num_queries, + commit_proof_of_work_bits: fri_parameters.commit_proof_of_work_bits, + query_proof_of_work_bits: fri_parameters.query_proof_of_work_bits, + mmcs: ExtensionMmcs::::new(rec_mmcs.clone()), + }; + + // The stored FRI proof, re-typed over the instrumented mmcs. The + // commitment, digest and pruned-path types are identical; only the + // phantom mmcs parameter changes. + let fri = &proof.opening_proof; + let rec_proof: FriProof>> = + FriProof { + commit_phase_commits: fri.commit_phase_commits.clone(), + commit_pow_witnesses: fri.commit_pow_witnesses.clone(), + input_openings: fri + .input_openings + .iter() + .map(|o| BatchMultiOpening { + opened_values: o.opened_values.clone(), + opening_proof: o.opening_proof.clone(), + }) + .collect(), + commit_phase_openings: fri + .commit_phase_openings + .iter() + .map(|s| CommitPhaseMultiStep { + log_arity: s.log_arity, + sibling_values: s.sibling_values.clone(), + opening_proof: s.opening_proof.clone(), + }) + .collect(), + final_poly: fri.final_poly.clone(), + query_pow_witness: fri.query_pow_witness, + }; + + // Run the PCS phase exactly as `TwoAdicFriPcs::verify` would — the + // evaluation observations, then `verify_fri` — against the context's + // rounds and transcript state, with the instrumented components. + let mut challenger = RecordingChallenger:: { + inner: ctx.challenger, + samples: Arc::new(Mutex::new(Vec::new())), + }; + for (_, round) in &ctx.rounds { + for (_, mat) in round { + for (_, point) in mat { + challenger.observe_algebra_slice(point); + } + } + } + let folding = TwoAdicFriFolding::< + Vec>, + >::Error, + >(PhantomData); + verify_fri( + &folding, + &rec_params, + &rec_proof, + &mut challenger, + &ctx.rounds, + &rec_mmcs, + ) + .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; + + // The query indices are the last `num_queries` recorded samples. + let samples = challenger.samples.lock().unwrap(); + let query_indices: Vec = samples + .len() + .checked_sub(fri_parameters.num_queries) + .map(|start| samples[start..].to_vec()) + .ok_or(AdviceError::MissingQueryIndices)?; + drop(samples); + let map = log.lock().unwrap(); + + let log_arities: Vec = fri + .commit_phase_openings + .iter() + .map(|s| usize::from(s.log_arity)) + .collect(); + let total_log_reduction: usize = log_arities.iter().sum(); + let log_global_max_height = + total_log_reduction + commitment_parameters.log_blowup + fri_parameters.log_final_poly_len; + + let hasher = SerializingHasher::new(Blake3); + + // Input rounds: per round, per query, opened rows + expanded path. + // Dimensions and reduced indices mirror p3's `open_inputs`. + let mut input_openings_per_query: Vec> = (0..fri_parameters + .num_queries) + .map(|_| Vec::new()) + .collect(); + for ((commit, mats), batch) in ctx.rounds.iter().zip(&fri.input_openings) { + let heights: Vec = mats + .iter() + .map(|(domain, _)| domain.size() << commitment_parameters.log_blowup) + .collect(); + let dims: Vec = heights + .iter() + .zip(mats) + .map(|(&height, (_, points))| Dimensions { + width: points.first().map_or(0, |(_, values)| values.len()), + height, + }) + .collect(); + let max_height = heights.iter().copied().max().unwrap_or(1); + let bits_reduced = log_global_max_height - log2_strict_usize(max_height); + let schedule = rec_mmcs + .proof_arity_schedule(&dims) + .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; + + // Matrices whose padded height matches the tallest are hashed into + // the leaf digest, in tallest-first order — mirroring `verify_batch`. + let leaf_npt = max_height.next_power_of_two(); + let mut order: Vec = (0..dims.len()).collect(); + order.sort_by_key(|&i| std::cmp::Reverse(dims[i].height)); + let leaf_matrices: Vec = order + .iter() + .copied() + .take_while(|&i| dims[i].height.next_power_of_two() == leaf_npt) + .collect(); + + for (q, per_query) in input_openings_per_query.iter_mut().enumerate() { + let reduced_index = query_indices[q] >> bits_reduced; + let opened_values = batch.opened_values[q].clone(); + let leaf = hasher + .hash_iter_slices(leaf_matrices.iter().map(|&mi| opened_values[mi].as_slice())); + let opening_proof = expand_path( + &map, + commit, + &schedule, + &dims, + max_height, + reduced_index, + Some(leaf), + )?; + per_query.push(AdviceBatchOpening { + opened_values, + opening_proof, + }); + } + } + + // Commit-phase rounds: single-matrix trees of the folded codewords, + // flattened to base columns by the extension mmcs. The wire sibling + // values come straight off the multiproof; only the paths are read + // back. No leaf check: the folded row is not materialized here — the + // walk's landing digest is pinned by the functional map and the cap. + let ext_d = >::DIMENSION; + let mut steps_per_query: Vec> = (0..fri_parameters.num_queries) + .map(|_| Vec::new()) + .collect(); + let mut log_current_height = log_global_max_height; + for (round, step) in fri.commit_phase_openings.iter().enumerate() { + let log_arity = usize::from(step.log_arity); + let arity = 1 << log_arity; + let log_folded_height = log_current_height - log_arity; + let dims = [Dimensions { + width: arity * ext_d, + height: 1 << log_folded_height, + }]; + let schedule = rec_mmcs + .proof_arity_schedule(&dims) + .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; + let commit = &fri.commit_phase_commits[round]; + let bits_consumed: usize = log_arities[..=round].iter().sum(); + for (q, per_query) in steps_per_query.iter_mut().enumerate() { + let group_index = query_indices[q] >> bits_consumed; + let opening_proof = expand_path( + &map, + commit, + &schedule, + &dims, + 1 << log_folded_height, + group_index, + None, + )?; + per_query.push(AdviceCommitPhaseStep { + log_arity: step.log_arity, + sibling_values: step.sibling_values[q].clone(), + opening_proof, + }); + } + log_current_height = log_folded_height; + } + + let query_proofs: Vec = input_openings_per_query + .into_iter() + .zip(steps_per_query) + .map(|(input_proof, commit_phase_openings)| AdviceQueryProof { + input_proof, + commit_phase_openings, + }) + .collect(); + + let advice = AdviceProof { + active: proof.active.clone(), + commitments: Commitments { + stage_1_trace: proof.commitments.stage_1_trace.clone(), + stage_2_trace: proof.commitments.stage_2_trace.clone(), + quotient_chunks: proof.commitments.quotient_chunks.clone(), + }, + intermediate_accumulators: proof.intermediate_accumulators.clone(), + log_degrees: proof.log_degrees.clone(), + opening_proof: AdviceFriProof { + commit_phase_commits: fri.commit_phase_commits.clone(), + commit_pow_witnesses: fri.commit_pow_witnesses.clone(), + query_proofs, + final_poly: fri.final_poly.clone(), + query_pow_witness: fri.query_pow_witness, + }, + quotient_opened_values: proof.quotient_opened_values.clone(), + preprocessed_opened_values: proof.preprocessed_opened_values.clone(), + stage_1_opened_values: proof.stage_1_opened_values.clone(), + stage_2_opened_values: proof.stage_2_opened_values.clone(), + }; + encode_to_vec(&advice, Proof::::serde_config()) + .map_err(AdviceError::Encode) +} + +#[cfg(all(test, not(feature = "cuda")))] +mod tests { + use bincode::serde::decode_from_slice; + use p3_field::PrimeCharacteristicRing; + use p3_matrix::dense::RowMajorMatrix; + + use super::*; + use crate::system::ProverKey; + use crate::system::{CircuitInputs, SystemWitness}; + + fn parameters(cap_height: usize) -> (CommitmentParameters, FriParameters) { + ( + CommitmentParameters { + log_blowup: 1, + cap_height, + }, + FriParameters { + log_final_poly_len: 0, + max_log_arity: 1, + num_queries: 8, + commit_proof_of_work_bits: 0, + query_proof_of_work_bits: 0, + }, + ) + } + + /// Two constraint-free circuits at different trace heights (so the + /// stage-1 tree exercises matrix injection), one with a preprocessed + /// matrix (so the preprocessed round exists). + fn test_system( + cap_height: usize, + ) -> ( + System, + ProverKey, + ) { + let (cp, fp) = parameters(cap_height); + let config = GoldilocksBlake3Config::new(cp, fp); + let preprocessed = RowMajorMatrix::new((0..8u32).map(Val::from_u32).collect::>(), 1); + let inputs = [ + CircuitInputs { + main_width: 2, + preprocessed: Some(preprocessed), + ..Default::default() + }, + CircuitInputs { + main_width: 3, + ..Default::default() + }, + ]; + System::new(config, inputs) + } + + fn test_proof( + system: &System, + key: &ProverKey, + ) -> Proof { + let trace_1 = RowMajorMatrix::new((0..16u32).map(Val::from_u32).collect::>(), 2); + let trace_2 = RowMajorMatrix::new( + (0..12u32) + .map(|i| Val::from_u32(7 * i + 3)) + .collect::>(), + 3, + ); + let witness = SystemWitness::from_stage_1(vec![trace_1, trace_2], system); + system.prove_multiple_claims(key, &[], witness) + } + + #[test] + fn advice_expands_and_paths_verify_per_query() { + for cap_height in [0, 1] { + let (cp, fp) = parameters(cap_height); + let (system, key) = test_system(cap_height); + let proof = test_proof(&system, &key); + let bytes = proof_to_advice_bytes(&system, cp, fp, &[], &proof) + .unwrap_or_else(|e| panic!("advice expansion failed: {e:?}")); + + let (advice, consumed): (AdviceProof, usize) = + decode_from_slice(&bytes, Proof::::serde_config()) + .expect("advice bytes decode under the same serde config"); + assert_eq!(consumed, bytes.len(), "no trailing bytes"); + assert_eq!(advice.opening_proof.query_proofs.len(), fp.num_queries); + + // Recover the query indices the same way the expansion did. + let ctx = system.pcs_verification_context(&[], &proof).unwrap(); + let plain_mmcs = MerkleTreeMmcs::::new( + SerializingHasher::new(Blake3), + Blake3Compress::new(Blake3), + cp.cap_height, + ); + let plain_params = InnerFriParameters { + log_blowup: cp.log_blowup, + log_final_poly_len: fp.log_final_poly_len, + max_log_arity: fp.max_log_arity, + num_queries: fp.num_queries, + commit_proof_of_work_bits: fp.commit_proof_of_work_bits, + query_proof_of_work_bits: fp.query_proof_of_work_bits, + mmcs: ExtensionMmcs::::new(plain_mmcs.clone()), + }; + let mut challenger = RecordingChallenger:: { + inner: ctx.challenger, + samples: Arc::new(Mutex::new(Vec::new())), + }; + for (_, round) in &ctx.rounds { + for (_, mat) in round { + for (_, point) in mat { + challenger.observe_algebra_slice(point); + } + } + } + let folding = TwoAdicFriFolding::< + Vec>, + , Blake3Compress, 2, 32> as Mmcs< + Val, + >>::Error, + >(PhantomData); + verify_fri( + &folding, + &plain_params, + &proof.opening_proof, + &mut challenger, + &ctx.rounds, + &plain_mmcs, + ) + .unwrap(); + let samples = challenger.samples.lock().unwrap(); + let indices = &samples[samples.len() - fp.num_queries..]; + + let log_arities: Vec = proof + .opening_proof + .commit_phase_openings + .iter() + .map(|s| usize::from(s.log_arity)) + .collect(); + let log_global_max_height = + log_arities.iter().sum::() + cp.log_blowup + fp.log_final_poly_len; + + // Oracle: every expanded input-round path must satisfy p3's own + // per-path verifier. + for (round, (commit, mats)) in ctx.rounds.iter().enumerate() { + let heights: Vec = mats + .iter() + .map(|(domain, _)| domain.size() << cp.log_blowup) + .collect(); + let dims: Vec = heights + .iter() + .zip(mats) + .map(|(&height, (_, points))| Dimensions { + width: points.first().map_or(0, |(_, values)| values.len()), + height, + }) + .collect(); + let max_height = heights.iter().copied().max().unwrap(); + let bits_reduced = log_global_max_height - log2_strict_usize(max_height); + for (q, &index) in indices.iter().enumerate() { + let opening = &advice.opening_proof.query_proofs[q].input_proof[round]; + plain_mmcs + .verify_batch( + commit, + &dims, + index >> bits_reduced, + p3_commit::BatchOpeningRef::new( + &opening.opened_values, + &opening.opening_proof, + ), + ) + .unwrap_or_else(|e| { + panic!("round {round} query {q}: expanded path rejected: {e:?}") + }); + } + } + + // Commit-phase paths: length and sibling-value parity with the + // multiproof (the Lean verifier is the end-to-end oracle here). + let mut log_current = log_global_max_height; + for (round, step) in proof.opening_proof.commit_phase_openings.iter().enumerate() { + let log_arity = usize::from(step.log_arity); + let log_folded = log_current - log_arity; + let expected_levels = log_folded.saturating_sub(cp.cap_height); + for (q, qp) in advice.opening_proof.query_proofs.iter().enumerate() { + let s = &qp.commit_phase_openings[round]; + assert_eq!(usize::from(s.log_arity), log_arity); + assert_eq!(s.sibling_values, step.sibling_values[q]); + assert_eq!( + s.opening_proof.len(), + expected_levels, + "round {round} query {q}: phase path length" + ); + } + log_current = log_folded; + } + } + } + + #[test] + fn advice_refuses_invalid_proof() { + let cap_height = 0; + let (cp, fp) = parameters(cap_height); + let (system, key) = test_system(cap_height); + let mut proof = test_proof(&system, &key); + proof.log_degrees[0] ^= 1; + assert!(matches!( + proof_to_advice_bytes(&system, cp, fp, &[], &proof), + Err(AdviceError::Verification(_)) + )); + } +} 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 86e2f89..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, @@ -2699,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()) }) @@ -2746,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 d1f4049..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,10 +27,7 @@ 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::{ @@ -44,10 +40,10 @@ 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 @@ -1158,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()) @@ -1228,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( @@ -1276,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/lib.rs b/src/lib.rs index a3afd0a..3f18e65 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1,3 +1,4 @@ +pub mod advice; pub mod config; #[cfg(feature = "cuda")] pub mod cuda; diff --git a/src/prover.rs b/src/prover.rs index 57a574c..5276604 100644 --- a/src/prover.rs +++ b/src/prover.rs @@ -255,7 +255,7 @@ impl Clone for Proof { } impl Proof { - fn serde_config() -> Configuration { + pub(crate) fn serde_config() -> Configuration { standard().with_little_endian().with_fixed_int_encoding() } diff --git a/src/types.rs b/src/types.rs index 01121e9..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}; @@ -330,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") @@ -347,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 a99af01..bf1dc3f 100644 --- a/src/verifier.rs +++ b/src/verifier.rs @@ -159,7 +159,7 @@ //! actual low-degree-extension values of the witness. Do not use it when the //! witness must remain hidden from the verifier. -use crate::config::{PcsError, StarkGenericConfig, Val}; +use crate::config::{Com, Domain, PcsError, StarkGenericConfig, Val}; use crate::ensure_eq; use crate::eval::VarValues; use crate::lookup::fingerprint; @@ -191,6 +191,29 @@ pub enum VerificationError { UnbalancedChannel, } +pub(crate) type OpeningRounds = Vec<( + Com, + Vec<( + Domain, + Vec<( + ::Challenge, + Vec<::Challenge>, + )>, + )>, +)>; + +pub(crate) struct PcsVerificationContext { + pub rounds: OpeningRounds, + pub challenger: SC::Challenger, + pub zeta: SC::Challenge, + pub lookup_argument_challenge: SC::Challenge, + pub fingerprint_challenge: SC::Challenge, + pub constraint_challenge: SC::Challenge, + pub claim_accumulator: SC::Challenge, + pub active_indices: Vec, + pub quotient_degrees: Vec, +} + impl System { /// Verifies a STARK proof against a single claim. pub fn verify( @@ -204,12 +227,11 @@ impl System { self.verify_multiple_claims(&[claim], proof) } - /// Verifies a STARK proof against multiple claims. - pub fn verify_multiple_claims( + pub(crate) fn pcs_verification_context( &self, claims: &[&[Val]], proof: &Proof, - ) -> Result<(), VerificationError>> + ) -> Result, VerificationError>> where Val: TwoAdicField, { @@ -218,7 +240,7 @@ impl System { commitments, intermediate_accumulators, log_degrees, - opening_proof, + opening_proof: _, quotient_opened_values, preprocessed_opened_values, stage_1_opened_values, @@ -406,11 +428,57 @@ impl System { preprocessed_trace_evaluations, )); } + Ok(PcsVerificationContext { + rounds: coms_to_verify, + challenger, + zeta, + lookup_argument_challenge, + fingerprint_challenge, + constraint_challenge, + claim_accumulator: acc, + active_indices, + quotient_degrees, + }) + } + + /// Verifies a STARK proof against multiple claims. + pub fn verify_multiple_claims( + &self, + claims: &[&[Val]], + proof: &Proof, + ) -> Result<(), VerificationError>> + where + Val: TwoAdicField, + { + let Proof { + intermediate_accumulators, + log_degrees, + opening_proof, + quotient_opened_values, + preprocessed_opened_values, + stage_1_opened_values, + stage_2_opened_values, + .. + } = proof; + let PcsVerificationContext { + rounds, + mut challenger, + zeta, + lookup_argument_challenge, + fingerprint_challenge, + constraint_challenge, + claim_accumulator, + active_indices, + quotient_degrees, + } = self.pcs_verification_context(claims, proof)?; + let mut acc = claim_accumulator; + let pcs = self.config.pcs(); + // Soundness: FRI proximity test. Verifies that the committed polynomials // are close to low-degree polynomials and that the claimed evaluations are // consistent with the commitments. Soundness error ≤ ρ^num_queries, where // ρ = 2^(-log_blowup). This is the dominant term in the overall bound. - pcs.verify(coms_to_verify, opening_proof, &mut challenger) + pcs.verify(rounds, opening_proof, &mut challenger) .map_err(VerificationError::InvalidOpeningArgument)?; // use the opened values to compute the composition polynomial for each circuit From fac47fa17b8640d708949a6fb64550f64bb27cc5 Mon Sep 17 00:00:00 2001 From: Arthur Paulino Date: Tue, 1 Sep 2026 12:39:31 +0000 Subject: [PATCH 3/4] Fix the Nix Rust 1.98 toolchain hash Update the fixed-output hash used by Fenix to match the Rust 1.98 channel manifest selected by rust-toolchain.toml. This restores both the packaged Nix build and the development shell without changing the pinned toolchain or flake inputs. --- flake.nix | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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; From 2a1d3a59efc47a521c8cab3311b20fb6859676e3 Mon Sep 17 00:00:00 2001 From: Samuel Burnham <45365069+samuelburnham@users.noreply.github.com> Date: Tue, 1 Sep 2026 20:23:14 -0400 Subject: [PATCH 4/4] Remove the unused per-query advice module (#78) `advice.rs` re-encoded a pruned FRI multiproof into one authentication path per query, for a recursive verifier that consumed the legacy per-query transport. That verifier now consumes the native pruned multiproof directly, so the expansion has no caller: nothing in this crate references `advice`, and the only downstream user (aiur) imported just its `AdviceError` type, which it can define locally. Drop the module and its `pub mod advice;` declaration. The library is self-consistent without it (lib + tests build clean; 35 tests pass). The frontier expansion remains recoverable from history if a fallback to the per-query transport is ever wanted. --- src/advice.rs | 795 ------------------------------------------------ src/lib.rs | 1 - src/prover.rs | 2 +- src/verifier.rs | 80 +---- 4 files changed, 7 insertions(+), 871 deletions(-) delete mode 100644 src/advice.rs diff --git a/src/advice.rs b/src/advice.rs deleted file mode 100644 index d9d9dbf..0000000 --- a/src/advice.rs +++ /dev/null @@ -1,795 +0,0 @@ -//! Re-encoding of a proof's FRI opening transport for per-query verifiers. -//! -//! Plonky3 v0.6.0 ships FRI query openings as *pruned Merkle multiproofs*: -//! per commitment, one flat list of boundary sibling digests shared by all -//! queries, verified by an amortized bottom-up walk -//! (`MerkleTreeMmcs::verify_batch_pruned`). The in-circuit recursive -//! verifier instead consumes one full authentication path per query — the -//! legacy per-query layout — because its per-query control flow is a far -//! smaller circuit than the amortized walk's sort/merge bookkeeping. -//! -//! [`proof_to_advice_bytes`] converts a verified [`Proof`] into that -//! per-query **advice encoding**: the outer proof fields unchanged, the FRI -//! opening transport expanded from pruned multiproofs to per-query paths. -//! -//! # Soundness: encoding freedom -//! -//! The advice bytes are untrusted prover input to the recursive verifier — -//! never digest-bound, never observed into the transcript. What the -//! transcript binds are the *commitments* (read from the advice and -//! observed), and every expanded sibling digest is authenticated against -//! them by the per-query Merkle checks. Pruning is transport compression: -//! a pruned proof and its expansion authenticate identical opened values -//! against identical commitments, so a valid advice encoding exists iff a -//! valid pruned proof exists for the same statement. Per-query -//! verification is at least as strong as the amortized walk — duplicate -//! queries that disagree on opened values would need a hash collision to -//! both authenticate. -//! -//! # How the expansion recovers interior digests -//! -//! No tree data is available (the input is a proof, not prover data), and -//! the walk that recomputes interior digests lives inside p3. Rather than -//! reimplement it, the expansion *runs* it: `verify_fri` is invoked with a -//! [`MerkleTreeMmcs`] whose compression function records every -//! `(inputs → output)` call. The recorded map is functional (a collision -//! would break Blake3), so each query's full path is read back by walking -//! the tree top-down from its cap entry: at every level the map yields the -//! two children — one continues the path, the other is the sibling the -//! legacy wire format carries. Matrix-injection levels (a shorter matrix's -//! row hash compressed into the running digest) consume no wire sibling, -//! mirroring the per-path verifier's schedule. Query indices come from the -//! same instrumented run, via a challenger wrapper that records -//! `sample_bits` results. - -use std::collections::HashMap; -use std::marker::PhantomData; -use std::sync::{Arc, Mutex}; - -use bincode::serde::encode_to_vec; -use p3_blake3::Blake3; -use p3_challenger::{CanObserve, CanSample, CanSampleBits, FieldChallenger, GrindingChallenger}; -use p3_commit::{ExtensionMmcs, Mmcs, OpenedValuesForRound}; -use p3_field::{BasedVectorSpace, Field}; -use p3_fri::verifier::verify_fri; -use p3_fri::{BatchMultiOpening, CommitPhaseMultiStep, FriProof}; -use p3_fri::{FriParameters as InnerFriParameters, TwoAdicFriFolding}; -use p3_matrix::Dimensions; -use p3_merkle_tree::MerkleTreeMmcs; -use p3_symmetric::{ - CompressionFunctionFromHasher, CryptographicHasher, PseudoCompressionFunction, - SerializingHasher, -}; -use p3_util::log2_strict_usize; -use serde::{Deserialize, Serialize}; - -use crate::prover::{Commitments, Proof}; -use crate::system::System; -use crate::types::{ - Challenger, Commitment, CommitmentParameters, ExtVal, FriParameters, GoldilocksBlake3Config, - Val, -}; - -type Digest = [u8; 32]; -type Blake3Compress = CompressionFunctionFromHasher; -type RecMmcs = MerkleTreeMmcs, RecordingCompress, 2, 32>; -type RecExtMmcs = ExtensionMmcs; - -/// Why an advice re-encoding could not be produced. Every variant except -/// [`AdviceError::Encode`] indicates an invalid proof or a bug: expansion -/// only runs after native verification succeeds. -#[derive(Debug)] -pub enum AdviceError { - /// Native verification of the proof failed; nothing was expanded. - Verification(String), - /// The instrumented PCS run failed (unreachable after native - /// verification passed, absent a bug). - Recording(String), - /// A digest needed for path read-back was never computed by the - /// instrumented run. - MissingDigest, - /// A cap index fell outside the commitment. - CapIndexOutOfBounds, - /// The recorded `sample_bits` log did not contain the query indices. - MissingQueryIndices, - /// Serialization of the legacy layout failed. - Encode(bincode::error::EncodeError), -} - -// --------------------------------------------------------------------------- -// Instrumentation -// --------------------------------------------------------------------------- - -/// A compression function that forwards to Blake3 and records every call. -/// The map is keyed by output: Blake3 collisions aside, each digest has a -/// unique preimage pair, so read-back is unambiguous. -#[derive(Clone)] -struct RecordingCompress { - inner: Blake3Compress, - log: Arc>>, -} - -impl PseudoCompressionFunction for RecordingCompress { - fn compress(&self, input: [Digest; 2]) -> Digest { - let output = self.inner.compress(input); - self.log.lock().unwrap().insert(output, input); - output - } -} - -/// A challenger that forwards everything and records `sample_bits` calls. -/// FRI query indices are the last `num_queries` recorded samples: they are -/// drawn after the query proof-of-work check, and nothing samples bits -/// after them. -#[derive(Clone)] -struct RecordingChallenger { - inner: C, - samples: Arc>>, -} - -impl> CanObserve for RecordingChallenger { - fn observe(&mut self, value: T) { - self.inner.observe(value); - } -} - -impl> CanSample for RecordingChallenger { - fn sample(&mut self) -> T { - self.inner.sample() - } -} - -impl> CanSampleBits for RecordingChallenger { - fn sample_bits(&mut self, bits: usize) -> usize { - let value = self.inner.sample_bits(bits); - self.samples.lock().unwrap().push(value); - value - } -} - -impl> FieldChallenger for RecordingChallenger {} - -impl GrindingChallenger for RecordingChallenger { - type Witness = C::Witness; - - fn grind(&mut self, bits: usize) -> Self::Witness { - self.inner.grind(bits) - } - - // Delegate so the inner challenger's own `sample_bits` runs unrecorded: - // proof-of-work samples never pollute the query-index log. - fn check_witness(&mut self, bits: usize, witness: Self::Witness) -> bool { - self.inner.check_witness(bits, witness) - } -} - -// --------------------------------------------------------------------------- -// Legacy wire layout (the advice encoding) -// --------------------------------------------------------------------------- - -/// One round's input opening for one query: the opened rows and a full -/// Merkle authentication path. -#[derive(Serialize, Deserialize)] -pub struct AdviceBatchOpening { - pub opened_values: Vec>, - pub opening_proof: Vec, -} - -/// One FRI folding step for one query: the folding arity, the `arity - 1` -/// sibling evaluations, and a full path into that round's commitment. -#[derive(Serialize, Deserialize)] -pub struct AdviceCommitPhaseStep { - pub log_arity: u8, - pub sibling_values: Vec, - pub opening_proof: Vec, -} - -/// One query's complete opening data: input openings per round, then one -/// folding step per commit phase. -#[derive(Serialize, Deserialize)] -pub struct AdviceQueryProof { - pub input_proof: Vec, - pub commit_phase_openings: Vec, -} - -/// The FRI proof in per-query transport. -#[derive(Serialize, Deserialize)] -pub struct AdviceFriProof { - pub commit_phase_commits: Vec, - pub commit_pow_witnesses: Vec, - pub query_proofs: Vec, - pub final_poly: Vec, - pub query_pow_witness: Val, -} - -/// The full advice proof: every field of [`Proof`] unchanged except the -/// opening transport. -#[derive(Serialize, Deserialize)] -pub struct AdviceProof { - pub active: Vec, - pub commitments: Commitments, - pub intermediate_accumulators: Vec, - pub log_degrees: Vec, - pub opening_proof: AdviceFriProof, - pub quotient_opened_values: OpenedValuesForRound, - pub preprocessed_opened_values: Option>, - pub stage_1_opened_values: OpenedValuesForRound, - pub stage_2_opened_values: OpenedValuesForRound, -} - -// --------------------------------------------------------------------------- -// Path read-back -// --------------------------------------------------------------------------- - -/// One step of the per-path verifier's ascent, derived from the matrix -/// dimensions exactly as `verify_batch` derives it. -enum LevelOp { - /// An arity-`step` compression; the path consumes `step - 1` wire - /// siblings here. - Fold { step: usize }, - /// A shorter matrix's row hash is compressed into the running digest; - /// nothing on the wire. - Inject, -} - -/// The ascent schedule (leaf to cap) for a tree of the given dimensions, -/// mirroring `verify_batch`'s traversal: one `Fold` per arity-schedule -/// step, an `Inject` after any step whose folded height picks up shorter -/// matrices. -fn ascent_schedule( - schedule: &[usize], - dimensions: &[Dimensions], - max_height: usize, -) -> Vec { - let mut heights: Vec = dimensions.iter().map(|d| d.height).collect(); - heights.sort_unstable_by(|a, b| b.cmp(a)); - let leaf_npt = max_height.next_power_of_two(); - // Heights hashed into the leaf layer are consumed up front. - let mut next = heights - .iter() - .position(|h| h.next_power_of_two() != leaf_npt) - .unwrap_or(heights.len()); - - let mut ops = Vec::new(); - let mut curr_height_padded = max_height.next_multiple_of(2); - for &step in schedule { - ops.push(LevelOp::Fold { step }); - let logical_next = curr_height_padded / step; - curr_height_padded = logical_next.next_multiple_of(2); - let logical_next_npt = logical_next.next_power_of_two(); - if next < heights.len() && heights[next].next_power_of_two() == logical_next_npt { - ops.push(LevelOp::Inject); - while next < heights.len() && heights[next].next_power_of_two() == logical_next_npt { - next += 1; - } - } - } - ops -} - -/// Reads one query's full authentication path out of the recorded -/// compression map by walking its tree top-down from the cap entry, -/// emitting siblings in the bottom-up order the legacy wire carries. -/// -/// `leaf_check`, when supplied, is the expected leaf digest (the hash of -/// the query's opened rows); the walk must land exactly there. -fn expand_path( - map: &HashMap, - commitment: &Commitment, - schedule: &[usize], - dimensions: &[Dimensions], - max_height: usize, - index: usize, - leaf_check: Option, -) -> Result, AdviceError> { - let ops = ascent_schedule(schedule, dimensions, max_height); - - // Per-fold node index at each level, bottom-up. - let mut level_indices = Vec::new(); - let mut idx = index; - for op in &ops { - if let LevelOp::Fold { step } = op { - level_indices.push(idx); - idx /= step; - } - } - let cap_index = idx; - if cap_index >= commitment.num_roots() { - return Err(AdviceError::CapIndexOutOfBounds); - } - - let mut digest: Digest = commitment[cap_index]; - let mut fold_level = level_indices.len(); - // Wire order is bottom-up; the walk is top-down, so collect reversed. - let mut siblings_rev: Vec = Vec::new(); - for op in ops.iter().rev() { - let inputs = map.get(&digest).ok_or(AdviceError::MissingDigest)?; - match op { - LevelOp::Inject => { - // inject_inputs = [running digest, row hash]; descend left. - digest = inputs[0]; - } - LevelOp::Fold { step } => { - fold_level -= 1; - let pos_in_group = level_indices[fold_level] % step; - // The recording mmcs is binary (N = 2), so `step` is 2. - for k in (0..*step).rev() { - if k != pos_in_group { - siblings_rev.push(inputs[k]); - } - } - digest = inputs[pos_in_group]; - } - } - } - if let Some(expected) = leaf_check - && digest != expected - { - return Err(AdviceError::MissingDigest); - } - siblings_rev.reverse(); - Ok(siblings_rev) -} - -// --------------------------------------------------------------------------- -// Re-encoding -// --------------------------------------------------------------------------- - -/// Converts a natively-verified proof into the per-query advice encoding. -/// -/// `commitment_parameters` and `fri_parameters` must be the ones the -/// system's config was built from (the config does not expose them back). -/// Verifies the proof natively first and refuses to expand on failure. -pub fn proof_to_advice_bytes( - system: &System, - commitment_parameters: CommitmentParameters, - fri_parameters: FriParameters, - claims: &[&[Val]], - proof: &Proof, -) -> Result, AdviceError> { - system - .verify_multiple_claims(claims, proof) - .map_err(|e| AdviceError::Verification(format!("{e:?}")))?; - let ctx = system - .pcs_verification_context(claims, proof) - .map_err(|e| AdviceError::Verification(format!("{e:?}")))?; - - // Instrumented components sharing one compression log. - let log: Arc>> = Arc::new(Mutex::new(HashMap::new())); - let rec_compress = RecordingCompress { - inner: Blake3Compress::new(Blake3), - log: Arc::clone(&log), - }; - let rec_mmcs = RecMmcs::new( - SerializingHasher::new(Blake3), - rec_compress, - commitment_parameters.cap_height, - ); - let rec_params = InnerFriParameters { - log_blowup: commitment_parameters.log_blowup, - log_final_poly_len: fri_parameters.log_final_poly_len, - max_log_arity: fri_parameters.max_log_arity, - num_queries: fri_parameters.num_queries, - commit_proof_of_work_bits: fri_parameters.commit_proof_of_work_bits, - query_proof_of_work_bits: fri_parameters.query_proof_of_work_bits, - mmcs: ExtensionMmcs::::new(rec_mmcs.clone()), - }; - - // The stored FRI proof, re-typed over the instrumented mmcs. The - // commitment, digest and pruned-path types are identical; only the - // phantom mmcs parameter changes. - let fri = &proof.opening_proof; - let rec_proof: FriProof>> = - FriProof { - commit_phase_commits: fri.commit_phase_commits.clone(), - commit_pow_witnesses: fri.commit_pow_witnesses.clone(), - input_openings: fri - .input_openings - .iter() - .map(|o| BatchMultiOpening { - opened_values: o.opened_values.clone(), - opening_proof: o.opening_proof.clone(), - }) - .collect(), - commit_phase_openings: fri - .commit_phase_openings - .iter() - .map(|s| CommitPhaseMultiStep { - log_arity: s.log_arity, - sibling_values: s.sibling_values.clone(), - opening_proof: s.opening_proof.clone(), - }) - .collect(), - final_poly: fri.final_poly.clone(), - query_pow_witness: fri.query_pow_witness, - }; - - // Run the PCS phase exactly as `TwoAdicFriPcs::verify` would — the - // evaluation observations, then `verify_fri` — against the context's - // rounds and transcript state, with the instrumented components. - let mut challenger = RecordingChallenger:: { - inner: ctx.challenger, - samples: Arc::new(Mutex::new(Vec::new())), - }; - for (_, round) in &ctx.rounds { - for (_, mat) in round { - for (_, point) in mat { - challenger.observe_algebra_slice(point); - } - } - } - let folding = TwoAdicFriFolding::< - Vec>, - >::Error, - >(PhantomData); - verify_fri( - &folding, - &rec_params, - &rec_proof, - &mut challenger, - &ctx.rounds, - &rec_mmcs, - ) - .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; - - // The query indices are the last `num_queries` recorded samples. - let samples = challenger.samples.lock().unwrap(); - let query_indices: Vec = samples - .len() - .checked_sub(fri_parameters.num_queries) - .map(|start| samples[start..].to_vec()) - .ok_or(AdviceError::MissingQueryIndices)?; - drop(samples); - let map = log.lock().unwrap(); - - let log_arities: Vec = fri - .commit_phase_openings - .iter() - .map(|s| usize::from(s.log_arity)) - .collect(); - let total_log_reduction: usize = log_arities.iter().sum(); - let log_global_max_height = - total_log_reduction + commitment_parameters.log_blowup + fri_parameters.log_final_poly_len; - - let hasher = SerializingHasher::new(Blake3); - - // Input rounds: per round, per query, opened rows + expanded path. - // Dimensions and reduced indices mirror p3's `open_inputs`. - let mut input_openings_per_query: Vec> = (0..fri_parameters - .num_queries) - .map(|_| Vec::new()) - .collect(); - for ((commit, mats), batch) in ctx.rounds.iter().zip(&fri.input_openings) { - let heights: Vec = mats - .iter() - .map(|(domain, _)| domain.size() << commitment_parameters.log_blowup) - .collect(); - let dims: Vec = heights - .iter() - .zip(mats) - .map(|(&height, (_, points))| Dimensions { - width: points.first().map_or(0, |(_, values)| values.len()), - height, - }) - .collect(); - let max_height = heights.iter().copied().max().unwrap_or(1); - let bits_reduced = log_global_max_height - log2_strict_usize(max_height); - let schedule = rec_mmcs - .proof_arity_schedule(&dims) - .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; - - // Matrices whose padded height matches the tallest are hashed into - // the leaf digest, in tallest-first order — mirroring `verify_batch`. - let leaf_npt = max_height.next_power_of_two(); - let mut order: Vec = (0..dims.len()).collect(); - order.sort_by_key(|&i| std::cmp::Reverse(dims[i].height)); - let leaf_matrices: Vec = order - .iter() - .copied() - .take_while(|&i| dims[i].height.next_power_of_two() == leaf_npt) - .collect(); - - for (q, per_query) in input_openings_per_query.iter_mut().enumerate() { - let reduced_index = query_indices[q] >> bits_reduced; - let opened_values = batch.opened_values[q].clone(); - let leaf = hasher - .hash_iter_slices(leaf_matrices.iter().map(|&mi| opened_values[mi].as_slice())); - let opening_proof = expand_path( - &map, - commit, - &schedule, - &dims, - max_height, - reduced_index, - Some(leaf), - )?; - per_query.push(AdviceBatchOpening { - opened_values, - opening_proof, - }); - } - } - - // Commit-phase rounds: single-matrix trees of the folded codewords, - // flattened to base columns by the extension mmcs. The wire sibling - // values come straight off the multiproof; only the paths are read - // back. No leaf check: the folded row is not materialized here — the - // walk's landing digest is pinned by the functional map and the cap. - let ext_d = >::DIMENSION; - let mut steps_per_query: Vec> = (0..fri_parameters.num_queries) - .map(|_| Vec::new()) - .collect(); - let mut log_current_height = log_global_max_height; - for (round, step) in fri.commit_phase_openings.iter().enumerate() { - let log_arity = usize::from(step.log_arity); - let arity = 1 << log_arity; - let log_folded_height = log_current_height - log_arity; - let dims = [Dimensions { - width: arity * ext_d, - height: 1 << log_folded_height, - }]; - let schedule = rec_mmcs - .proof_arity_schedule(&dims) - .map_err(|e| AdviceError::Recording(format!("{e:?}")))?; - let commit = &fri.commit_phase_commits[round]; - let bits_consumed: usize = log_arities[..=round].iter().sum(); - for (q, per_query) in steps_per_query.iter_mut().enumerate() { - let group_index = query_indices[q] >> bits_consumed; - let opening_proof = expand_path( - &map, - commit, - &schedule, - &dims, - 1 << log_folded_height, - group_index, - None, - )?; - per_query.push(AdviceCommitPhaseStep { - log_arity: step.log_arity, - sibling_values: step.sibling_values[q].clone(), - opening_proof, - }); - } - log_current_height = log_folded_height; - } - - let query_proofs: Vec = input_openings_per_query - .into_iter() - .zip(steps_per_query) - .map(|(input_proof, commit_phase_openings)| AdviceQueryProof { - input_proof, - commit_phase_openings, - }) - .collect(); - - let advice = AdviceProof { - active: proof.active.clone(), - commitments: Commitments { - stage_1_trace: proof.commitments.stage_1_trace.clone(), - stage_2_trace: proof.commitments.stage_2_trace.clone(), - quotient_chunks: proof.commitments.quotient_chunks.clone(), - }, - intermediate_accumulators: proof.intermediate_accumulators.clone(), - log_degrees: proof.log_degrees.clone(), - opening_proof: AdviceFriProof { - commit_phase_commits: fri.commit_phase_commits.clone(), - commit_pow_witnesses: fri.commit_pow_witnesses.clone(), - query_proofs, - final_poly: fri.final_poly.clone(), - query_pow_witness: fri.query_pow_witness, - }, - quotient_opened_values: proof.quotient_opened_values.clone(), - preprocessed_opened_values: proof.preprocessed_opened_values.clone(), - stage_1_opened_values: proof.stage_1_opened_values.clone(), - stage_2_opened_values: proof.stage_2_opened_values.clone(), - }; - encode_to_vec(&advice, Proof::::serde_config()) - .map_err(AdviceError::Encode) -} - -#[cfg(all(test, not(feature = "cuda")))] -mod tests { - use bincode::serde::decode_from_slice; - use p3_field::PrimeCharacteristicRing; - use p3_matrix::dense::RowMajorMatrix; - - use super::*; - use crate::system::ProverKey; - use crate::system::{CircuitInputs, SystemWitness}; - - fn parameters(cap_height: usize) -> (CommitmentParameters, FriParameters) { - ( - CommitmentParameters { - log_blowup: 1, - cap_height, - }, - FriParameters { - log_final_poly_len: 0, - max_log_arity: 1, - num_queries: 8, - commit_proof_of_work_bits: 0, - query_proof_of_work_bits: 0, - }, - ) - } - - /// Two constraint-free circuits at different trace heights (so the - /// stage-1 tree exercises matrix injection), one with a preprocessed - /// matrix (so the preprocessed round exists). - fn test_system( - cap_height: usize, - ) -> ( - System, - ProverKey, - ) { - let (cp, fp) = parameters(cap_height); - let config = GoldilocksBlake3Config::new(cp, fp); - let preprocessed = RowMajorMatrix::new((0..8u32).map(Val::from_u32).collect::>(), 1); - let inputs = [ - CircuitInputs { - main_width: 2, - preprocessed: Some(preprocessed), - ..Default::default() - }, - CircuitInputs { - main_width: 3, - ..Default::default() - }, - ]; - System::new(config, inputs) - } - - fn test_proof( - system: &System, - key: &ProverKey, - ) -> Proof { - let trace_1 = RowMajorMatrix::new((0..16u32).map(Val::from_u32).collect::>(), 2); - let trace_2 = RowMajorMatrix::new( - (0..12u32) - .map(|i| Val::from_u32(7 * i + 3)) - .collect::>(), - 3, - ); - let witness = SystemWitness::from_stage_1(vec![trace_1, trace_2], system); - system.prove_multiple_claims(key, &[], witness) - } - - #[test] - fn advice_expands_and_paths_verify_per_query() { - for cap_height in [0, 1] { - let (cp, fp) = parameters(cap_height); - let (system, key) = test_system(cap_height); - let proof = test_proof(&system, &key); - let bytes = proof_to_advice_bytes(&system, cp, fp, &[], &proof) - .unwrap_or_else(|e| panic!("advice expansion failed: {e:?}")); - - let (advice, consumed): (AdviceProof, usize) = - decode_from_slice(&bytes, Proof::::serde_config()) - .expect("advice bytes decode under the same serde config"); - assert_eq!(consumed, bytes.len(), "no trailing bytes"); - assert_eq!(advice.opening_proof.query_proofs.len(), fp.num_queries); - - // Recover the query indices the same way the expansion did. - let ctx = system.pcs_verification_context(&[], &proof).unwrap(); - let plain_mmcs = MerkleTreeMmcs::::new( - SerializingHasher::new(Blake3), - Blake3Compress::new(Blake3), - cp.cap_height, - ); - let plain_params = InnerFriParameters { - log_blowup: cp.log_blowup, - log_final_poly_len: fp.log_final_poly_len, - max_log_arity: fp.max_log_arity, - num_queries: fp.num_queries, - commit_proof_of_work_bits: fp.commit_proof_of_work_bits, - query_proof_of_work_bits: fp.query_proof_of_work_bits, - mmcs: ExtensionMmcs::::new(plain_mmcs.clone()), - }; - let mut challenger = RecordingChallenger:: { - inner: ctx.challenger, - samples: Arc::new(Mutex::new(Vec::new())), - }; - for (_, round) in &ctx.rounds { - for (_, mat) in round { - for (_, point) in mat { - challenger.observe_algebra_slice(point); - } - } - } - let folding = TwoAdicFriFolding::< - Vec>, - , Blake3Compress, 2, 32> as Mmcs< - Val, - >>::Error, - >(PhantomData); - verify_fri( - &folding, - &plain_params, - &proof.opening_proof, - &mut challenger, - &ctx.rounds, - &plain_mmcs, - ) - .unwrap(); - let samples = challenger.samples.lock().unwrap(); - let indices = &samples[samples.len() - fp.num_queries..]; - - let log_arities: Vec = proof - .opening_proof - .commit_phase_openings - .iter() - .map(|s| usize::from(s.log_arity)) - .collect(); - let log_global_max_height = - log_arities.iter().sum::() + cp.log_blowup + fp.log_final_poly_len; - - // Oracle: every expanded input-round path must satisfy p3's own - // per-path verifier. - for (round, (commit, mats)) in ctx.rounds.iter().enumerate() { - let heights: Vec = mats - .iter() - .map(|(domain, _)| domain.size() << cp.log_blowup) - .collect(); - let dims: Vec = heights - .iter() - .zip(mats) - .map(|(&height, (_, points))| Dimensions { - width: points.first().map_or(0, |(_, values)| values.len()), - height, - }) - .collect(); - let max_height = heights.iter().copied().max().unwrap(); - let bits_reduced = log_global_max_height - log2_strict_usize(max_height); - for (q, &index) in indices.iter().enumerate() { - let opening = &advice.opening_proof.query_proofs[q].input_proof[round]; - plain_mmcs - .verify_batch( - commit, - &dims, - index >> bits_reduced, - p3_commit::BatchOpeningRef::new( - &opening.opened_values, - &opening.opening_proof, - ), - ) - .unwrap_or_else(|e| { - panic!("round {round} query {q}: expanded path rejected: {e:?}") - }); - } - } - - // Commit-phase paths: length and sibling-value parity with the - // multiproof (the Lean verifier is the end-to-end oracle here). - let mut log_current = log_global_max_height; - for (round, step) in proof.opening_proof.commit_phase_openings.iter().enumerate() { - let log_arity = usize::from(step.log_arity); - let log_folded = log_current - log_arity; - let expected_levels = log_folded.saturating_sub(cp.cap_height); - for (q, qp) in advice.opening_proof.query_proofs.iter().enumerate() { - let s = &qp.commit_phase_openings[round]; - assert_eq!(usize::from(s.log_arity), log_arity); - assert_eq!(s.sibling_values, step.sibling_values[q]); - assert_eq!( - s.opening_proof.len(), - expected_levels, - "round {round} query {q}: phase path length" - ); - } - log_current = log_folded; - } - } - } - - #[test] - fn advice_refuses_invalid_proof() { - let cap_height = 0; - let (cp, fp) = parameters(cap_height); - let (system, key) = test_system(cap_height); - let mut proof = test_proof(&system, &key); - proof.log_degrees[0] ^= 1; - assert!(matches!( - proof_to_advice_bytes(&system, cp, fp, &[], &proof), - Err(AdviceError::Verification(_)) - )); - } -} diff --git a/src/lib.rs b/src/lib.rs index 3f18e65..a3afd0a 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1,4 +1,3 @@ -pub mod advice; pub mod config; #[cfg(feature = "cuda")] pub mod cuda; diff --git a/src/prover.rs b/src/prover.rs index 5276604..57a574c 100644 --- a/src/prover.rs +++ b/src/prover.rs @@ -255,7 +255,7 @@ impl Clone for Proof { } impl Proof { - pub(crate) fn serde_config() -> Configuration { + fn serde_config() -> Configuration { standard().with_little_endian().with_fixed_int_encoding() } diff --git a/src/verifier.rs b/src/verifier.rs index bf1dc3f..a99af01 100644 --- a/src/verifier.rs +++ b/src/verifier.rs @@ -159,7 +159,7 @@ //! actual low-degree-extension values of the witness. Do not use it when the //! witness must remain hidden from the verifier. -use crate::config::{Com, Domain, PcsError, StarkGenericConfig, Val}; +use crate::config::{PcsError, StarkGenericConfig, Val}; use crate::ensure_eq; use crate::eval::VarValues; use crate::lookup::fingerprint; @@ -191,29 +191,6 @@ pub enum VerificationError { UnbalancedChannel, } -pub(crate) type OpeningRounds = Vec<( - Com, - Vec<( - Domain, - Vec<( - ::Challenge, - Vec<::Challenge>, - )>, - )>, -)>; - -pub(crate) struct PcsVerificationContext { - pub rounds: OpeningRounds, - pub challenger: SC::Challenger, - pub zeta: SC::Challenge, - pub lookup_argument_challenge: SC::Challenge, - pub fingerprint_challenge: SC::Challenge, - pub constraint_challenge: SC::Challenge, - pub claim_accumulator: SC::Challenge, - pub active_indices: Vec, - pub quotient_degrees: Vec, -} - impl System { /// Verifies a STARK proof against a single claim. pub fn verify( @@ -227,11 +204,12 @@ impl System { self.verify_multiple_claims(&[claim], proof) } - pub(crate) fn pcs_verification_context( + /// Verifies a STARK proof against multiple claims. + pub fn verify_multiple_claims( &self, claims: &[&[Val]], proof: &Proof, - ) -> Result, VerificationError>> + ) -> Result<(), VerificationError>> where Val: TwoAdicField, { @@ -240,7 +218,7 @@ impl System { commitments, intermediate_accumulators, log_degrees, - opening_proof: _, + opening_proof, quotient_opened_values, preprocessed_opened_values, stage_1_opened_values, @@ -428,57 +406,11 @@ impl System { preprocessed_trace_evaluations, )); } - Ok(PcsVerificationContext { - rounds: coms_to_verify, - challenger, - zeta, - lookup_argument_challenge, - fingerprint_challenge, - constraint_challenge, - claim_accumulator: acc, - active_indices, - quotient_degrees, - }) - } - - /// Verifies a STARK proof against multiple claims. - pub fn verify_multiple_claims( - &self, - claims: &[&[Val]], - proof: &Proof, - ) -> Result<(), VerificationError>> - where - Val: TwoAdicField, - { - let Proof { - intermediate_accumulators, - log_degrees, - opening_proof, - quotient_opened_values, - preprocessed_opened_values, - stage_1_opened_values, - stage_2_opened_values, - .. - } = proof; - let PcsVerificationContext { - rounds, - mut challenger, - zeta, - lookup_argument_challenge, - fingerprint_challenge, - constraint_challenge, - claim_accumulator, - active_indices, - quotient_degrees, - } = self.pcs_verification_context(claims, proof)?; - let mut acc = claim_accumulator; - let pcs = self.config.pcs(); - // Soundness: FRI proximity test. Verifies that the committed polynomials // are close to low-degree polynomials and that the claimed evaluations are // consistent with the commitments. Soundness error ≤ ρ^num_queries, where // ρ = 2^(-log_blowup). This is the dominant term in the overall bound. - pcs.verify(rounds, opening_proof, &mut challenger) + pcs.verify(coms_to_verify, opening_proof, &mut challenger) .map_err(VerificationError::InvalidOpeningArgument)?; // use the opened values to compute the composition polynomial for each circuit