refactor(stark): unify & clean up the commitment layer (#735)

* refactor(stark): extract Merkle commitment into a commitment module

Unifies the two near-identical bit-reversed leaf hashers (per-row + per-row-pair)
into one keccak_leaves_bit_reversed_grouped(columns, rows_per_leaf), and the two
near-identical commit fns (commit_columns_bit_reversed + commit_composition_polynomial)
into commit_bit_reversed(columns, rows_per_leaf), in a new crypto/stark/src/commitment.rs.
Removes them from the IsStarkProver trait (they used no self). prover.rs re-exports
the named leaf hashers for the math-cuda GPU parity test. Byte-identical (stark 128/128).

* perf(stark): row-pair the trace commitment (one Merkle path per query)

Every FRI query opens a value and its symmetric counterpart (LDE positions
2*iota, 2*iota+1). The composition-poly commitment already grouped that pair into
one leaf; the trace (main/aux/precomputed) committed one row per leaf and opened
TWO leaves (proof + proof_sym) per query. Commit the trace with rows_per_leaf=2
too, so each query opens ONE leaf/path; drop the now-redundant proof_sym from
PolynomialOpenings (also shrinks the composition opening, which stored the path
twice). verify_opening_pair now reconstructs the paired leaf and verifies once
(mirrors verify_composition_poly_opening); the dead verify_opening is removed.

Halves trace Merkle authentication-path data per query (smaller proofs + less
verifier hashing). stark 128/128 (prove+verify). NOTE: proof FORMAT change (not
byte-identical). cuda follow-up: the GPU trace leaf+tree builders
(gpu_lde::try_expand_leaf_and_tree_batched_keep/_ext3_keep + math-cuda kernels)
still build 1-row leaves and must switch to the row-pair pattern (the GPU
composition builder already pairs) or cuda proofs will fail verification.

* refactor(stark): prover cleanup — par helpers, ROWS_PER_LEAF, error propagation, doc/log fixes

* docs(stark): fix commitment.rs leaf-layout docs after trace pairing

The trace commitment now uses the row-pair leaf layout (ROWS_PER_LEAF=2), same as
composition; rows_per_leaf=1 is only kept for the GPU parity test. Update the
module/const/wrapper docs that still described the pre-pairing per-row trace.

* refactor(prover): dedup commit pipeline (commit_plain + spill_tree) (B)

Extract two helpers on IsStarkProver, collapsing the near-duplicate
main-trace and aux-trace commit code:

- spill_tree<C>: the identical-except-label disk-spill block, shared by
  the main / preprocessed-split / aux commit sites (4 call sites).
- commit_plain<C>: commit_bit_reversed + spill_tree + TableCommit::plain,
  shared by the main-trace (non-preprocessed None arm) and aux-trace
  plain-commit paths.

Proof output is byte-identical (stark 128/128, +disk-spill 133/133).
The only behavioral delta is in the instruments profiling feature: the
aux commit's timing bucket now includes its (tiny) disk spill, matching
what the main path already measured.

clippy clean on default / disk-spill / instruments; builds on the
combined feature set.

* fix(prover): row-pair the preprocessed-table commitments (CI fix)

The trace Merkle commitment moved to a row-pair leaf layout
(ROWS_PER_LEAF=2), but the 5 preprocessed-table commitment computers
(bitwise/keccak_rc/page/decode/register) still built a 1-row-per-leaf
tree manually, so the prover's row-pair precomputed root no longer
matched the computed/hardcoded one -> PrecomputedCommitmentMismatch at
prove time.

- Route all 5 compute_*commitment fns through the shared
  stark::commitment::commit_bit_reversed(.., ROWS_PER_LEAF), dropping the
  manual bit-reverse + columns2rows + 1-row BatchedMerkleTree::build.
- Regenerate the hardcoded bitwise/keccak_rc/zero_page static commitments
  for the row-pair layout (via compute_static_commitments).
- cargo fmt (prover.rs + touched tables).

Verified: static_commitments drift tests 5/5, stark 128/128, clippy +
fmt clean, no PrecomputedCommitmentMismatch. Full ELF prove/verify runs
in CI (guest artifacts absent locally).

* test(prover): regenerate SUB_DECODE_COMMITMENT_BLOWUP_2 for row-pair layout

The compile-time decode-commitment const for sub.elf shifted with the
row-pair preprocessed commitment; regenerated via commitment_from_elf
(the print_decode_commitment_for_sub regen path).

* test(prover): TEMP print actual decode commitment to regenerate const from CI

* test(prover): set SUB_DECODE_COMMITMENT_BLOWUP_2 to CI-computed row-pair value

Regenerated from CI's sub.elf (local riscv toolchain unavailable);
removed the temporary print instrumentation.

* mplement changes in GPU

* refactor

* Fix CUDA LDE clippy lint

* fix(cuda): align review cleanup with row-pair commits (#723)

* fix(stark): silence dead_code on par_for_each_mut (debug-checks-only after merge)

* style(stark): cargo fmt after merge resolution

* test(cuda): focused GPU row-pair commitment prove+verify test

* fix(cuda): drop stale R2 parts-LDE asserts (#700 fused path), silence GPU column-LDE dead_code

The R2 parts-LDE / comp-poly-tree GPU dispatches no longer fire since #699/#700
route degree-3 tables through the 2-part fused coset_lde_full path (no AIR has
number_of_parts > 2). try_expand_columns_batched* are debug-checks-only after
the #650 row-major LDE became production.

* review(stark): address PR #735 review — coverage, dead code, cleanups (#740)

Follow-up to the row-pair commitment PR. Excludes the intentional
proof-format break (proof_sym removal / leaf-count change), which is
by design.

Test coverage (in their own files under crypto/stark/src/tests/, per the
crate's test structure):
- tests/commitment_tests.rs: direct unit tests pinning the row-pair leaf
  layout (R=1 and R=2) against an independent reference, wrapper
  agreement, commit-root consistency, and empty-input short-circuit.
  Previously the leaf layout was only covered transitively via full
  prove->verify, and GPU parity tests compared against an inline reimpl
  rather than this module.
- tests/row_pair_opening_tests.rs: two negative tests for the row-pair
  verify_opening_pair — a tampered symmetric trace evaluation and a
  corrupted Merkle authentication path must both be rejected. Removing
  proof_sym deleted the old "symmetric opening mismatch" rejection class;
  these restore it (an impl ignoring evaluations_sym / the auth path
  would otherwise pass every existing test). Reuses the now pub(crate)
  make_valid_simple_proof helper.
- cuda_path_integration.rs: restore assert!(gpu_comp_poly_tree_calls() > 0).
  try_build_comp_poly_tree_gpu is dispatched unconditionally (round 2,
  after the parts-count branch), so it fires for the common
  number_of_parts == 2 (degree-3) case — it was NOT obsolete. Keep the
  genuinely-dead parts-LDE assertion dropped.

Cleanups:
- Delete dead fn commit_plain (zero callers; main/aux commit inline
  commit_rows_bit_reversed + spill_tree, which are row-major and
  incompatible with its column-major signature — the dedup never landed).
- Delete orphaned pub fn columns2rows (all callers removed by #735).
- Add ProvingError::Fft and map FFTError to it instead of WrongParameter
  (internal FFT failure is not a caller-supplied-parameter error).
- Fix stale profiler label commit_composition_poly -> commit_bit_reversed.
- Drop the rot-prone "// = 2" comment on the local ROWS_PER_LEAF alias.

* review(stark): tidy test layout + remove AGENTS.md (follow-up to #740) (#744)

Two items that landed too late for #740:

- Move the shared make_valid_simple_proof helper out of small_trace_tests.rs
  into tests/trace_test_helpers.rs, where the crate keeps shared test
  helpers (matching how prover_tests sources get_trace_evaluations).
  row_pair_opening_tests.rs and small_trace_tests.rs now both import it
  from there instead of one test file reaching sideways into another.
- Delete AGENTS.md (added by #735).

---------

Co-authored-by: Joaquin Carletti <joaquin.carletti@lambdaclass.com>
Co-authored-by: MauroFab <maurotoscano2@gmail.com>
Co-authored-by: Mauro Toscano <12560266+MauroToscano@users.noreply.github.com>

Author: 3 people

crypto/math-cuda/kernels/keccak.cu

@@ -159,8 +159,8 @@ extern "C" __global__ void keccak256_leaves_base_batched(
     uint64_t tid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x;
     if (tid >= num_rows) return;
 
-    // Bit-reverse the row index so we read columns at `br` but write the
-    // hashed leaf at `tid` — matching the CPU `commit_columns_bit_reversed`.
+    // Bit-reverse the row index so we read columns at `br` but write the hashed
+    // leaf at `tid` — matching the CPU per-row `commit_bit_reversed(.., 1)`.
     uint64_t br = __brevll(tid) >> (64 - log_num_rows);
 
     uint64_t st[25];
@@ -181,6 +181,51 @@ extern "C" __global__ void keccak256_leaves_base_batched(
     finalize_keccak256(st, rate_pos, hashed_leaves_out + tid * 32);
 }
 
+// ---------------------------------------------------------------------------
+// Goldilocks BASE-FIELD row-pair leaf hashing.
+//
+// Leaf `leaf_idx` hashes TWO consecutive bit-reversed rows
+//   br_0 = reverse_index(2*leaf_idx),  br_1 = reverse_index(2*leaf_idx + 1)
+// each written column-by-column in canonical BE (same per-row byte layout as
+// `keccak256_leaves_base_batched`), in (br_0 row: col 0..K-1) then (br_1 row:
+// col 0..K-1) order. `num_leaves = num_rows / 2`; writes 32 bytes to
+// `hashed_leaves_out[leaf_idx * 32 ..]`. Matches the CPU
+// `keccak_leaves_row_pair_bit_reversed` (rows_per_leaf = 2) — the base-field
+// analog of `keccak_comp_poly_leaves_ext3`.
+// ---------------------------------------------------------------------------
+extern "C" __global__ void keccak256_leaves_base_row_pair_batched(
+    const uint64_t *columns_base_ptr,
+    uint64_t col_stride,
+    uint64_t num_cols,
+    uint64_t num_rows,
+    uint64_t log_num_rows,
+    uint8_t *hashed_leaves_out) {
+    uint64_t tid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x;
+    uint64_t num_leaves = num_rows >> 1;
+    if (tid >= num_leaves) return;
+
+    uint64_t br_0 = __brevll(2 * tid) >> (64 - log_num_rows);
+    uint64_t br_1 = __brevll(2 * tid + 1) >> (64 - log_num_rows);
+
+    uint64_t st[25];
+    #pragma unroll
+    for (int i = 0; i < 25; ++i) st[i] = 0;
+
+    uint32_t rate_pos = 0;
+    // First row (br_0): col 0..K-1.
+    for (uint64_t c = 0; c < num_cols; ++c) {
+        uint64_t v = columns_base_ptr[c * col_stride + br_0];
+        absorb_lane(st, rate_pos, bswap64(goldilocks::canonical(v)));
+    }
+    // Second row (br_1): col 0..K-1.
+    for (uint64_t c = 0; c < num_cols; ++c) {
+        uint64_t v = columns_base_ptr[c * col_stride + br_1];
+        absorb_lane(st, rate_pos, bswap64(goldilocks::canonical(v)));
+    }
+
+    finalize_keccak256(st, rate_pos, hashed_leaves_out + tid * 32);
+}
+
 // ---------------------------------------------------------------------------
 // Goldilocks EXT3 leaf hashing (3 base-field components per ext3 element).
 //
@@ -349,35 +394,49 @@ extern "C" __global__ void keccak_merkle_level(
 }
 
 // ---------------------------------------------------------------------------
-// Row-major base leaf hashing.
+// Row-major ROW-PAIR leaf hashing.
+//
+// Row-major analog of `keccak256_leaves_base_row_pair_batched` (which reads a
+// column-major slab): each leaf hashes TWO consecutive bit-reversed rows.
+// Leaf `tid` hashes row `reverse_index(2*tid)` followed by row
+// `reverse_index(2*tid + 1)`, each as `m` canonical big-endian lanes read from
+// the contiguous row-major buffer (`data + br * m`). `num_leaves = num_rows/2`;
+// writes 32 bytes to `hashed_leaves_out[tid*32 ..]`.
 //
-// Input layout: data[row * m + col] for `num_rows` rows and `m` columns.
-// For leaf `tid`, reads the bit-reversed row `br(tid)` — a contiguous slice
-// of `m` elements starting at data[br * m]. Coalesced when multiple threads
-// in the same warp process consecutive `tid` values (they read non-overlapping
-// rows, each a contiguous block of m u64s in order).
+// `m` is the row stride in u64s: base trace = num columns; ext3 trace = 3 *
+// num columns (an ext3 element's components c0,c1,c2 are consecutive, matching
+// the CPU `write_bytes_be`). Byte layout therefore equals the CPU
+// `commit_bit_reversed(.., ROWS_PER_LEAF=2)` and the verifier's
+// `verify_opening_pair` (queried row ‖ its symmetric counterpart, one leaf).
 // ---------------------------------------------------------------------------
-extern "C" __global__ void keccak256_leaves_base_row_major(
+extern "C" __global__ void keccak256_leaves_base_row_major_row_pair(
     const uint64_t *data,
     uint64_t m,
     uint64_t num_rows,
     uint64_t log_num_rows,
     uint8_t *hashed_leaves_out)
 {
     uint64_t tid = (uint64_t)blockIdx.x * blockDim.x + threadIdx.x;
-    if (tid >= num_rows) return;
-    uint64_t br = __brevll(tid) >> (64 - log_num_rows);
-    const uint64_t *row = data + br * m;
+    uint64_t num_leaves = num_rows >> 1;
+    if (tid >= num_leaves) return;
+
+    uint64_t br_0 = __brevll(2 * tid) >> (64 - log_num_rows);
+    uint64_t br_1 = __brevll(2 * tid + 1) >> (64 - log_num_rows);
+    const uint64_t *row_0 = data + br_0 * m;
+    const uint64_t *row_1 = data + br_1 * m;
 
     uint64_t st[25];
     #pragma unroll
     for (int i = 0; i < 25; ++i) st[i] = 0;
 
     uint32_t rate_pos = 0;
+    // First row (br_0): cols 0..m-1.
     for (uint64_t c = 0; c < m; ++c) {
-        uint64_t canon = goldilocks::canonical(row[c]);
-        uint64_t lane  = bswap64(canon);
-        absorb_lane(st, rate_pos, lane);
+        absorb_lane(st, rate_pos, bswap64(goldilocks::canonical(row_0[c])));
+    }
+    // Second row (br_1): cols 0..m-1.
+    for (uint64_t c = 0; c < m; ++c) {
+        absorb_lane(st, rate_pos, bswap64(goldilocks::canonical(row_1[c])));
     }
     finalize_keccak256(st, rate_pos, hashed_leaves_out + tid * 32);
 }

crypto/math-cuda/src/device.rs

@@ -147,8 +147,9 @@ pub struct Backend {
     pub matrix_transpose_strided: CudaFunction,
 
     // keccak.ptx
-    pub keccak256_leaves_base_row_major: CudaFunction,
+    pub keccak256_leaves_base_row_major_row_pair: CudaFunction,
     pub keccak256_leaves_base_batched: CudaFunction,
+    pub keccak256_leaves_base_row_pair_batched: CudaFunction,
     pub keccak256_leaves_ext3_batched: CudaFunction,
     pub keccak_comp_poly_leaves_ext3: CudaFunction,
     pub keccak_fri_leaves_ext3: CudaFunction,
@@ -247,9 +248,11 @@ impl Backend {
             ntt_dit_level_row_major: ntt.load_function("ntt_dit_level_row_major")?,
             pointwise_mul_row_major: ntt.load_function("pointwise_mul_row_major")?,
             matrix_transpose_strided: ntt.load_function("matrix_transpose_strided")?,
-            keccak256_leaves_base_row_major: keccak
-                .load_function("keccak256_leaves_base_row_major")?,
+            keccak256_leaves_base_row_major_row_pair: keccak
+                .load_function("keccak256_leaves_base_row_major_row_pair")?,
             keccak256_leaves_base_batched: keccak.load_function("keccak256_leaves_base_batched")?,
+            keccak256_leaves_base_row_pair_batched: keccak
+                .load_function("keccak256_leaves_base_row_pair_batched")?,
             keccak256_leaves_ext3_batched: keccak.load_function("keccak256_leaves_ext3_batched")?,
             keccak_comp_poly_leaves_ext3: keccak.load_function("keccak_comp_poly_leaves_ext3")?,
             keccak_fri_leaves_ext3: keccak.load_function("keccak_fri_leaves_ext3")?,