diff --git a/.github/workflows/continuous-integration-workflow.yml b/.github/workflows/continuous-integration-workflow.yml
index f38bba611..4b51661d1 100644
--- a/.github/workflows/continuous-integration-workflow.yml
+++ b/.github/workflows/continuous-integration-workflow.yml
@@ -64,9 +64,8 @@ jobs:
         uses: actions/checkout@v4
 
       - name: Install nightly toolchain
-        uses: dtolnay/rust-toolchain@master
+        uses: dtolnay/rust-toolchain@nightly
         with:
-          toolchain: nightly-2024-06-23
           components: rustfmt, clippy
 
       - name: Rust Cache
diff --git a/field/src/goldilocks_field.rs b/field/src/goldilocks_field.rs
index b0191ca59..a4da1fee0 100644
--- a/field/src/goldilocks_field.rs
+++ b/field/src/goldilocks_field.rs
@@ -160,6 +160,11 @@ impl Field for GoldilocksField {
         reduce96((n_lo, n_hi))
     }
 
+    fn from_noncanonical_u128_with_96_bits(n: u128) -> Self {
+        debug_assert!(n < (1u128 << 96));
+        reduce128_with_96_bits(n)
+    }
+
     fn from_noncanonical_u128(n: u128) -> Self {
         reduce128(n)
     }
@@ -396,6 +401,14 @@ fn reduce96((x_lo, x_hi): (u64, u32)) -> GoldilocksField {
     GoldilocksField(t2)
 }
 
+#[inline]
+fn reduce128_with_96_bits(x: u128) -> GoldilocksField {
+    let (x_lo, x_hi) = split(x); // This is a no-op
+    let t1 = x_hi * EPSILON;
+    let t2 = unsafe { add_no_canonicalize_trashing_input(x_lo, t1) };
+    GoldilocksField(t2)
+}
+
 /// Reduces to a 64-bit value. The result might not be in canonical form; it could be in between the
 /// field order and `2^64`.
 #[inline]
diff --git a/field/src/lib.rs b/field/src/lib.rs
index c713db885..9a2ea4f9c 100644
--- a/field/src/lib.rs
+++ b/field/src/lib.rs
@@ -4,7 +4,6 @@
 #![deny(rustdoc::broken_intra_doc_links)]
 #![deny(missing_debug_implementations)]
 #![feature(specialization)]
-#![cfg_attr(target_arch = "x86_64", feature(stdarch_x86_avx512))]
 #![cfg_attr(not(test), no_std)]
 
 extern crate alloc;
diff --git a/field/src/types.rs b/field/src/types.rs
index d714b7a84..153231c9a 100644
--- a/field/src/types.rs
+++ b/field/src/types.rs
@@ -356,6 +356,12 @@ pub trait Field:
         Self::from_noncanonical_u128(n)
     }
 
+    fn from_noncanonical_u128_with_96_bits(n: u128) -> Self {
+        // Default implementation.
+        debug_assert!(n < (1u128 << 96));
+        Self::from_noncanonical_u128(n)
+    }
+
     fn exp_power_of_2(&self, power_log: usize) -> Self {
         let mut res = *self;
         for _ in 0..power_log {
diff --git a/plonky2/Cargo.toml b/plonky2/Cargo.toml
index 83ff08519..36f26fb79 100644
--- a/plonky2/Cargo.toml
+++ b/plonky2/Cargo.toml
@@ -38,6 +38,11 @@ plonky2_field = { version = "1.0.0", path = "../field", default-features = false
 plonky2_maybe_rayon = { version = "1.0.0", path = "../maybe_rayon", default-features = false }
 plonky2_util = { version = "1.0.0", path = "../util", default-features = false }
 
+# Plonky3 dependencies for Poseidon2
+p3-poseidon2 = { git = "https://github.com/Plonky3/Plonky3.git", rev = "eeb4e37b20127c4daa871b2bad0df30a7c7380db" }
+p3-goldilocks = { git = "https://github.com/Plonky3/Plonky3.git", rev = "eeb4e37b20127c4daa871b2bad0df30a7c7380db" }
+p3-field = { git = "https://github.com/Plonky3/Plonky3.git", rev = "eeb4e37b20127c4daa871b2bad0df30a7c7380db" }
+p3-symmetric = { git = "https://github.com/Plonky3/Plonky3.git", rev = "eeb4e37b20127c4daa871b2bad0df30a7c7380db" }
 
 [target.'cfg(all(target_arch = "wasm32", target_os = "unknown"))'.dependencies]
 getrandom = { version = "0.2", default-features = false, features = ["js"] }
diff --git a/plonky2/benches/hashing.rs b/plonky2/benches/hashing.rs
index cf64b764a..05e8731cb 100644
--- a/plonky2/benches/hashing.rs
+++ b/plonky2/benches/hashing.rs
@@ -1,11 +1,16 @@
 mod allocator;
 
 use criterion::{criterion_group, criterion_main, BatchSize, Criterion};
+use p3_field::AbstractField;
+use p3_goldilocks::{DiffusionMatrixGoldilocks, Goldilocks};
+use p3_poseidon2::{Poseidon2 as P3Poseidon2, Poseidon2ExternalMatrixGeneral};
+use p3_symmetric::Permutation;
 use plonky2::field::goldilocks_field::GoldilocksField;
 use plonky2::field::types::Sample;
 use plonky2::hash::hash_types::{BytesHash, RichField};
 use plonky2::hash::keccak::KeccakHash;
 use plonky2::hash::poseidon::{Poseidon, SPONGE_WIDTH};
+use plonky2::hash::poseidon2::hash::Poseidon2;
 use plonky2::plonk::config::Hasher;
 use tynm::type_name;
 
@@ -32,8 +37,82 @@ pub(crate) fn bench_poseidon<F: Poseidon>(c: &mut Criterion) {
     );
 }
 
+pub(crate) fn bench_poseidon2<F: Poseidon2>(c: &mut Criterion) {
+    c.bench_function(
+        &format!("optimized poseidon2<{}, {SPONGE_WIDTH}>", type_name::<F>()),
+        |b| {
+            b.iter_batched(
+                || F::rand_array::<SPONGE_WIDTH>(),
+                |state| F::poseidon2(state),
+                BatchSize::SmallInput,
+            )
+        },
+    );
+}
+
+pub(crate) fn bench_p3_poseidon2(c: &mut Criterion) {
+    const WIDTH: usize = 12;
+    const D: u64 = 7;
+    const ROUNDS_F: usize = 8;
+    const ROUNDS_P: usize = 22;
+
+    // Create the Poseidon2 instance
+    let external_linear_layer = Poseidon2ExternalMatrixGeneral;
+    let internal_linear_layer = DiffusionMatrixGoldilocks;
+
+    let external_constants = plonky2::hash::poseidon2::config::EXTERNAL_CONSTANTS
+        .iter()
+        .map(|v| {
+            v.iter()
+                .map(|&x| Goldilocks::from_canonical_u64(x))
+                .collect::<Vec<Goldilocks>>()
+                .try_into()
+                .unwrap()
+        })
+        .collect::<Vec<[Goldilocks; WIDTH]>>();
+
+    let internal_constants = plonky2::hash::poseidon2::config::INTERNAL_CONSTANTS
+        .iter()
+        .map(|&x| Goldilocks::from_canonical_u64(x))
+        .collect::<Vec<Goldilocks>>();
+
+    let poseidon = P3Poseidon2::<
+        Goldilocks,
+        Poseidon2ExternalMatrixGeneral,
+        DiffusionMatrixGoldilocks,
+        WIDTH,
+        D,
+    >::new(
+        ROUNDS_F,
+        external_constants,
+        external_linear_layer,
+        ROUNDS_P,
+        internal_constants,
+        internal_linear_layer,
+    );
+
+    c.bench_function("plonky3's poseidon2", |b| {
+        b.iter_batched(
+            || {
+                let mut state = [Goldilocks::zero(); WIDTH];
+                state.iter_mut().for_each(|item| {
+                    *item = Goldilocks::from_canonical_u64(rand::random::<u64>());
+                });
+                state
+            },
+            |mut state| {
+                poseidon.permute_mut(&mut state);
+                state
+            },
+            BatchSize::SmallInput,
+        )
+    });
+}
+
 fn criterion_benchmark(c: &mut Criterion) {
     bench_poseidon::<GoldilocksField>(c);
+    bench_poseidon2::<GoldilocksField>(c);
+    bench_p3_poseidon2(c);
     bench_keccak::<GoldilocksField>(c);
 }
 
diff --git a/plonky2/examples/fibonacci.rs b/plonky2/examples/fibonacci.rs
index 578dc2424..be491aa21 100644
--- a/plonky2/examples/fibonacci.rs
+++ b/plonky2/examples/fibonacci.rs
@@ -1,27 +1,38 @@
+use std::time::Instant;
+
 use anyhow::Result;
+use log::Level;
 use plonky2::field::types::Field;
 use plonky2::iop::witness::{PartialWitness, WitnessWrite};
 use plonky2::plonk::circuit_builder::CircuitBuilder;
 use plonky2::plonk::circuit_data::CircuitConfig;
-use plonky2::plonk::config::{GenericConfig, PoseidonGoldilocksConfig};
+use plonky2::plonk::config::{GenericConfig, Poseidon2GoldilocksConfig, PoseidonGoldilocksConfig};
+use plonky2::plonk::prover::prove;
+use plonky2::util::timing::TimingTree;
 
 /// An example of using Plonky2 to prove a statement of the form
 /// "I know the 100th element of the Fibonacci sequence, starting with constants a and b."
 /// When a == 0 and b == 1, this is proving knowledge of the 100th (standard) Fibonacci number.
 fn main() -> Result<()> {
+    env_logger::Builder::from_default_env()
+        .filter_level(log::LevelFilter::Debug)
+        .init();
+    work::<PoseidonGoldilocksConfig>()?;
+    work::<Poseidon2GoldilocksConfig>()
+}
+
+fn work<C: GenericConfig<2>>() -> Result<()> {
     const D: usize = 2;
-    type C = PoseidonGoldilocksConfig;
-    type F = <C as GenericConfig<D>>::F;
 
     let config = CircuitConfig::standard_recursion_config();
-    let mut builder = CircuitBuilder::<F, D>::new(config);
+    let mut builder = CircuitBuilder::<C::F, D>::new(config);
 
     // The arithmetic circuit.
     let initial_a = builder.add_virtual_target();
     let initial_b = builder.add_virtual_target();
     let mut prev_target = initial_a;
     let mut cur_target = initial_b;
-    for _ in 0..99 {
+    for _ in 0..999999 {
         let temp = builder.add(prev_target, cur_target);
         prev_target = cur_target;
         cur_target = temp;
@@ -33,17 +44,29 @@ fn main() -> Result<()> {
     builder.register_public_input(cur_target);
 
     // Provide initial values.
+    let timer1 = Instant::now();
     let mut pw = PartialWitness::new();
-    pw.set_target(initial_a, F::ZERO)?;
-    pw.set_target(initial_b, F::ONE)?;
+    pw.set_target(initial_a, C::F::ZERO)?;
+    pw.set_target(initial_b, C::F::ONE)?;
 
     let data = builder.build::<C>();
-    let proof = data.prove(pw)?;
+    let timer2 = Instant::now();
+
+    // Create a TimingTree to track detailed timing information
+    let mut timing = TimingTree::new("prove", Level::Debug);
+    let proof = prove::<C::F, C, D>(&data.prover_only, &data.common, pw, &mut timing)?;
+    let timer3 = Instant::now();
+
+    // Print the timing tree
+    timing.print();
 
     println!(
         "100th Fibonacci number mod |F| (starting with {}, {}) is: {}",
         proof.public_inputs[0], proof.public_inputs[1], proof.public_inputs[2]
     );
 
+    println!("Build time: {:?}", timer2.duration_since(timer1));
+    println!("Prove time: {:?}", timer3.duration_since(timer2));
+
     data.verify(proof)
 }
diff --git a/plonky2/src/batch_fri/oracle.rs b/plonky2/src/batch_fri/oracle.rs
index 58deeaa3c..1f34cfef3 100644
--- a/plonky2/src/batch_fri/oracle.rs
+++ b/plonky2/src/batch_fri/oracle.rs
@@ -460,7 +460,7 @@ mod test {
             &fri_instances,
             &fri_openings,
             &fri_challenges,
-            &[merkle_cap.clone()],
+            std::slice::from_ref(&merkle_cap),
             &proof,
             &fri_params,
         )?;
diff --git a/plonky2/src/fri/validate_shape.rs b/plonky2/src/fri/validate_shape.rs
index be675ed61..8805478f6 100644
--- a/plonky2/src/fri/validate_shape.rs
+++ b/plonky2/src/fri/validate_shape.rs
@@ -20,7 +20,7 @@ where
     F: RichField + Extendable<D>,
     C: GenericConfig<D, F = F>,
 {
-    validate_batch_fri_proof_shape::<F, C, D>(proof, &[instance.clone()], params)
+    validate_batch_fri_proof_shape::<F, C, D>(proof, std::slice::from_ref(instance), params)
 }
 
 pub(crate) fn validate_batch_fri_proof_shape<F, C, const D: usize>(
diff --git a/plonky2/src/gadgets/interpolation.rs b/plonky2/src/gadgets/interpolation.rs
index 39b048af4..9aedf7143 100644
--- a/plonky2/src/gadgets/interpolation.rs
+++ b/plonky2/src/gadgets/interpolation.rs
@@ -86,7 +86,7 @@ mod tests {
 
         let value_targets = values
             .iter()
-            .map(|&v| (builder.constant_extension(v)))
+            .map(|&v| builder.constant_extension(v))
             .collect::<Vec<_>>();
 
         let zt = builder.constant_extension(z);
diff --git a/plonky2/src/hash/arch/aarch64/poseidon_goldilocks_neon.rs b/plonky2/src/hash/arch/aarch64/poseidon_goldilocks_neon.rs
index a9328d069..17078ff99 100644
--- a/plonky2/src/hash/arch/aarch64/poseidon_goldilocks_neon.rs
+++ b/plonky2/src/hash/arch/aarch64/poseidon_goldilocks_neon.rs
@@ -58,7 +58,7 @@ const_assert!(check_mds_matrix());
 
 /// Addition modulo ORDER accounting for wraparound. Correct only when a + b < 2**64 + ORDER.
 #[inline(always)]
-unsafe fn add_with_wraparound(a: u64, b: u64) -> u64 {
+pub unsafe fn add_with_wraparound(a: u64, b: u64) -> u64 {
     let res: u64;
     let adj: u64;
     asm!(
@@ -872,12 +872,12 @@ unsafe fn partial_rounds(
 */
 
 #[inline(always)]
-fn unwrap_state(state: [GoldilocksField; 12]) -> [u64; 12] {
+pub(crate) fn unwrap_state(state: [GoldilocksField; 12]) -> [u64; 12] {
     state.map(|s| s.0)
 }
 
 #[inline(always)]
-fn wrap_state(state: [u64; 12]) -> [GoldilocksField; 12] {
+pub(crate) fn wrap_state(state: [u64; 12]) -> [GoldilocksField; 12] {
     state.map(GoldilocksField)
 }
 
@@ -915,3 +915,28 @@ pub unsafe fn mds_layer(state: &[GoldilocksField; WIDTH]) -> [GoldilocksField; W
     let state = mds_layer_full(state);
     wrap_state(state)
 }
+
+#[inline(always)]
+#[unroll::unroll_for_loops]
+pub unsafe fn vector_add(a: &[u64; WIDTH], b: &[u64; WIDTH]) -> [u64; WIDTH] {
+    let mut res = [0u64; WIDTH];
+    // Process 2 elements at a time using NEON
+    for i in (0..WIDTH).step_by(2) {
+        let a_vec = vld1q_u64(a[i..].as_ptr());
+        let b_vec = vld1q_u64(b[i..].as_ptr());
+
+        // Add the round constants
+        let sum = vaddq_u64(a_vec, b_vec);
+
+        // Check for overflow (if sum < state_vec, we wrapped around)
+        let overflow_mask = vcltq_u64(sum, a_vec);
+
+        // Add EPSILON (0xffffffff) where overflow occurred
+        let epsilon = vdupq_n_u64(0xffffffff);
+        let adjustment = vandq_u64(overflow_mask, epsilon);
+        let result = vaddq_u64(sum, adjustment);
+
+        vst1q_u64(res[i..].as_mut_ptr(), result);
+    }
+    res
+}
diff --git a/plonky2/src/hash/merkle_proofs.rs b/plonky2/src/hash/merkle_proofs.rs
index 424e03ae6..74963a2a4 100644
--- a/plonky2/src/hash/merkle_proofs.rs
+++ b/plonky2/src/hash/merkle_proofs.rs
@@ -59,7 +59,7 @@ pub fn verify_merkle_proof_to_cap<F: RichField, H: Hasher<F>>(
     proof: &MerkleProof<F, H>,
 ) -> Result<()> {
     verify_batch_merkle_proof_to_cap(
-        &[leaf_data.clone()],
+        &[leaf_data],
         &[proof.siblings.len()],
         leaf_index,
         merkle_cap,
diff --git a/plonky2/src/hash/mod.rs b/plonky2/src/hash/mod.rs
index 0e4bb8a59..0d3436973 100644
--- a/plonky2/src/hash/mod.rs
+++ b/plonky2/src/hash/mod.rs
@@ -10,4 +10,5 @@ pub mod merkle_proofs;
 pub mod merkle_tree;
 pub mod path_compression;
 pub mod poseidon;
+pub mod poseidon2;
 pub mod poseidon_goldilocks;
diff --git a/plonky2/src/hash/poseidon2/config.rs b/plonky2/src/hash/poseidon2/config.rs
new file mode 100644
index 000000000..96f25099b
--- /dev/null
+++ b/plonky2/src/hash/poseidon2/config.rs
@@ -0,0 +1,167 @@
+pub const WIDTH: usize = 12;
+pub const D: u64 = 7;
+pub const RATE: usize = 8;
+pub const OUT: usize = 4;
+
+/// Generated by `poseidon2_round_numbers_128`
+pub const ROUNDS_F: usize = 8;
+pub const ROUNDS_F_HALF: usize = 4;
+pub const ROUNDS_P: usize = 22;
+
+/// Generated randomly for ROUNDS_F
+pub const EXTERNAL_CONSTANTS: [[u64; WIDTH]; ROUNDS_F] = [
+    [
+        15492826721047263190,
+        11728330187201910315,
+        8836021247773420868,
+        16777404051263952451,
+        5510875212538051896,
+        6173089941271892285,
+        2927757366422211339,
+        10340958981325008808,
+        8541987352684552425,
+        9739599543776434497,
+        15073950188101532019,
+        12084856431752384512,
+    ],
+    [
+        4584713381960671270,
+        8807052963476652830,
+        54136601502601741,
+        4872702333905478703,
+        5551030319979516287,
+        12889366755535460989,
+        16329242193178844328,
+        412018088475211848,
+        10505784623379650541,
+        9758812378619434837,
+        7421979329386275117,
+        375240370024755551,
+    ],
+    [
+        3331431125640721931,
+        15684937309956309981,
+        578521833432107983,
+        14379242000670861838,
+        17922409828154900976,
+        8153494278429192257,
+        15904673920630731971,
+        11217863998460634216,
+        3301540195510742136,
+        9937973023749922003,
+        3059102938155026419,
+        1895288289490976132,
+    ],
+    [
+        5580912693628927540,
+        10064804080494788323,
+        9582481583369602410,
+        10186259561546797986,
+        247426333829703916,
+        13193193905461376067,
+        6386232593701758044,
+        17954717245501896472,
+        1531720443376282699,
+        2455761864255501970,
+        11234429217864304495,
+        4746959618548874102,
+    ],
+    [
+        13571697342473846203,
+        17477857865056504753,
+        15963032953523553760,
+        16033593225279635898,
+        14252634232868282405,
+        8219748254835277737,
+        7459165569491914711,
+        15855939513193752003,
+        16788866461340278896,
+        7102224659693946577,
+        3024718005636976471,
+        13695468978618890430,
+    ],
+    [
+        8214202050877825436,
+        2670727992739346204,
+        16259532062589659211,
+        11869922396257088411,
+        3179482916972760137,
+        13525476046633427808,
+        3217337278042947412,
+        14494689598654046340,
+        15837379330312175383,
+        8029037639801151344,
+        2153456285263517937,
+        8301106462311849241,
+    ],
+    [
+        13294194396455217955,
+        17394768489610594315,
+        12847609130464867455,
+        14015739446356528640,
+        5879251655839607853,
+        9747000124977436185,
+        8950393546890284269,
+        10765765936405694368,
+        14695323910334139959,
+        16366254691123000864,
+        15292774414889043182,
+        10910394433429313384,
+    ],
+    [
+        17253424460214596184,
+        3442854447664030446,
+        3005570425335613727,
+        10859158614900201063,
+        9763230642109343539,
+        6647722546511515039,
+        909012944955815706,
+        18101204076790399111,
+        11588128829349125809,
+        15863878496612806566,
+        5201119062417750399,
+        176665553780565743,
+    ],
+];
+
+/// Generated randomly for ROUNDS_P
+pub const INTERNAL_CONSTANTS: [u64; ROUNDS_P] = [
+    11921381764981422944,
+    10318423381711320787,
+    8291411502347000766,
+    229948027109387563,
+    9152521390190983261,
+    7129306032690285515,
+    15395989607365232011,
+    8641397269074305925,
+    17256848792241043600,
+    6046475228902245682,
+    12041608676381094092,
+    12785542378683951657,
+    14546032085337914034,
+    3304199118235116851,
+    16499627707072547655,
+    10386478025625759321,
+    13475579315436919170,
+    16042710511297532028,
+    1411266850385657080,
+    9024840976168649958,
+    14047056970978379368,
+    838728605080212101,
+];
+
+/// Taken from Plonk3 Poseidon2 implementation. https://github.com/Plonky3/Plonky3/blob/eeb4e37b20127c4daa871b2bad0df30a7c7380db/goldilocks/src/poseidon2.rs#L28
+pub const MATRIX_DIAG_12_U64: [u64; WIDTH] = [
+    0xc3b6c08e23ba9300,
+    0xd84b5de94a324fb6,
+    0x0d0c371c5b35b84f,
+    0x7964f570e7188037,
+    0x5daf18bbd996604b,
+    0x6743bc47b9595257,
+    0x5528b9362c59bb70,
+    0xac45e25b7127b68b,
+    0xa2077d7dfbb606b5,
+    0xf3faac6faee378ae,
+    0x0c6388b51545e883,
+    0xd27dbb6944917b60,
+];
diff --git a/plonky2/src/hash/poseidon2/gate.rs b/plonky2/src/hash/poseidon2/gate.rs
new file mode 100644
index 000000000..ff53335e2
--- /dev/null
+++ b/plonky2/src/hash/poseidon2/gate.rs
@@ -0,0 +1,602 @@
+//! Implementation of a Plonky2 gate for an entire Poseidon2 permutation over a
+//! state of width 12
+use core::marker::PhantomData;
+
+use anyhow::Result;
+
+use super::config::*;
+use super::hash::Poseidon2;
+use crate::field::extension::Extendable;
+use crate::field::types::Field;
+use crate::gates::gate::Gate;
+use crate::gates::util::StridedConstraintConsumer;
+use crate::hash::hash_types::RichField;
+use crate::iop::ext_target::ExtensionTarget;
+use crate::iop::generator::{GeneratedValues, SimpleGenerator, WitnessGeneratorRef};
+use crate::iop::target::Target;
+use crate::iop::wire::Wire;
+use crate::iop::witness::{PartitionWitness, Witness, WitnessWrite};
+use crate::plonk::circuit_builder::CircuitBuilder;
+use crate::plonk::circuit_data::CommonCircuitData;
+use crate::plonk::vars::{EvaluationTargets, EvaluationVars, EvaluationVarsBase};
+use crate::util::serialization::{Buffer, IoResult, Read, Write};
+
+/// Evaluates a full Poseidon2 permutation with 12 state elements.
+///
+/// This also has some extra features to make it suitable for efficiently
+/// verifying Merkle proofs. It has a flag which can be used to swap the first
+/// four inputs with the next four, for ordering sibling digests.
+#[derive(Debug, Default)]
+pub struct Poseidon2Gate<F: RichField + Extendable<D>, const D: usize> {
+    _phantom: PhantomData<F>,
+}
+
+impl<F: RichField + Extendable<D> + Poseidon2, const D: usize> Poseidon2Gate<F, D> {
+    pub fn new() -> Self {
+        Poseidon2Gate {
+            _phantom: PhantomData,
+        }
+    }
+
+    /// The wire index for the `i`th input to the permutation.
+    pub fn wire_input(i: usize) -> usize {
+        i
+    }
+
+    /// The wire index for the `i`th output to the permutation.
+    pub fn wire_output(i: usize) -> usize {
+        WIDTH + i
+    }
+
+    /// If this is set to 1, the first four inputs will be swapped with the next
+    /// four inputs. This is useful for ordering hashes in Merkle proofs.
+    /// Otherwise, this should be set to 0.
+    pub const WIRE_SWAP: usize = 2 * WIDTH;
+
+    const START_DELTA: usize = 2 * WIDTH + 1;
+
+    /// A wire which stores `swap * (input[i + 4] - input[i])`; used to compute
+    /// the swapped inputs.
+    fn wire_delta(i: usize) -> usize {
+        assert!(i < 4);
+        Self::START_DELTA + i
+    }
+
+    const START_ROUND_F_BEGIN: usize = Self::START_DELTA + 4;
+
+    /// A wire which stores the input of the `i`-th S-box of the `round`-th
+    /// round of the first set of full rounds.
+    fn wire_full_sbox_0(round: usize, i: usize) -> usize {
+        debug_assert!(
+            round != 0,
+            "First round S-box inputs are not stored as wires"
+        );
+        debug_assert!(round < ROUNDS_F_HALF);
+        debug_assert!(i < WIDTH);
+        Self::START_ROUND_F_BEGIN + WIDTH * (round - 1) + i
+    }
+
+    const START_PARTIAL: usize = Self::START_ROUND_F_BEGIN + WIDTH * (ROUNDS_F_HALF - 1);
+
+    /// A wire which stores the input of the S-box of the `round`-th round of
+    /// the partial rounds.
+    const fn wire_partial_sbox(round: usize) -> usize {
+        debug_assert!(round < ROUNDS_P);
+        Self::START_PARTIAL + round
+    }
+
+    const START_ROUND_F_END: usize = Self::START_PARTIAL + ROUNDS_P;
+
+    /// A wire which stores the input of the `i`-th S-box of the `round`-th
+    /// round of the second set of full rounds.
+    const fn wire_full_sbox_1(round: usize, i: usize) -> usize {
+        debug_assert!(round < ROUNDS_F_HALF);
+        debug_assert!(i < WIDTH);
+        Self::START_ROUND_F_END + WIDTH * round + i
+    }
+
+    /// End of wire indices, exclusive.
+    const fn end() -> usize {
+        Self::START_ROUND_F_END + WIDTH * ROUNDS_F_HALF
+    }
+}
+
+impl<F: RichField + Extendable<D> + Poseidon2, const D: usize> Gate<F, D> for Poseidon2Gate<F, D> {
+    fn id(&self) -> String {
+        format!("{:?}<WIDTH={}>", self, WIDTH)
+    }
+
+    fn serialize(
+        &self,
+        _dst: &mut Vec<u8>,
+        _common_data: &CommonCircuitData<F, D>,
+    ) -> IoResult<()> {
+        Ok(())
+    }
+
+    fn deserialize(_src: &mut Buffer, _common_data: &CommonCircuitData<F, D>) -> IoResult<Self> {
+        Ok(Poseidon2Gate::new())
+    }
+
+    fn eval_unfiltered(&self, vars: EvaluationVars<F, D>) -> Vec<F::Extension> {
+        let mut constraints = Vec::with_capacity(self.num_constraints());
+
+        // Assert that `swap` is binary.
+        let swap = vars.local_wires[Self::WIRE_SWAP];
+        constraints.push(swap * (swap - F::Extension::ONE));
+
+        // Assert that each delta wire is set properly: `delta_i = swap * (rhs - lhs)`.
+        for i in 0..4 {
+            let input_lhs = vars.local_wires[Self::wire_input(i)];
+            let input_rhs = vars.local_wires[Self::wire_input(i + 4)];
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            constraints.push(swap * (input_rhs - input_lhs) - delta_i);
+        }
+
+        // Compute the possibly-swapped input layer.
+        let mut state = [F::Extension::ZERO; WIDTH];
+        for i in 0..4 {
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            let input_lhs = Self::wire_input(i);
+            let input_rhs = Self::wire_input(i + 4);
+            state[i] = vars.local_wires[input_lhs] + delta_i;
+            state[i + 4] = vars.local_wires[input_rhs] - delta_i;
+        }
+        for i in 8..WIDTH {
+            state[i] = vars.local_wires[Self::wire_input(i)];
+        }
+
+        // The initial linear layer.
+        <F as Poseidon2>::external_linear_layer_extension(&mut state);
+
+        // The first half of the external rounds.
+        for r in 0..ROUNDS_F_HALF {
+            <F as Poseidon2>::add_rc_extension(&mut state, r);
+            if r != 0 {
+                for i in 0..WIDTH {
+                    let sbox_in = vars.local_wires[Self::wire_full_sbox_0(r, i)];
+                    constraints.push(state[i] - sbox_in);
+                    state[i] = sbox_in;
+                }
+            }
+            <F as Poseidon2>::sbox_extension(&mut state);
+            <F as Poseidon2>::external_linear_layer_extension(&mut state);
+        }
+
+        // The internal rounds.
+        for r in 0..ROUNDS_P {
+            state[0] += F::Extension::from_canonical_u64(INTERNAL_CONSTANTS[r]);
+            let sbox_in = vars.local_wires[Self::wire_partial_sbox(r)];
+            constraints.push(state[0] - sbox_in);
+            state[0] = sbox_in;
+            state[0] = <F as Poseidon2>::sbox_p_extension(&state[0]);
+            <F as Poseidon2>::internal_linear_layer_extension(&mut state);
+        }
+
+        // The second half of the external rounds.
+        for r in ROUNDS_F_HALF..ROUNDS_F {
+            <F as Poseidon2>::add_rc_extension(&mut state, r);
+            for i in 0..WIDTH {
+                let sbox_in = vars.local_wires[Self::wire_full_sbox_1(r - ROUNDS_F_HALF, i)];
+                constraints.push(state[i] - sbox_in);
+                state[i] = sbox_in;
+            }
+            <F as Poseidon2>::sbox_extension(&mut state);
+            <F as Poseidon2>::external_linear_layer_extension(&mut state);
+        }
+
+        for i in 0..WIDTH {
+            constraints.push(state[i] - vars.local_wires[Self::wire_output(i)]);
+        }
+
+        constraints
+    }
+
+    fn eval_unfiltered_base_one(
+        &self,
+        vars: EvaluationVarsBase<F>,
+        mut yield_constr: StridedConstraintConsumer<F>,
+    ) {
+        // Assert that `swap` is binary.
+        let swap = vars.local_wires[Self::WIRE_SWAP];
+        yield_constr.one(swap * swap.sub_one());
+
+        // Assert that each delta wire is set properly: `delta_i = swap * (rhs - lhs)`.
+        for i in 0..4 {
+            let input_lhs = vars.local_wires[Self::wire_input(i)];
+            let input_rhs = vars.local_wires[Self::wire_input(i + 4)];
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            yield_constr.one(swap * (input_rhs - input_lhs) - delta_i);
+        }
+
+        // Compute the possibly-swapped input layer.
+        let mut state = [F::ZERO; WIDTH];
+        for i in 0..4 {
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            let input_lhs = Self::wire_input(i);
+            let input_rhs = Self::wire_input(i + 4);
+            state[i] = vars.local_wires[input_lhs] + delta_i;
+            state[i + 4] = vars.local_wires[input_rhs] - delta_i;
+        }
+        for i in 8..WIDTH {
+            state[i] = vars.local_wires[Self::wire_input(i)];
+        }
+
+        // The initial linear layer.
+        <F as Poseidon2>::external_linear_layer(&mut state);
+
+        // The first half of the external rounds.
+        for r in 0..ROUNDS_F_HALF {
+            <F as Poseidon2>::add_rc(&mut state, r);
+            if r != 0 {
+                for i in 0..WIDTH {
+                    let sbox_in = vars.local_wires[Self::wire_full_sbox_0(r, i)];
+                    yield_constr.one(state[i] - sbox_in);
+                    state[i] = sbox_in;
+                }
+            }
+            <F as Poseidon2>::sbox(&mut state);
+            <F as Poseidon2>::external_linear_layer(&mut state);
+        }
+
+        // The internal rounds.
+        for r in 0..ROUNDS_P {
+            state[0] += F::from_canonical_u64(INTERNAL_CONSTANTS[r]);
+            let sbox_in = vars.local_wires[Self::wire_partial_sbox(r)];
+            yield_constr.one(state[0] - sbox_in);
+            state[0] = sbox_in;
+            state[0] = <F as Poseidon2>::sbox_p(&state[0]);
+            <F as Poseidon2>::internal_linear_layer(&mut state);
+        }
+
+        // The second half of the external rounds.
+        for r in ROUNDS_F_HALF..ROUNDS_F {
+            <F as Poseidon2>::add_rc(&mut state, r);
+            for i in 0..WIDTH {
+                let sbox_in = vars.local_wires[Self::wire_full_sbox_1(r - ROUNDS_F_HALF, i)];
+                yield_constr.one(state[i] - sbox_in);
+                state[i] = sbox_in;
+            }
+            <F as Poseidon2>::sbox(&mut state);
+            <F as Poseidon2>::external_linear_layer(&mut state);
+        }
+
+        for i in 0..WIDTH {
+            yield_constr.one(state[i] - vars.local_wires[Self::wire_output(i)]);
+        }
+    }
+
+    fn eval_unfiltered_circuit(
+        &self,
+        builder: &mut CircuitBuilder<F, D>,
+        vars: EvaluationTargets<D>,
+    ) -> Vec<ExtensionTarget<D>> {
+        let mut constraints = Vec::with_capacity(self.num_constraints());
+
+        // Assert that `swap` is binary.
+        let swap = vars.local_wires[Self::WIRE_SWAP];
+        constraints.push(builder.mul_sub_extension(swap, swap, swap));
+
+        // Assert that each delta wire is set properly: `delta_i = swap * (rhs - lhs)`.
+        for i in 0..4 {
+            let input_lhs = vars.local_wires[Self::wire_input(i)];
+            let input_rhs = vars.local_wires[Self::wire_input(i + 4)];
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            let diff = builder.sub_extension(input_rhs, input_lhs);
+            constraints.push(builder.mul_sub_extension(swap, diff, delta_i));
+        }
+
+        // Compute the possibly-swapped input layer.
+        let mut state = [builder.zero_extension(); WIDTH];
+        for i in 0..4 {
+            let delta_i = vars.local_wires[Self::wire_delta(i)];
+            let input_lhs = vars.local_wires[Self::wire_input(i)];
+            let input_rhs = vars.local_wires[Self::wire_input(i + 4)];
+            state[i] = builder.add_extension(input_lhs, delta_i);
+            state[i + 4] = builder.sub_extension(input_rhs, delta_i);
+        }
+        for i in 8..WIDTH {
+            state[i] = vars.local_wires[Self::wire_input(i)];
+        }
+
+        // The initial linear layer.
+        <F as Poseidon2>::external_linear_layer_circuit(builder, &mut state);
+
+        // The first half of the external rounds.
+        for r in 0..ROUNDS_F_HALF {
+            <F as Poseidon2>::add_rc_circuit(builder, &mut state, r);
+            if r != 0 {
+                for i in 0..WIDTH {
+                    let sbox_in = vars.local_wires[Self::wire_full_sbox_0(r, i)];
+                    constraints.push(builder.sub_extension(state[i], sbox_in));
+                    state[i] = sbox_in;
+                }
+            }
+            <F as Poseidon2>::sbox_circuit(builder, &mut state);
+            <F as Poseidon2>::external_linear_layer_circuit(builder, &mut state);
+        }
+
+        // The internal rounds.
+        for r in 0..ROUNDS_P {
+            let round_constant = F::Extension::from_canonical_u64(INTERNAL_CONSTANTS[r]);
+            let round_constant = builder.constant_extension(round_constant);
+            state[0] = builder.add_extension(state[0], round_constant);
+
+            let sbox_in = vars.local_wires[Self::wire_partial_sbox(r)];
+            constraints.push(builder.sub_extension(state[0], sbox_in));
+            state[0] = sbox_in;
+            state[0] = <F as Poseidon2>::sbox_p_circuit(builder, state[0]);
+            <F as Poseidon2>::internal_linear_layer_circuit(builder, &mut state);
+        }
+
+        // The second half of the external rounds.
+        for r in ROUNDS_F_HALF..ROUNDS_F {
+            <F as Poseidon2>::add_rc_circuit(builder, &mut state, r);
+
+            for i in 0..WIDTH {
+                let sbox_in = vars.local_wires[Self::wire_full_sbox_1(r - ROUNDS_F_HALF, i)];
+                constraints.push(builder.sub_extension(state[i], sbox_in));
+                state[i] = sbox_in;
+            }
+            <F as Poseidon2>::sbox_circuit(builder, &mut state);
+            <F as Poseidon2>::external_linear_layer_circuit(builder, &mut state);
+        }
+
+        for i in 0..WIDTH {
+            constraints
+                .push(builder.sub_extension(state[i], vars.local_wires[Self::wire_output(i)]));
+        }
+
+        constraints
+    }
+
+    fn generators(&self, row: usize, _local_constants: &[F]) -> Vec<WitnessGeneratorRef<F, D>> {
+        let g = Poseidon2Generator::<F, D> {
+            row,
+            _phantom: PhantomData,
+        };
+        vec![WitnessGeneratorRef::new(g.adapter())]
+    }
+
+    fn num_wires(&self) -> usize {
+        Self::end()
+    }
+
+    fn num_constants(&self) -> usize {
+        0
+    }
+
+    fn degree(&self) -> usize {
+        7
+    }
+
+    fn num_constraints(&self) -> usize {
+        WIDTH * (ROUNDS_F - 1) + ROUNDS_P + WIDTH + 1 + 4
+    }
+}
+
+#[derive(Debug, Default)]
+pub struct Poseidon2Generator<F: RichField + Extendable<D> + Poseidon2, const D: usize> {
+    row: usize,
+    _phantom: PhantomData<F>,
+}
+
+impl<F: RichField + Extendable<D> + Poseidon2, const D: usize> SimpleGenerator<F, D>
+    for Poseidon2Generator<F, D>
+{
+    fn id(&self) -> String {
+        "Poseidon2Generator".to_string()
+    }
+
+    fn dependencies(&self) -> Vec<Target> {
+        (0..WIDTH)
+            .map(|i| Poseidon2Gate::<F, D>::wire_input(i))
+            .chain(Some(Poseidon2Gate::<F, D>::WIRE_SWAP))
+            .map(|column| Target::wire(self.row, column))
+            .collect()
+    }
+
+    fn run_once(
+        &self,
+        witness: &PartitionWitness<F>,
+        out_buffer: &mut GeneratedValues<F>,
+    ) -> Result<()> {
+        let local_wire = |column| Wire {
+            row: self.row,
+            column,
+        };
+
+        let mut state = (0..WIDTH)
+            .map(|i| witness.get_wire(local_wire(Poseidon2Gate::<F, D>::wire_input(i))))
+            .collect::<Vec<_>>();
+
+        let swap_value = witness.get_wire(local_wire(Poseidon2Gate::<F, D>::WIRE_SWAP));
+        debug_assert!(swap_value == F::ZERO || swap_value == F::ONE);
+
+        for i in 0..4 {
+            let delta_i = swap_value * (state[i + 4] - state[i]);
+            out_buffer.set_wire(local_wire(Poseidon2Gate::<F, D>::wire_delta(i)), delta_i)?;
+        }
+
+        if swap_value == F::ONE {
+            for i in 0..4 {
+                state.swap(i, 4 + i);
+            }
+        }
+
+        let mut state: [F; WIDTH] = state.try_into().unwrap();
+
+        <F as Poseidon2>::external_linear_layer(&mut state);
+
+        // The first half of the external rounds.
+        for r in 0..ROUNDS_F_HALF {
+            <F as Poseidon2>::add_rc(&mut state, r);
+            if r != 0 {
+                for i in 0..WIDTH {
+                    out_buffer.set_wire(
+                        local_wire(Poseidon2Gate::<F, D>::wire_full_sbox_0(r, i)),
+                        state[i],
+                    )?;
+                }
+            }
+            <F as Poseidon2>::sbox(&mut state);
+            <F as Poseidon2>::external_linear_layer(&mut state);
+        }
+
+        // The internal rounds.
+        for r in 0..ROUNDS_P {
+            state[0] += F::from_canonical_u64(INTERNAL_CONSTANTS[r]);
+            out_buffer.set_wire(
+                local_wire(Poseidon2Gate::<F, D>::wire_partial_sbox(r)),
+                state[0],
+            )?;
+            state[0] = <F as Poseidon2>::sbox_p(&state[0]);
+            <F as Poseidon2>::internal_linear_layer(&mut state);
+        }
+
+        // The second half of the external rounds.
+        for r in ROUNDS_F_HALF..ROUNDS_F {
+            <F as Poseidon2>::add_rc(&mut state, r);
+            for i in 0..WIDTH {
+                out_buffer.set_wire(
+                    local_wire(Poseidon2Gate::<F, D>::wire_full_sbox_1(
+                        r - ROUNDS_F_HALF,
+                        i,
+                    )),
+                    state[i],
+                )?;
+            }
+            <F as Poseidon2>::sbox(&mut state);
+            <F as Poseidon2>::external_linear_layer(&mut state);
+        }
+
+        for i in 0..WIDTH {
+            out_buffer.set_wire(local_wire(Poseidon2Gate::<F, D>::wire_output(i)), state[i])?;
+        }
+
+        Ok(())
+    }
+
+    fn serialize(&self, dst: &mut Vec<u8>, _common_data: &CommonCircuitData<F, D>) -> IoResult<()> {
+        dst.write_usize(self.row)
+    }
+
+    fn deserialize(src: &mut Buffer, _common_data: &CommonCircuitData<F, D>) -> IoResult<Self> {
+        let row = src.read_usize()?;
+        Ok(Self {
+            row,
+            _phantom: PhantomData,
+        })
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use anyhow::Result;
+
+    use super::{Poseidon2Gate, *};
+    use crate::field::goldilocks_field::GoldilocksField;
+    use crate::gates::gate_testing::{test_eval_fns, test_low_degree};
+    use crate::gates::poseidon::PoseidonGate;
+    use crate::iop::generator::generate_partial_witness;
+    use crate::iop::witness::PartialWitness;
+    use crate::plonk::circuit_data::CircuitConfig;
+    use crate::plonk::config::{GenericConfig, Poseidon2GoldilocksConfig};
+
+    #[test]
+    fn wire_indices() {
+        type F = GoldilocksField;
+        type Gate = Poseidon2Gate<F, 4>;
+
+        assert_eq!(Gate::wire_input(0), 0);
+        assert_eq!(Gate::wire_input(11), 11);
+        assert_eq!(Gate::wire_output(0), 12);
+        assert_eq!(Gate::wire_output(11), 23);
+        assert_eq!(Gate::WIRE_SWAP, 24);
+        assert_eq!(Gate::wire_delta(0), 25);
+        assert_eq!(Gate::wire_delta(3), 28);
+        assert_eq!(Gate::wire_full_sbox_0(1, 0), 29);
+        assert_eq!(Gate::wire_full_sbox_0(3, 0), 53);
+        assert_eq!(Gate::wire_full_sbox_0(3, 11), 64);
+        assert_eq!(Gate::wire_partial_sbox(0), 65);
+        assert_eq!(Gate::wire_partial_sbox(21), 86);
+        assert_eq!(Gate::wire_full_sbox_1(0, 0), 87);
+        assert_eq!(Gate::wire_full_sbox_1(3, 0), 123);
+        assert_eq!(Gate::wire_full_sbox_1(3, 11), 134);
+    }
+
+    #[test]
+    fn generated_output() {
+        const D: usize = 2;
+        type C = Poseidon2GoldilocksConfig;
+        type F = <C as GenericConfig<D>>::F;
+
+        let config = CircuitConfig {
+            num_wires: 143,
+            ..CircuitConfig::standard_recursion_config()
+        };
+        let mut builder = CircuitBuilder::new(config);
+        type Gate = Poseidon2Gate<F, D>;
+        let gate = Gate::new();
+        let row = builder.add_gate(gate, vec![]);
+        let circuit = builder.build_prover::<C>();
+
+        let permutation_inputs = (0..WIDTH).map(F::from_canonical_usize).collect::<Vec<_>>();
+
+        let mut inputs = PartialWitness::new();
+        inputs
+            .set_wire(
+                Wire {
+                    row,
+                    column: Gate::WIRE_SWAP,
+                },
+                F::ZERO,
+            )
+            .unwrap();
+        for i in 0..WIDTH {
+            inputs
+                .set_wire(
+                    Wire {
+                        row,
+                        column: Gate::wire_input(i),
+                    },
+                    permutation_inputs[i],
+                )
+                .unwrap();
+        }
+
+        let witness =
+            generate_partial_witness(inputs, &circuit.prover_only, &circuit.common).unwrap();
+
+        let expected_outputs: [F; WIDTH] = F::poseidon2(permutation_inputs.try_into().unwrap());
+        for i in 0..WIDTH {
+            let out = witness.get_wire(Wire {
+                row: 0,
+                column: Gate::wire_output(i),
+            });
+            assert_eq!(out, expected_outputs[i]);
+        }
+    }
+
+    #[test]
+    fn low_degree() {
+        type F = GoldilocksField;
+        let gate = Poseidon2Gate::<F, 4>::new();
+        test_low_degree(gate);
+
+        let gate = PoseidonGate::<F, 4>::new();
+        test_low_degree(gate)
+    }
+
+    #[test]
+    fn eval_fns() -> Result<()> {
+        const D: usize = 2;
+        type C = Poseidon2GoldilocksConfig;
+        type F = <C as GenericConfig<D>>::F;
+        let gate = Poseidon2Gate::<F, D>::new();
+        test_eval_fns::<F, C, _, D>(gate)?;
+
+        let gate = PoseidonGate::<F, D>::new();
+        test_eval_fns::<F, C, _, D>(gate)
+    }
+}
diff --git a/plonky2/src/hash/poseidon2/hash.rs b/plonky2/src/hash/poseidon2/hash.rs
new file mode 100644
index 000000000..5f9553fad
--- /dev/null
+++ b/plonky2/src/hash/poseidon2/hash.rs
@@ -0,0 +1,676 @@
+use core::fmt::Debug;
+
+use plonky2_field::ops::Square;
+
+use super::config::*;
+use super::gate::Poseidon2Gate;
+use crate::field::extension::{Extendable, FieldExtension};
+use crate::field::goldilocks_field::GoldilocksField as F;
+use crate::field::types::{Field, PrimeField64};
+use crate::hash::hash_types::{HashOut, RichField};
+use crate::hash::hashing::{compress, hash_n_to_hash_no_pad, PlonkyPermutation};
+use crate::iop::ext_target::ExtensionTarget;
+use crate::iop::target::{BoolTarget, Target};
+use crate::plonk::circuit_builder::CircuitBuilder;
+use crate::plonk::config::{AlgebraicHasher, Hasher};
+
+pub trait Poseidon2: PrimeField64 {
+    #[inline]
+    fn poseidon2(input: [Self; WIDTH]) -> [Self; WIDTH] {
+        let mut state = input;
+
+        Self::external_linear_layer(&mut state);
+
+        Self::full_rounds(&mut state, 0);
+        Self::partial_rounds(&mut state);
+        Self::full_rounds(&mut state, ROUNDS_F_HALF);
+
+        state
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn full_rounds(state: &mut [Self; WIDTH], start: usize) {
+        for r in start..(start + ROUNDS_F_HALF) {
+            Self::add_rc(state, r);
+            Self::sbox(state);
+            Self::external_linear_layer(state);
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn partial_rounds(state: &mut [Self; WIDTH]) {
+        for r in 0..ROUNDS_P {
+            state[0] += Self::from_canonical_u64(INTERNAL_CONSTANTS[r]);
+            state[0] = Self::sbox_p(&state[0]);
+            Self::internal_linear_layer(state);
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn external_linear_layer(state: &mut [Self; WIDTH]) {
+        let mut state_u128: [u128; WIDTH] = [0u128; WIDTH];
+        for i in 0..WIDTH {
+            state_u128[i] = state[i].to_noncanonical_u64() as u128;
+        }
+        external_linear_layer_u128(&mut state_u128);
+        for i in 0..WIDTH {
+            state[i] = Self::from_noncanonical_u128_with_96_bits(state_u128[i]);
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn external_linear_layer_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(
+        state: &mut [F; WIDTH],
+    ) {
+        // First, we apply M_4 to each consecutive four elements of the state.
+        // In Appendix B's terminology, this replaces each x_i with x_i'.
+        for i in (0..WIDTH).step_by(4) {
+            // Would be nice to find a better way to do this.
+            let mut state_4 = [state[i], state[i + 1], state[i + 2], state[i + 3]];
+            Self::apply_mat4_mut_extension(&mut state_4);
+            state[i..i + 4].clone_from_slice(&state_4);
+        }
+        // Now, we apply the outer circulant matrix (to compute the y_i values).
+
+        // We first precompute the four sums of every four elements.
+        let sums: [F; 4] =
+            core::array::from_fn(|k| (0..WIDTH).step_by(4).map(|j| state[j + k]).sum::<F>());
+
+        // The formula for each y_i involves 2x_i' term and x_j' terms for each j that equals i mod 4.
+        // In other words, we can add a single copy of x_i' to the appropriate one of our precomputed sums
+        for i in 0..WIDTH {
+            state[i] += sums[i % 4];
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn internal_linear_layer(state: &mut [Self; WIDTH]) {
+        let sum = sum_12(state); // hard coded for WIDTH = 12
+        for i in 0..WIDTH {
+            state[i] =
+                sum.multiply_accumulate(state[i], Self::from_canonical_u64(MATRIX_DIAG_12_U64[i]));
+        }
+    }
+
+    #[inline]
+    fn internal_linear_layer_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(
+        state: &mut [F; WIDTH],
+    ) {
+        let sum: F = state.iter().cloned().sum();
+        state
+            .iter_mut()
+            .zip(MATRIX_DIAG_12_U64.iter())
+            .for_each(|(x, &m)| {
+                *x = sum.multiply_accumulate(*x, F::from_canonical_u64(m));
+            });
+    }
+
+    fn add_rc(state: &mut [Self; WIDTH], external_round: usize);
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn add_rc_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(
+        state: &mut [F; WIDTH],
+        external_round: usize,
+    ) {
+        debug_assert!(external_round < EXTERNAL_CONSTANTS.len());
+
+        for i in 0..WIDTH {
+            state[i] += F::from_canonical_u64(EXTERNAL_CONSTANTS[external_round][i]);
+        }
+    }
+
+    fn sbox(state: &mut [Self; WIDTH]);
+
+    #[inline]
+    fn sbox_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(
+        state: &mut [F; WIDTH],
+    ) {
+        state
+            .iter_mut()
+            .for_each(|a| *a = Self::sbox_p_extension(a));
+    }
+
+    fn sbox_p(a: &Self) -> Self;
+
+    fn sbox_p_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(a: &F) -> F;
+
+    #[inline]
+    fn apply_mat4_mut_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(
+        x: &mut [F; 4],
+    ) {
+        let t01 = x[0] + x[1];
+        let t23 = x[2] + x[3];
+        let t0123 = t01 + t23;
+        let t01123 = t0123 + x[1];
+        let t01233 = t0123 + x[3];
+        // The order here is important. Need to overwrite x[0] and x[2] after x[1] and x[3].
+        x[3] = t01233 + x[0].double(); // 3*x[0] + x[1] + x[2] + 2*x[3]
+        x[1] = t01123 + x[2].double(); // x[0] + 2*x[1] + 3*x[2] + x[3]
+        x[0] = t01123 + t01; // 2*x[0] + 3*x[1] + x[2] + x[3]
+        x[2] = t01233 + t23; // x[0] + x[1] + 2*x[2] + 3*x[3]
+    }
+
+    // In circuit functions
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn external_linear_layer_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        state: &mut [ExtensionTarget<D>; WIDTH],
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        // First, we apply M_4 to each consecutive four elements of the state.
+        // In Appendix B's terminology, this replaces each x_i with x_i'.
+        for i in (0..WIDTH).step_by(4) {
+            Self::apply_mat4_mut_circuit(builder, (&mut state[i..i + 4]).try_into().unwrap());
+        }
+        // Now, we apply the outer circulant matrix (to compute the y_i values).
+
+        // We first precompute the four sums of every four elements.
+        let sums: [ExtensionTarget<D>; 4] = core::array::from_fn(|k| {
+            (0..WIDTH)
+                .step_by(4)
+                .map(|j| state[j + k])
+                .reduce(|acc, t| builder.add_extension(acc, t))
+                .unwrap()
+        });
+
+        // The formula for each y_i involves 2x_i' term and x_j' terms for each j that equals i mod 4.
+        // In other words, we can add a single copy of x_i' to the appropriate one of our precomputed sums
+        for i in 0..WIDTH {
+            state[i] = builder.add_extension(state[i], sums[i % 4]);
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn apply_mat4_mut_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        x: &mut [ExtensionTarget<D>; 4],
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        let two = builder.constant_extension(Self::Extension::from_canonical_u64(2));
+
+        let t01 = builder.add_extension(x[0], x[1]);
+        let t23 = builder.add_extension(x[2], x[3]);
+        let t0123 = builder.add_extension(t01, t23);
+        let t01123 = builder.add_extension(t0123, x[1]);
+        let t01233 = builder.add_extension(t0123, x[3]);
+        // The order here is important. Need to overwrite x[0] and x[2] after x[1] and x[3].
+        let dx0 = builder.mul_extension(x[0], two);
+        let dx2 = builder.mul_extension(x[2], two);
+        x[3] = builder.add_extension(t01233, dx0); // 3*x[0] + x[1] + x[2] + 2*x[3]
+        x[1] = builder.add_extension(t01123, dx2); // x[0] + 2*x[1] + 3*x[2] + x[3]
+        x[0] = builder.add_extension(t01123, t01); // 2*x[0] + 3*x[1] + x[2] + x[3]
+        x[2] = builder.add_extension(t01233, t23); // x[0] + x[1] + 2*x[2] + 3*x[3]
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn matmul_m4_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        input: &mut [ExtensionTarget<D>; WIDTH],
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        for i in 0..3 {
+            let t_0 = builder.mul_const_add_extension(Self::ONE, input[i * 4], input[i * 4 + 1]);
+            let t_1 =
+                builder.mul_const_add_extension(Self::ONE, input[i * 4 + 2], input[i * 4 + 3]);
+            let t_2 = builder.mul_const_add_extension(Self::TWO, input[i * 4 + 1], t_1);
+            let t_3 = builder.mul_const_add_extension(Self::TWO, input[i * 4 + 3], t_0);
+
+            let four = Self::TWO + Self::TWO;
+
+            let t_4 = builder.mul_const_add_extension(four, t_1, t_3);
+            let t_5 = builder.mul_const_add_extension(four, t_0, t_2);
+            let t_6 = builder.mul_const_add_extension(Self::ONE, t_3, t_5);
+            let t_7 = builder.mul_const_add_extension(Self::ONE, t_2, t_4);
+
+            input[i * 4] = t_6;
+            input[i * 4 + 1] = t_5;
+            input[i * 4 + 2] = t_7;
+            input[i * 4 + 3] = t_4;
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn add_rc_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        input: &mut [ExtensionTarget<D>; WIDTH],
+        rc_index: usize,
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        for i in 0..WIDTH {
+            let round_constant =
+                Self::Extension::from_canonical_u64(EXTERNAL_CONSTANTS[rc_index][i]);
+            let round_constant = builder.constant_extension(round_constant);
+            input[i] = builder.add_extension(input[i], round_constant);
+        }
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn sbox_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        input: &mut [ExtensionTarget<D>; WIDTH],
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        for i in 0..WIDTH {
+            input[i] = Self::sbox_p_circuit(builder, input[i]);
+        }
+    }
+
+    #[inline]
+    fn sbox_p_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        input: ExtensionTarget<D>,
+    ) -> ExtensionTarget<D>
+    where
+        Self: RichField + Extendable<D>,
+    {
+        builder.exp_u64_extension(input, super::config::D)
+    }
+
+    #[inline]
+    #[unroll::unroll_for_loops]
+    fn internal_linear_layer_circuit<const D: usize>(
+        builder: &mut CircuitBuilder<Self, D>,
+        input: &mut [ExtensionTarget<D>; WIDTH],
+    ) where
+        Self: RichField + Extendable<D>,
+    {
+        let sum = builder.add_many_extension([
+            input[0], input[1], input[2], input[3], input[4], input[5], input[6], input[7],
+            input[8], input[9], input[10], input[11],
+        ]);
+
+        for i in 0..WIDTH {
+            let round_constant = Self::Extension::from_canonical_u64(MATRIX_DIAG_12_U64[i]);
+            let round_constant = builder.constant_extension(round_constant);
+
+            input[i] = builder.mul_add_extension(round_constant, input[i], sum);
+        }
+    }
+}
+
+#[inline]
+#[unroll::unroll_for_loops]
+fn external_linear_layer_u128(state: &mut [u128; WIDTH]) {
+    // First, we apply M_4 to each consecutive four elements of the state.
+    // In Appendix B's terminology, this replaces each x_i with x_i'.
+    for i in (0..WIDTH).step_by(4) {
+        // Multiply a 4-element vector x by:
+        // [ 2 3 1 1 ]
+        // [ 1 2 3 1 ]
+        // [ 1 1 2 3 ]
+        // [ 3 1 1 2 ].
+        let t01 = state[i] + state[i + 1];
+        let t23 = state[i + 2] + state[i + 3];
+        let t0123 = t01 + t23;
+
+        let x0 = state[i];
+        let x2 = state[i + 2];
+
+        state[i] = t0123 + t01 + state[i + 1]; // 2*x[0] + 3*x[1] + x[2] + x[3]
+        state[i + 1] = t0123 + state[i + 1] + x2 + x2; // x[0] + 2*x[1] + 3*x[2] + x[3]
+        state[i + 2] = t0123 + t23 + state[i + 3]; // x[0] + x[1] + 2*x[2] + 3*x[3]
+        state[i + 3] = t0123 + state[i + 3] + x0 + x0; // 3*x[0] + x[1] + x[2] + 2*x[3]
+    }
+    // Now, we apply the outer circulant matrix (to compute the y_i values).
+
+    // We first precompute the four sums of every four elements.
+    let mut sums = [0u128; 4];
+    for i in 0..4 {
+        sums[i] = state[i] + state[i + 4] + state[i + 8];
+    }
+
+    // The formula for each y_i involves 2x_i' term and x_j' terms for each j that equals i mod 4.
+    // In other words, we can add a single copy of x_i' to the appropriate one of our precomputed sums
+    for i in 0..WIDTH {
+        state[i] += sums[i % 4];
+    }
+}
+
+impl Poseidon2 for F {
+    #[inline]
+    fn sbox_p(a: &Self) -> Self {
+        let a2 = a.square();
+        let a4 = a2.square();
+        let a3 = *a * a2;
+        a3 * a4
+    }
+
+    #[inline]
+    fn sbox_p_extension<F: FieldExtension<D, BaseField = Self>, const D: usize>(a: &F) -> F {
+        let a2 = a.square();
+        let a4 = a2.square();
+        let a3 = *a * a2;
+        a3 * a4
+    }
+
+    #[inline]
+    #[cfg(not(all(target_arch = "aarch64", target_feature = "neon")))]
+    fn add_rc(state: &mut [Self; WIDTH], external_round: usize) {
+        use plonky2_field::types::Field64;
+        debug_assert!(external_round < EXTERNAL_CONSTANTS.len());
+        state
+            .iter_mut()
+            .zip(EXTERNAL_CONSTANTS[external_round].iter())
+            .for_each(|(x, &m)| {
+                *x = unsafe { x.add_canonical_u64(m) };
+            });
+    }
+
+    #[inline]
+    #[cfg(all(target_arch = "aarch64", target_feature = "neon"))]
+    fn add_rc(state: &mut [Self; WIDTH], external_round: usize) {
+        debug_assert!(external_round < EXTERNAL_CONSTANTS.len());
+
+        unsafe {
+            use core::mem::transmute;
+
+            use crate::hash::arch::aarch64::poseidon_goldilocks_neon::vector_add;
+
+            let state_u64 = transmute::<[Self; WIDTH], [u64; WIDTH]>(*state);
+            let round_constants = &EXTERNAL_CONSTANTS[external_round];
+
+            let res = vector_add(&state_u64, round_constants);
+            *state = transmute::<[u64; WIDTH], [Self; WIDTH]>(res);
+        }
+    }
+
+    #[inline]
+    #[cfg(not(all(target_arch = "aarch64", target_feature = "neon")))]
+    fn sbox(state: &mut [Self; WIDTH]) {
+        state.iter_mut().for_each(|a| *a = Self::sbox_p(a));
+    }
+
+    #[inline(always)]
+    #[cfg(all(target_arch = "aarch64", target_feature = "neon"))]
+    fn sbox(state: &mut [Self; WIDTH]) {
+        unsafe {
+            crate::hash::arch::aarch64::poseidon_goldilocks_neon::sbox_layer(state);
+        }
+    }
+}
+
+#[derive(Copy, Clone, Default, Debug, PartialEq)]
+pub struct Poseidon2Permutation<T> {
+    state: [T; WIDTH],
+}
+
+impl<T: Eq> Eq for Poseidon2Permutation<T> {}
+
+impl<T> AsRef<[T]> for Poseidon2Permutation<T> {
+    fn as_ref(&self) -> &[T] {
+        &self.state
+    }
+}
+
+trait Permuter: Sized {
+    fn permute(input: [Self; WIDTH]) -> [Self; WIDTH];
+}
+
+impl<F: Poseidon2> Permuter for F {
+    fn permute(input: [Self; WIDTH]) -> [Self; WIDTH] {
+        <F as Poseidon2>::poseidon2(input)
+    }
+}
+
+impl Permuter for Target {
+    fn permute(_input: [Self; WIDTH]) -> [Self; WIDTH] {
+        panic!("Call `permute_swapped()` instead of `permute()`");
+    }
+}
+
+impl<T: Copy + Debug + Default + Eq + Permuter + Send + Sync> PlonkyPermutation<T>
+    for Poseidon2Permutation<T>
+{
+    const RATE: usize = RATE;
+    const WIDTH: usize = WIDTH;
+
+    fn new<I: IntoIterator<Item = T>>(elts: I) -> Self {
+        let mut perm = Self {
+            state: [T::default(); WIDTH],
+        };
+        perm.set_from_iter(elts, 0);
+        perm
+    }
+
+    fn set_elt(&mut self, elt: T, idx: usize) {
+        self.state[idx] = elt;
+    }
+
+    fn set_from_slice(&mut self, elts: &[T], start_idx: usize) {
+        let begin = start_idx;
+        let end = start_idx + elts.len();
+        self.state[begin..end].copy_from_slice(elts);
+    }
+
+    fn set_from_iter<I: IntoIterator<Item = T>>(&mut self, elts: I, start_idx: usize) {
+        for (s, e) in self.state[start_idx..].iter_mut().zip(elts) {
+            *s = e;
+        }
+    }
+
+    fn permute(&mut self) {
+        self.state = T::permute(self.state);
+    }
+
+    fn squeeze(&self) -> &[T] {
+        &self.state[..Self::RATE]
+    }
+}
+
+#[inline]
+/// Sum of 12 elements to u128; unrolled for performance.
+fn sum_12<F: PrimeField64>(inputs: &[F]) -> F {
+    debug_assert!(inputs.len() == 12);
+    let tmp = inputs[0].to_noncanonical_u64() as u128
+        + inputs[1].to_noncanonical_u64() as u128
+        + inputs[2].to_noncanonical_u64() as u128
+        + inputs[3].to_noncanonical_u64() as u128
+        + inputs[4].to_noncanonical_u64() as u128
+        + inputs[5].to_noncanonical_u64() as u128
+        + inputs[6].to_noncanonical_u64() as u128
+        + inputs[7].to_noncanonical_u64() as u128
+        + inputs[8].to_noncanonical_u64() as u128
+        + inputs[9].to_noncanonical_u64() as u128
+        + inputs[10].to_noncanonical_u64() as u128
+        + inputs[11].to_noncanonical_u64() as u128;
+
+    F::from_noncanonical_u128_with_96_bits(tmp)
+}
+
+/// Poseidon2 hash function.
+#[derive(Copy, Clone, Debug, Eq, PartialEq)]
+pub struct Poseidon2Hash;
+impl<F: RichField + Poseidon2> Hasher<F> for Poseidon2Hash {
+    const HASH_SIZE: usize = 4 * 8;
+    type Hash = HashOut<F>;
+    type Permutation = Poseidon2Permutation<F>;
+
+    fn hash_no_pad(input: &[F]) -> Self::Hash {
+        hash_n_to_hash_no_pad::<F, Self::Permutation>(input)
+    }
+
+    fn two_to_one(left: Self::Hash, right: Self::Hash) -> Self::Hash {
+        compress::<F, Self::Permutation>(left, right)
+    }
+}
+
+impl Poseidon2Hash {
+    #[inline]
+    #[unroll::unroll_for_loops]
+    pub fn hash_n_to_one(
+        input: &[<Poseidon2Hash as Hasher<F>>::Hash],
+    ) -> <Poseidon2Hash as Hasher<F>>::Hash {
+        assert!(!input.is_empty());
+
+        if input.len() == 1 {
+            return input[0];
+        }
+
+        let mut result = <Poseidon2Hash as Hasher<F>>::two_to_one(input[0], input[1]);
+
+        for i in 2..input.len() {
+            result = <Poseidon2Hash as Hasher<F>>::two_to_one(result, input[i]);
+        }
+
+        result
+    }
+}
+
+impl<F: RichField + Poseidon2> AlgebraicHasher<F> for Poseidon2Hash {
+    type AlgebraicPermutation = Poseidon2Permutation<Target>;
+
+    fn permute_swapped<const D: usize>(
+        inputs: Self::AlgebraicPermutation,
+        swap: BoolTarget,
+        builder: &mut CircuitBuilder<F, D>,
+    ) -> Self::AlgebraicPermutation
+    where
+        F: RichField + Extendable<D>,
+    {
+        let gate_type = Poseidon2Gate::<F, D>::new();
+        let gate = builder.add_gate(gate_type, vec![]);
+
+        let swap_wire = Poseidon2Gate::<F, D>::WIRE_SWAP;
+        let swap_wire = Target::wire(gate, swap_wire);
+        builder.connect(swap.target, swap_wire);
+
+        // Route input wires.
+        let inputs = inputs.as_ref();
+        for i in 0..WIDTH {
+            let in_wire = Poseidon2Gate::<F, D>::wire_input(i);
+            let in_wire = Target::wire(gate, in_wire);
+            builder.connect(inputs[i], in_wire);
+        }
+
+        // Collect output wires.
+        Self::AlgebraicPermutation::new(
+            (0..WIDTH).map(|i| Target::wire(gate, Poseidon2Gate::<F, D>::wire_output(i))),
+        )
+    }
+}
+
+#[cfg(test)]
+mod test {
+    use anyhow::Result;
+    use num::{BigUint, One};
+    use p3_field::{AbstractField, PrimeField64 as _};
+    use p3_goldilocks::Goldilocks;
+    use rand::{thread_rng, RngCore};
+
+    use super::*;
+    use crate::field::types::PrimeField64;
+    use crate::hash::hashing::hash_n_to_m_no_pad;
+    use crate::hash::poseidon2::p3::p3_poseidon2_hash_n_to_m_no_pad;
+    use crate::iop::witness::{PartialWitness, WitnessWrite};
+    use crate::plonk::circuit_data::CircuitConfig;
+    use crate::plonk::config::PoseidonGoldilocksConfig;
+
+    #[test]
+    fn test_poseidon2_with_plonky3() {
+        let mut rng = thread_rng();
+
+        let input: [u32; 12] = core::array::from_fn(|_| rng.next_u32());
+
+        let input_f = input
+            .iter()
+            .map(|&x| F::from_canonical_u64((x as u64) + 1073741824))
+            .collect::<Vec<F>>();
+        let expected_output_f = hash_n_to_m_no_pad::<F, Poseidon2Permutation<F>>(&input_f, 12);
+
+        let input_f3 = input
+            .iter()
+            .map(|&x| Goldilocks::from_canonical_u64((x as u64) + 1073741824))
+            .collect::<Vec<Goldilocks>>();
+        let expected_output_f3 = p3_poseidon2_hash_n_to_m_no_pad(&input_f3, 12);
+
+        for i in 0..4 {
+            assert_eq!(
+                expected_output_f[i].to_canonical_u64(),
+                expected_output_f3[i].as_canonical_u64()
+            );
+        }
+    }
+
+    #[test]
+    fn test_poseidon2_gate() -> Result<()> {
+        let mut rng = thread_rng();
+
+        let input: [u32; 12] = core::array::from_fn(|_| rng.next_u32());
+        let input_f = input
+            .iter()
+            .map(|&x| F::from_canonical_u64((x as u64) + 1073741824))
+            .collect::<Vec<F>>();
+
+        let expected_output = hash_n_to_m_no_pad::<F, Poseidon2Permutation<F>>(&input_f[0..8], 4);
+
+        let mut builder = CircuitBuilder::<F, 2>::new(CircuitConfig::standard_recursion_config());
+
+        let input_target: [Target; 12] = input_f
+            .iter()
+            .map(|&x| builder.constant(x))
+            .collect::<Vec<Target>>()
+            .try_into()
+            .unwrap();
+        let output_target =
+            builder.hash_n_to_m_no_pad::<Poseidon2Hash>(input_target[0..8].to_vec(), 4);
+
+        let expected_output_target = builder.add_virtual_target_arr::<4>();
+        for i in 0..4 {
+            builder.connect(expected_output_target[i], output_target[i]);
+        }
+
+        let circuit = builder.build::<PoseidonGoldilocksConfig>();
+        let mut pw = PartialWitness::new();
+        pw.set_target_arr(&expected_output_target, &expected_output)?;
+
+        let proof = circuit.prove(pw).unwrap();
+        circuit.verify(proof.clone())
+    }
+
+    #[test]
+    fn test_poseidon2_gate_big() -> Result<()> {
+        let input_f: [F; 12] =
+            core::array::from_fn(|_| F::from_noncanonical_biguint(F::order() - BigUint::one()));
+
+        let expected_output = hash_n_to_m_no_pad::<F, Poseidon2Permutation<F>>(&input_f[0..8], 4);
+
+        let mut builder = CircuitBuilder::<F, 2>::new(CircuitConfig::standard_recursion_config());
+
+        let input_target: [Target; 12] = input_f
+            .iter()
+            .map(|&x| builder.constant(x))
+            .collect::<Vec<Target>>()
+            .try_into()
+            .unwrap();
+        let output_target =
+            builder.hash_n_to_m_no_pad::<Poseidon2Hash>(input_target[0..8].to_vec(), 4);
+
+        let expected_output_target = builder.add_virtual_target_arr::<4>();
+        for i in 0..4 {
+            builder.connect(expected_output_target[i], output_target[i]);
+        }
+
+        let circuit = builder.build::<PoseidonGoldilocksConfig>();
+        let mut pw = PartialWitness::new();
+        pw.set_target_arr(&expected_output_target, &expected_output)?;
+
+        let proof = circuit.prove(pw).unwrap();
+        circuit.verify(proof.clone())
+    }
+}
diff --git a/plonky2/src/hash/poseidon2/mod.rs b/plonky2/src/hash/poseidon2/mod.rs
new file mode 100644
index 000000000..52ede9a4d
--- /dev/null
+++ b/plonky2/src/hash/poseidon2/mod.rs
@@ -0,0 +1,6 @@
+pub mod config;
+pub mod gate;
+pub mod hash;
+
+#[cfg(test)]
+pub mod p3;
diff --git a/plonky2/src/hash/poseidon2/p3.rs b/plonky2/src/hash/poseidon2/p3.rs
new file mode 100644
index 000000000..345bce362
--- /dev/null
+++ b/plonky2/src/hash/poseidon2/p3.rs
@@ -0,0 +1,67 @@
+use p3_field::AbstractField;
+use p3_goldilocks::{DiffusionMatrixGoldilocks, Goldilocks};
+use p3_poseidon2::{Poseidon2, Poseidon2ExternalMatrixGeneral};
+use p3_symmetric::Permutation;
+
+use super::config::*;
+
+// Poseidon2 from plonky3
+pub fn p3_poseidon2_hash_n_to_m_no_pad(
+    inputs: &[Goldilocks],
+    num_outputs: usize,
+) -> Vec<Goldilocks> {
+    let external_linear_layer = Poseidon2ExternalMatrixGeneral;
+    let internal_linear_layer = DiffusionMatrixGoldilocks;
+
+    let external_constants = EXTERNAL_CONSTANTS
+        .iter()
+        .map(|v| {
+            v.iter()
+                .map(|&x| Goldilocks::from_canonical_u64(x))
+                .collect::<Vec<Goldilocks>>()
+                .try_into()
+                .unwrap()
+        })
+        .collect::<Vec<[Goldilocks; WIDTH]>>();
+
+    let internal_constants = INTERNAL_CONSTANTS
+        .iter()
+        .map(|&x| Goldilocks::from_canonical_u64(x))
+        .collect::<Vec<Goldilocks>>();
+
+    let poseidon = Poseidon2::<
+        Goldilocks,
+        Poseidon2ExternalMatrixGeneral,
+        DiffusionMatrixGoldilocks,
+        WIDTH,
+        D,
+    >::new(
+        ROUNDS_F,
+        external_constants,
+        external_linear_layer,
+        ROUNDS_P,
+        internal_constants,
+        internal_linear_layer,
+    );
+
+    let mut perm = [Goldilocks::zero(); WIDTH];
+
+    #[allow(clippy::manual_memcpy)]
+    for input_chunk in inputs.chunks(RATE) {
+        for i in 0..RATE.min(input_chunk.len()) {
+            perm[i] = input_chunk[i];
+        }
+        poseidon.permute_mut(&mut perm);
+    }
+
+    let mut outputs: Vec<Goldilocks> = Vec::new();
+    loop {
+        for &item in perm[0..RATE].iter() {
+            outputs.push(item);
+            if outputs.len() == num_outputs {
+                return outputs;
+            }
+        }
+        poseidon.permute_mut(&mut perm);
+    }
+}
diff --git a/plonky2/src/plonk/config.rs b/plonky2/src/plonk/config.rs
index 217c88976..4bbcbc4bc 100644
--- a/plonky2/src/plonk/config.rs
+++ b/plonky2/src/plonk/config.rs
@@ -20,6 +20,7 @@ use crate::hash::hash_types::{HashOut, RichField};
 use crate::hash::hashing::PlonkyPermutation;
 use crate::hash::keccak::KeccakHash;
 use crate::hash::poseidon::PoseidonHash;
+use crate::hash::poseidon2::hash::Poseidon2Hash;
 use crate::iop::target::{BoolTarget, Target};
 use crate::plonk::circuit_builder::CircuitBuilder;
 
@@ -51,7 +52,7 @@ pub trait Hasher<F: RichField>: Sized + Copy + Debug + Eq + PartialEq {
     fn hash_pad(input: &[F]) -> Self::Hash {
         let mut padded_input = input.to_vec();
         padded_input.push(F::ONE);
-        while (padded_input.len() + 1) % Self::Permutation::RATE != 0 {
+        while !(padded_input.len() + 1).is_multiple_of(Self::Permutation::RATE) {
             padded_input.push(F::ZERO);
         }
         padded_input.push(F::ONE);
@@ -115,6 +116,16 @@ impl GenericConfig<2> for PoseidonGoldilocksConfig {
     type InnerHasher = PoseidonHash;
 }
 
+/// Configuration using Poseidon over the Goldilocks field.
+#[derive(Debug, Copy, Clone, Default, Eq, PartialEq, Serialize)]
+pub struct Poseidon2GoldilocksConfig;
+impl GenericConfig<2> for Poseidon2GoldilocksConfig {
+    type F = GoldilocksField;
+    type FE = QuadraticExtension<Self::F>;
+    type Hasher = Poseidon2Hash;
+    type InnerHasher = Poseidon2Hash;
+}
+
 /// Configuration using truncated Keccak over the Goldilocks field.
 #[derive(Debug, Copy, Clone, Default, Eq, PartialEq)]
 pub struct KeccakGoldilocksConfig;
diff --git a/plonky2/src/recursion/conditional_recursive_verifier.rs b/plonky2/src/recursion/conditional_recursive_verifier.rs
index 0bca23f7f..6927866eb 100644
--- a/plonky2/src/recursion/conditional_recursive_verifier.rs
+++ b/plonky2/src/recursion/conditional_recursive_verifier.rs
@@ -349,28 +349,35 @@ mod tests {
 
     use anyhow::Result;
     use hashbrown::HashMap;
+    use plonky2_field::goldilocks_field::GoldilocksField;
 
     use super::*;
     use crate::field::types::Sample;
     use crate::gates::noop::NoopGate;
     use crate::iop::witness::{PartialWitness, WitnessWrite};
     use crate::plonk::circuit_data::CircuitConfig;
-    use crate::plonk::config::PoseidonGoldilocksConfig;
+    use crate::plonk::config::{Poseidon2GoldilocksConfig, PoseidonGoldilocksConfig};
     use crate::recursion::dummy_circuit::{dummy_circuit, dummy_proof};
 
     #[test]
     fn test_conditional_recursive_verifier() -> Result<()> {
         init_logger();
-        const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
+        conditional_recursive_verifier::<PoseidonGoldilocksConfig>()?;
+        conditional_recursive_verifier::<Poseidon2GoldilocksConfig>()
+    }
+
+    fn conditional_recursive_verifier<C: GenericConfig<2, F = GoldilocksField> + 'static>(
+    ) -> Result<()>
+    where
+        C::Hasher: AlgebraicHasher<GoldilocksField>,
+    {
         let config = CircuitConfig::standard_recursion_config();
 
         // Generate proof.
-        let mut builder = CircuitBuilder::<F, D>::new(config.clone());
+        let mut builder = CircuitBuilder::<C::F, 2>::new(config.clone());
         let mut pw = PartialWitness::new();
         let t = builder.add_virtual_target();
-        pw.set_target(t, F::rand())?;
+        pw.set_target(t, C::F::rand())?;
         builder.register_public_input(t);
         let _t2 = builder.square(t);
         for _ in 0..64 {
@@ -385,19 +392,19 @@ mod tests {
         let dummy_proof = dummy_proof(&dummy_data, HashMap::new())?;
 
         // Conditionally verify the two proofs.
-        let mut builder = CircuitBuilder::<F, D>::new(config);
+        let mut builder = CircuitBuilder::<C::F, 2>::new(config);
         let mut pw = PartialWitness::new();
         let pt = builder.add_virtual_proof_with_pis(&data.common);
         pw.set_proof_with_pis_target(&pt, &proof)?;
         let dummy_pt = builder.add_virtual_proof_with_pis(&data.common);
-        pw.set_proof_with_pis_target::<C, D>(&dummy_pt, &dummy_proof)?;
+        pw.set_proof_with_pis_target::<C, 2>(&dummy_pt, &dummy_proof)?;
         let inner_data =
             builder.add_virtual_verifier_data(data.common.config.fri_config.cap_height);
         pw.set_verifier_data_target(&inner_data, &data.verifier_only)?;
         let dummy_inner_data =
             builder.add_virtual_verifier_data(data.common.config.fri_config.cap_height);
         pw.set_verifier_data_target(&dummy_inner_data, &dummy_data.verifier_only)?;
-        let b = builder.constant_bool(F::rand().0 % 2 == 0);
+        let b = builder.constant_bool(C::F::rand().0 % 2 == 0);
         builder.conditionally_verify_proof::<C>(
             b,
             &pt,
diff --git a/plonky2/src/recursion/cyclic_recursion.rs b/plonky2/src/recursion/cyclic_recursion.rs
index df0fb95cd..956a8dba0 100644
--- a/plonky2/src/recursion/cyclic_recursion.rs
+++ b/plonky2/src/recursion/cyclic_recursion.rs
@@ -206,14 +206,18 @@ mod tests {
 
     use crate::field::extension::Extendable;
     use crate::field::types::{Field, PrimeField64};
+    use crate::gates::constant::ConstantGate;
     use crate::gates::noop::NoopGate;
     use crate::hash::hash_types::{HashOutTarget, RichField};
     use crate::hash::hashing::hash_n_to_hash_no_pad;
-    use crate::hash::poseidon::{PoseidonHash, PoseidonPermutation};
+    use crate::hash::poseidon::PoseidonPermutation;
+    use crate::hash::poseidon2::hash::{Poseidon2, Poseidon2Permutation};
     use crate::iop::witness::{PartialWitness, WitnessWrite};
     use crate::plonk::circuit_builder::CircuitBuilder;
     use crate::plonk::circuit_data::{CircuitConfig, CommonCircuitData};
-    use crate::plonk::config::{AlgebraicHasher, GenericConfig, PoseidonGoldilocksConfig};
+    use crate::plonk::config::{
+        AlgebraicHasher, GenericConfig, Poseidon2GoldilocksConfig, PoseidonGoldilocksConfig,
+    };
     use crate::recursion::cyclic_recursion::check_cyclic_proof_verifier_data;
     use crate::recursion::dummy_circuit::cyclic_base_proof;
 
@@ -222,7 +226,9 @@ mod tests {
         F: RichField + Extendable<D>,
         C: GenericConfig<D, F = F>,
         const D: usize,
-    >() -> CommonCircuitData<F, D>
+    >(
+        use_poseidon2: bool,
+    ) -> CommonCircuitData<F, D>
     where
         C::Hasher: AlgebraicHasher<F>,
     {
@@ -243,6 +249,17 @@ mod tests {
         let verifier_data =
             builder.add_virtual_verifier_data(data.common.config.fri_config.cap_height);
         builder.verify_proof::<C>(&proof, &verifier_data, &data.common);
+
+        if use_poseidon2 {
+            // Add a constant to ensure ConstantGate is included in the gate set
+            // This matches what the actual cyclic recursion circuit will do
+            builder.add_gate(
+                ConstantGate {
+                    num_consts: builder.config.num_constants,
+                },
+                vec![],
+            );
+        }
         while builder.num_gates() < 1 << 12 {
             builder.add_gate(NoopGate, vec![]);
         }
@@ -257,12 +274,19 @@ mod tests {
     /// - VK for cyclic recursion (?)
     #[test]
     fn test_cyclic_recursion() -> Result<()> {
-        const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
+        cyclic_recursion::<PoseidonGoldilocksConfig, 2>(false)?;
+        cyclic_recursion::<Poseidon2GoldilocksConfig, 2>(true)
+    }
 
+    fn cyclic_recursion<C: GenericConfig<D> + 'static, const D: usize>(
+        use_poseidon2: bool,
+    ) -> Result<()>
+    where
+        C::Hasher: AlgebraicHasher<C::F>,
+        C::F: Poseidon2,
+    {
         let config = CircuitConfig::standard_recursion_config();
-        let mut builder = CircuitBuilder::<F, D>::new(config);
+        let mut builder = CircuitBuilder::<C::F, D>::new(config);
         let one = builder.one();
 
         // Circuit that computes a repeated hash.
@@ -270,11 +294,11 @@ mod tests {
         builder.register_public_inputs(&initial_hash_target.elements);
         let current_hash_in = builder.add_virtual_hash();
         let current_hash_out =
-            builder.hash_n_to_hash_no_pad::<PoseidonHash>(current_hash_in.elements.to_vec());
+            builder.hash_n_to_hash_no_pad::<C::InnerHasher>(current_hash_in.elements.to_vec());
         builder.register_public_inputs(&current_hash_out.elements);
         let counter = builder.add_virtual_public_input();
 
-        let mut common_data = common_data_for_recursion::<F, C, D>();
+        let mut common_data = common_data_for_recursion::<C::F, C, D>(use_poseidon2);
         let verifier_data_target = builder.add_verifier_data_public_inputs();
         common_data.num_public_inputs = builder.num_public_inputs();
 
@@ -310,7 +334,12 @@ mod tests {
         let cyclic_circuit_data = builder.build::<C>();
 
         let mut pw = PartialWitness::new();
-        let initial_hash = [F::ZERO, F::ONE, F::TWO, F::from_canonical_usize(3)];
+        let initial_hash = [
+            C::F::ZERO,
+            C::F::ONE,
+            C::F::TWO,
+            C::F::from_canonical_usize(3),
+        ];
         let initial_hash_pis = initial_hash.into_iter().enumerate().collect();
         pw.set_bool_target(condition, false)?;
         pw.set_proof_with_pis_target::<C, D>(
@@ -359,19 +388,31 @@ mod tests {
         let initial_hash = &proof.public_inputs[..4];
         let hash = &proof.public_inputs[4..8];
         let counter = proof.public_inputs[8];
-        let expected_hash: [F; 4] = iterate_poseidon(
+        let expected_hash: [C::F; 4] = iterate_poseidon(
             initial_hash.try_into().unwrap(),
             counter.to_canonical_u64() as usize,
+            use_poseidon2,
         );
         assert_eq!(hash, expected_hash);
 
         cyclic_circuit_data.verify(proof)
     }
 
-    fn iterate_poseidon<F: RichField>(initial_state: [F; 4], n: usize) -> [F; 4] {
+    fn iterate_poseidon<F: RichField + Poseidon2>(
+        initial_state: [F; 4],
+        n: usize,
+        use_poseidon2: bool,
+    ) -> [F; 4] {
         let mut current = initial_state;
-        for _ in 0..n {
-            current = hash_n_to_hash_no_pad::<F, PoseidonPermutation<F>>(&current).elements;
+
+        if use_poseidon2 {
+            for _ in 0..n {
+                current = hash_n_to_hash_no_pad::<F, Poseidon2Permutation<F>>(&current).elements;
+            }
+        } else {
+            for _ in 0..n {
+                current = hash_n_to_hash_no_pad::<F, PoseidonPermutation<F>>(&current).elements;
+            }
         }
         current
     }
diff --git a/plonky2/src/recursion/dummy_circuit.rs b/plonky2/src/recursion/dummy_circuit.rs
index cd8559382..e9bdd4641 100644
--- a/plonky2/src/recursion/dummy_circuit.rs
+++ b/plonky2/src/recursion/dummy_circuit.rs
@@ -113,6 +113,7 @@ pub fn dummy_circuit<F: RichField + Extendable<D>, C: GenericConfig<D, F = F>, c
     }
 
     let circuit = builder.build::<C>();
+
     assert_eq!(&circuit.common, common_data);
     circuit
 }
diff --git a/plonky2/src/recursion/recursive_verifier.rs b/plonky2/src/recursion/recursive_verifier.rs
index 16a8ba85b..ebc0af9ca 100644
--- a/plonky2/src/recursion/recursive_verifier.rs
+++ b/plonky2/src/recursion/recursive_verifier.rs
@@ -205,6 +205,7 @@ mod tests {
     use anyhow::Result;
     use itertools::Itertools;
     use log::{info, Level};
+    use plonky2_field::goldilocks_field::GoldilocksField;
 
     use super::*;
     use crate::fri::reduction_strategies::FriReductionStrategy;
@@ -214,7 +215,9 @@ mod tests {
     use crate::gates::noop::NoopGate;
     use crate::iop::witness::{PartialWitness, WitnessWrite};
     use crate::plonk::circuit_data::{CircuitConfig, VerifierOnlyCircuitData};
-    use crate::plonk::config::{KeccakGoldilocksConfig, PoseidonGoldilocksConfig};
+    use crate::plonk::config::{
+        KeccakGoldilocksConfig, Poseidon2GoldilocksConfig, PoseidonGoldilocksConfig,
+    };
     use crate::plonk::proof::{CompressedProofWithPublicInputs, ProofWithPublicInputs};
     use crate::plonk::prover::prove;
     use crate::util::timing::TimingTree;
@@ -223,14 +226,27 @@ mod tests {
     fn test_recursive_verifier() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
-        let config = CircuitConfig::standard_recursion_zk_config();
+        {
+            type C = PoseidonGoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_zk_config();
+
+            let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 4_000)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
 
-        let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 4_000)?;
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
-        test_serialization(&proof, &vd, &common_data)?;
+        {
+            type C = Poseidon2GoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_zk_config();
+
+            let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 4_000)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
 
         Ok(())
     }
@@ -239,14 +255,26 @@ mod tests {
     fn test_recursive_verifier_one_lookup() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
-        let config = CircuitConfig::standard_recursion_zk_config();
-
-        let (proof, vd, common_data) = dummy_lookup_proof::<F, C, D>(&config, 10)?;
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
-        test_serialization(&proof, &vd, &common_data)?;
+        {
+            type C = PoseidonGoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_zk_config();
+
+            let (proof, vd, common_data) = dummy_lookup_proof::<F, C, D>(&config, 10)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
+        {
+            type C = Poseidon2GoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_zk_config();
+
+            let (proof, vd, common_data) = dummy_lookup_proof::<F, C, D>(&config, 10)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
 
         Ok(())
     }
@@ -255,15 +283,26 @@ mod tests {
     fn test_recursive_verifier_two_luts() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
-        let config = CircuitConfig::standard_recursion_config();
-
-        let (proof, vd, common_data) = dummy_two_luts_proof::<F, C, D>(&config)?;
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
-        test_serialization(&proof, &vd, &common_data)?;
-
+        {
+            type C = PoseidonGoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_config();
+
+            let (proof, vd, common_data) = dummy_two_luts_proof::<F, C, D>(&config)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
+        {
+            type C = Poseidon2GoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_config();
+
+            let (proof, vd, common_data) = dummy_two_luts_proof::<F, C, D>(&config)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
         Ok(())
     }
 
@@ -271,15 +310,26 @@ mod tests {
     fn test_recursive_verifier_too_many_rows() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
-        let config = CircuitConfig::standard_recursion_config();
-
-        let (proof, vd, common_data) = dummy_too_many_rows_proof::<F, C, D>(&config)?;
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
-        test_serialization(&proof, &vd, &common_data)?;
-
+        {
+            type C = PoseidonGoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_config();
+
+            let (proof, vd, common_data) = dummy_too_many_rows_proof::<F, C, D>(&config)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
+        {
+            type C = Poseidon2GoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+            let config = CircuitConfig::standard_recursion_config();
+
+            let (proof, vd, common_data) = dummy_too_many_rows_proof::<F, C, D>(&config)?;
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
         Ok(())
     }
 
@@ -287,27 +337,65 @@ mod tests {
     fn test_recursive_recursive_verifier() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
-
-        let config = CircuitConfig::standard_recursion_config();
-
-        // Start with a degree 2^14 proof
-        let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 16_000)?;
-        assert_eq!(common_data.degree_bits(), 14);
-
-        // Shrink it to 2^13.
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, Some(13), false, false)?;
-        assert_eq!(common_data.degree_bits(), 13);
-
-        // Shrink it to 2^12.
-        let (proof, vd, common_data) =
-            recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
-        assert_eq!(common_data.degree_bits(), 12);
-
-        test_serialization(&proof, &vd, &common_data)?;
+        {
+            type C = PoseidonGoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+
+            let config = CircuitConfig::standard_recursion_config();
+
+            // Start with a degree 2^14 proof
+            let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 16_000)?;
+            assert_eq!(common_data.degree_bits(), 14);
+
+            // Shrink it to 2^13.
+            let (proof, vd, common_data) = recursive_proof::<F, C, C, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                Some(13),
+                false,
+                false,
+            )?;
+            assert_eq!(common_data.degree_bits(), 13);
+
+            // Shrink it to 2^12.
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            assert_eq!(common_data.degree_bits(), 12);
+
+            test_serialization(&proof, &vd, &common_data)?;
+        }
 
+        {
+            type C = Poseidon2GoldilocksConfig;
+            type F = <C as GenericConfig<D>>::F;
+
+            let config = CircuitConfig::standard_recursion_config();
+
+            // Start with a degree 2^14 proof
+            let (proof, vd, common_data) = dummy_proof::<F, C, D>(&config, 16_000)?;
+            assert_eq!(common_data.degree_bits(), 14);
+
+            // Shrink it to 2^13.
+            let (proof, vd, common_data) = recursive_proof::<F, C, C, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                Some(13),
+                false,
+                false,
+            )?;
+            assert_eq!(common_data.degree_bits(), 13);
+
+            // Shrink it to 2^12.
+            let (proof, vd, common_data) =
+                recursive_proof::<F, C, C, D>(proof, vd, common_data, &config, None, true, true)?;
+            assert_eq!(common_data.degree_bits(), 12);
+
+            test_serialization(&proof, &vd, &common_data)?;
+        }
         Ok(())
     }
 
@@ -316,20 +404,26 @@ mod tests {
     #[test]
     #[ignore]
     fn test_size_optimized_recursion() -> Result<()> {
+        size_optimized_recursion::<PoseidonGoldilocksConfig>()?;
+        size_optimized_recursion::<Poseidon2GoldilocksConfig>()
+    }
+
+    fn size_optimized_recursion<C: GenericConfig<2, F = GoldilocksField>>() -> Result<()>
+    where
+        C::Hasher: AlgebraicHasher<GoldilocksField>,
+    {
         init_logger();
         const D: usize = 2;
-        type C = PoseidonGoldilocksConfig;
         type KC = KeccakGoldilocksConfig;
-        type F = <C as GenericConfig<D>>::F;
 
         let standard_config = CircuitConfig::standard_recursion_config();
 
         // An initial dummy proof.
-        let (proof, vd, common_data) = dummy_proof::<F, C, D>(&standard_config, 4_000)?;
+        let (proof, vd, common_data) = dummy_proof::<C::F, C, D>(&standard_config, 4_000)?;
         assert_eq!(common_data.degree_bits(), 12);
 
         // A standard recursive proof.
-        let (proof, vd, common_data) = recursive_proof::<F, C, C, D>(
+        let (proof, vd, common_data) = recursive_proof::<C::F, C, C, D>(
             proof,
             vd,
             common_data,
@@ -350,7 +444,7 @@ mod tests {
             },
             ..standard_config
         };
-        let (proof, vd, common_data) = recursive_proof::<F, C, C, D>(
+        let (proof, vd, common_data) = recursive_proof::<C::F, C, C, D>(
             proof,
             vd,
             common_data,
@@ -373,7 +467,7 @@ mod tests {
             },
             ..high_rate_config
         };
-        let (proof, vd, common_data) = recursive_proof::<F, KC, C, D>(
+        let (proof, vd, common_data) = recursive_proof::<C::F, KC, C, D>(
             proof,
             vd,
             common_data,
@@ -393,21 +487,66 @@ mod tests {
     fn test_recursive_verifier_multi_hash() -> Result<()> {
         init_logger();
         const D: usize = 2;
-        type PC = PoseidonGoldilocksConfig;
         type KC = KeccakGoldilocksConfig;
-        type F = <PC as GenericConfig<D>>::F;
-
-        let config = CircuitConfig::standard_recursion_config();
-        let (proof, vd, common_data) = dummy_proof::<F, PC, D>(&config, 4_000)?;
-
-        let (proof, vd, common_data) =
-            recursive_proof::<F, PC, PC, D>(proof, vd, common_data, &config, None, false, false)?;
-        test_serialization(&proof, &vd, &common_data)?;
-
-        let (proof, vd, common_data) =
-            recursive_proof::<F, KC, PC, D>(proof, vd, common_data, &config, None, false, false)?;
-        test_serialization(&proof, &vd, &common_data)?;
 
+        {
+            type PC = PoseidonGoldilocksConfig;
+            type F = <PC as GenericConfig<D>>::F;
+
+            let config = CircuitConfig::standard_recursion_config();
+            let (proof, vd, common_data) = dummy_proof::<F, PC, D>(&config, 4_000)?;
+
+            let (proof, vd, common_data) = recursive_proof::<F, PC, PC, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                None,
+                false,
+                false,
+            )?;
+            test_serialization(&proof, &vd, &common_data)?;
+
+            let (proof, vd, common_data) = recursive_proof::<F, KC, PC, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                None,
+                false,
+                false,
+            )?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
+        {
+            type PC = Poseidon2GoldilocksConfig;
+            type F = <PC as GenericConfig<D>>::F;
+
+            let config = CircuitConfig::standard_recursion_config();
+            let (proof, vd, common_data) = dummy_proof::<F, PC, D>(&config, 4_000)?;
+
+            let (proof, vd, common_data) = recursive_proof::<F, PC, PC, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                None,
+                false,
+                false,
+            )?;
+            test_serialization(&proof, &vd, &common_data)?;
+
+            let (proof, vd, common_data) = recursive_proof::<F, KC, PC, D>(
+                proof,
+                vd,
+                common_data,
+                &config,
+                None,
+                false,
+                false,
+            )?;
+            test_serialization(&proof, &vd, &common_data)?;
+        }
         Ok(())
     }
 
diff --git a/plonky2/src/util/reducing.rs b/plonky2/src/util/reducing.rs
index 89c8bd96c..1abe017b0 100644
--- a/plonky2/src/util/reducing.rs
+++ b/plonky2/src/util/reducing.rs
@@ -150,7 +150,7 @@ impl<const D: usize> ReducingFactorTarget<D> {
         let zero_ext = builder.zero_extension();
         let mut acc = zero_ext;
         let mut reversed_terms = terms.to_vec();
-        while reversed_terms.len() % max_coeffs_len != 0 {
+        while !reversed_terms.len().is_multiple_of(max_coeffs_len) {
             reversed_terms.push(zero);
         }
         reversed_terms.reverse();
@@ -200,7 +200,7 @@ impl<const D: usize> ReducingFactorTarget<D> {
         let zero_ext = builder.zero_extension();
         let mut acc = zero_ext;
         let mut reversed_terms = terms.to_vec();
-        while reversed_terms.len() % max_coeffs_len != 0 {
+        while !reversed_terms.len().is_multiple_of(max_coeffs_len) {
             reversed_terms.push(zero_ext);
         }
         reversed_terms.reverse();