test(DON'T MERGE): develop-v2.0.0-rc.1 by shuklaayush · Pull Request #2592 · openvm-org/openvm

shuklaayush · 2026-03-19T18:14:32Z

No description provided.

Resolves INT-7422.

Resolves INT-7433.

## Summary Refactors the SHA-2 block hasher CUDA tracegen kernel from a monolithic 1-thread-per-block design into a three-phase pipeline for better memory coalescing on column-major traces. | Phase | Kernel | Parallelism | Purpose | |-------|--------------------------|-----------------|----------------------------------------------------------------| | 1 | `sha2_first_pass_phase1` | 1 thread/block | Compute SHA-2 rounds, snapshot `{a..h, w_buf}` to scratch | | 2 | `sha2_first_pass_phase2` | 1 thread/row | Read snapshot, write trace columns (coalesced) | | 3 | `sha2_first_pass_phase3` | 1 thread/block | Cross-row dependencies (`intermed_4/8/12`, carry) | Also drops unused parameters (`ptr_max_bits`, `range_checker_ptr`/`num_bins`, `timestamp_max_bits`) from the block hasher kernel and its Rust-side wrappers, simplifying `Sha2BlockHasherChipGpu`. ## Changes - `extensions/sha2/circuit/cuda/src/sha2_hasher.cu`: split first-pass tracegen into three phases; tightened `CHECK_KERNEL` error propagation in `launch_sha2_second_pass_dependencies`. - `extensions/sha2/circuit/src/cuda/{mod.rs,cuda_abi.rs}`: wire the new scratch buffer and drop unused wrapper arguments. - `extensions/sha2/circuit/src/extension/cuda.rs`: update `Sha2BlockHasherChipGpu::new` call sites. - `extensions/sha2/circuit/src/sha2_chips/tests.rs`: update tests to match the new constructor signature. ## Correctness notes - Phase 1 snapshots state before each row's rounds; Phase 2 restores that snapshot and replays the rounds to compute intermediate trace values. The round computations are identical across phases. - Scratch is sized as `num_blocks * ROWS_PER_BLOCK * (8 + BLOCK_WORDS)` words, exactly matching what's written/read (no oversize). - Phase 2 launches `num_blocks * ROWS_PER_BLOCK` threads, so every thread maps to a valid block. - Padding rows are still handled by the existing `sha2_fill_first_dummy_row` + `sha2_fill_invalid_rows` kernels. - Digest row: Phase 1 snapshots the final state (after all rounds) for `row_in_block == ROUND_ROWS`; Phase 2 reads it to compute `final_hash = prev_hash + work_vars`. ## Test plan - [ ] `cargo nextest run --cargo-profile=fast -p openvm-sha2-circuit` - [ ] CUDA tracegen tests for both SHA-256 and SHA-512 variants (feature `cuda`) - [ ] Verify trace matrices match CPU reference for representative inputs resolves int-6985

Towards INT-7429. ## Summary - Mark `finalize_ptr`, `finalize`, and `native_keccak256` as `unsafe fn` in both `zkvm_impl.rs` and `host_impl.rs`. These functions write 32 bytes to an output pointer/slice via `copy_nonoverlapping` without runtime size checks (only `debug_assert!`), so callers must guarantee the buffer is at least 32 bytes. - Add `assert!` in the `tiny_keccak::Hasher` trait impl in `lib.rs` since the trait method must be safe and cannot propagate `unsafe` to callers. - Safe public APIs (`keccak256()`, `set_keccak256()`) are unchanged — they pass correctly-sized buffers internally. **TODO:** The `openvm-org/tiny-keccak` fork and `openvm-org/openvm-primitives` also call `Keccak256::finalize` / `native_keccak256` and will need matching updates (add `unsafe {}` blocks and/or `assert!` at their trait boundaries). --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-authored-by: claude[bot] <41898282+claude[bot]@users.noreply.github.com> Co-authored-by: Jonathan Wang <jonathanpwang@users.noreply.github.com>

Towards INT-7422.

Resolves INT-7439, INT-7422.

…ebra+ecc + update sha2 docs (#2710) Resolves INT-7439.

Towards INT-7429. ## Summary - Fix OOB read in `native_xorin` when `len` is not 4-byte aligned: `AlignedBuf::new` copied `adjusted_len` bytes from source buffers, reading 1-3 bytes past the valid region. Replace with `AlignedBuf::uninit` + `copy_nonoverlapping` of only `len` bytes. - Mark `native_xorin` and `native_keccakf` as `unsafe extern "C" fn` since they take raw pointers and require callers to guarantee pointer validity and buffer sizes. Add `unsafe` blocks in callers `xorin` and `keccakf` with safety comments. - Update keccak extension docs code example to reflect the new unsafe signatures. --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Towards INT-7478. - Replace `assert_eq!` with `Err(StaticProgramError::InvalidInstruction)` in multiplication executor `pre_compute_impl` for both bigint (mult.rs) and rv32im (mul/execution.rs). Functions returning `Result` should use error returns instead of panics, consistent with other executors in the codebase. - Add `/// # Safety` doc comments to all 12 `unsafe extern "C" fn` in bigint guest externs (externs.rs), documenting pointer validity, buffer size, and alignment preconditions. Follows precedent set by PRs #2706 (keccak) and #2708 (sha2). - Document unused branch opcode variants. --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Issue found and initially corrected by [Daniel Heim](https://github.com/dghelm)

Resolves INT-7541. Issue found and initially corrected by [Daniel Heim](https://github.com/dghelm).

) Co-authored-by: Ayush Shukla <ayush@axiom.xyz>

This PR introduces the `ColumnsAir` trait and makes it required by all AIRs. Via this trait, AIRs can (optionally) provide a list of column names. At powdr, we use this to generate (somewhat) human-readable autoprecompiles for OpenVM ([example](https://github.com/powdr-labs/powdr/blob/main/openvm-riscv/tests/apc_snapshots/complex/rotate.txt)). Unfortunately, this touches all AIRs and therefore a lot of files, but the changes are small and now change to `stark-backend` are needed. This change represents most of our [diff to OpenVM](https://github.com/openvm-org/openvm/compare/v2.0.0-beta.2...powdr-labs:openvm:v2-powdr-beta.2-reorganized?expand=1). --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

The metered segmentation memory formula hardcodes `main_cell_secondary_weight = 1.5`, which the doc comment on `DEFAULT_MAIN_CELL_SECONDARY_WEIGHT` derives assuming `need_rot = true` for every AIR (`mat_eval_bytes = (h / 2^l_skip) × (2 × w) × sizeof(EF)`). For AIRs with `need_rot = false`, `fold_ple` allocates `1 × w`, not `2 × w`, so their secondary contribution is half. This change threads per-AIR `need_rot` from `pk.vk.params` into `SegmentationCtx` and splits the secondary term: AIRs with `need_rot = true` use weight `1.5`, AIRs with `need_rot = false` use `0.75`. Main and interaction terms unchanged. [benchmark comparison](https://github.com/axiom-crypto/openvm-eth/actions/runs/25931043565) --------- Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

This is a manual re-application of 5b580ad to the new `Sha2BlockHasherSubAir` since it is in a new file. A negative test is added that fails without the new added constraint. ref INT-6416

ghost · 2026-05-18T18:52:19Z

Claude finished @jpw-axiom's task in 6m 59s —— View job

Code Review for PR #2592

Scope: 105 commits, 100 files, +37,627/−40,110 lines. develop-v2.0.0-rc.1 → develop-v2.0.0-beta.

All changes reviewed against origin/develop-v2.0.0-beta. This review covers the full PR including ~24 commits not covered by prior automated reviews.

Area	Status	Summary
SHA2-air generalization	✅ Excellent	FIPS 180-4 compliant, clean trait abstraction, soundness fixes applied
SHA2 soundness constraints	✅ Critical fixes	`is_enabled` monotonic-off + last-row padding constraint
Memory safety fixes	✅ Critical fixes	AlignedBuf use-after-free, OOB read in native_xorin
ColsRef derive macro	⚠️ Has issues	Copy-paste bug, unsafe casts, code duplication
Memory system refactoring	✅ Sound	Volatile memory + access adapters removed cleanly
CUDA/GPU safety fixes	✅ Excellent	CUB aliasing, keccakf UB, bounds checks — all correct
CUDA optimizations	✅ Excellent	Keccakf stack 1656B→384B, SHA2 3-phase trace, stream migration
SDK proof fixes	✅ Correct	Merkle path depth==0, recursion_flag mismatch, deferral verification
Metering performance	✅ Correct	Per-AIR `need_rot`, interaction weight tuning
CI/docs/benchmarks	✅ Correct	sccache, SHA2 rename consistent, docs updated

New Commits Since Prior Reviews (24 commits)

🔴 Critical Fixes — All Verified Correct

0110850 — re-apply sha2 fix to new Sha2BlockHasher (#2788)
Adds when_last_row().assert_one(local_is_padding_row) to Sha2BlockHasherSubAir (crates/circuits/sha2-air/src/air.rs:229-230). Without this, a malicious prover could submit a truncated 2-row trace with only round rows, bypassing all digest/finalization constraints. The negative test at extensions/sha2/circuit/src/sha2_chips/tests.rs validates this by constructing a crafted trace that satisfies compression constraints but omits padding rows. Sound.

0c32c0a — sha2 AlignedBuf bug (#2707)
Classic use-after-free: temporary AlignedBuf::new(...) values were consumed in if expressions and dropped before the intrinsic call, leaving dangling pointers. Fix correctly binds them to named variables (aligned_state, aligned_input, aligned_output) in the outer scope (extensions/sha2/guest/src/lib.rs:46-68). Correct.

5d357e0 — resolve OOB read in native_xorin (#2709)
AlignedBuf::new(ptr, adjusted_len, ...) would copy adjusted_len bytes (rounded up to 4-byte alignment) from the source, reading 1-3 bytes past the valid len. Fix uses AlignedBuf::uninit + copy_nonoverlapping(ptr, buf, len), copying only valid bytes. Functions correctly marked unsafe. Correct.

3cc550b — SDK merkle path generation depth==0 (#2731)
When depth == 0, the merkle proof needs the sibling at the leaf level. The fix at crates/sdk/src/prover/deferral/merkle.rs:46-50 correctly starts at leaf_idx + 1 (right sibling) when depth == 0 instead of leaf_idx >> 0 = leaf_idx. The debug_assert_eq!(leaf_idx % 2, 0) confirms the leaf is always a left child. Correct.

d9b189f — resolve recursion_flag mismatch in prove_mixed (#2720)
When VM and deferral proofs go through different numbers of internal recursion layers, prove_mixed would fail. Fix adds a check: if one is at layer 1 and the other is not, the lagging proof gets an extra wrap to equalize. Correct.

0d84b5e — verify-stark prover verifies deferral-enabled child proofs (#2717)
Changed prove_no_def → prove with vm_proof.deferral_merkle_proofs. Without this, child proofs with deferrals would fail verification. Correct.

3029c0e — Poseidon2 chip tracegen handles no records (#2729)
Fixes a bug where Poseidon2PeripheryBaseChip and DeferralPoseidon2Chip could produce non-empty traces when there were zero records. Adds explicit nonempty flag tracking with proper reset after generate_proving_ctx. Correct.

Performance Optimizations — Verified Sound

df460bc — per-AIR need_rot in segmentation memory formula (#2782)
Previously all AIRs assumed need_rot = true (weight 1.5). Now AIRs with need_rot = false get weight 0.75. The need_rot vector is threaded from pk.vk.params through MeteredCtx into SegmentationCtx with proper length assertions. Both calculate_cell_counts and should_segment updated consistently. Sound.

1d3c8dc — update interaction weight for virtual GKR padding (#2764)
Changes interaction cell weight from (2 * D_EF) * 1.04 to (2 * D_EF) * (1.0 + 2.0 * (1/16 + 1/256)) for more accurate scratch overhead modeling. Removes .next_power_of_two() from interaction_cnt since CUDA GKR virtualizes padding. Correct.

15b1b15 — optimize sha2 tracegen cuda kernel (#2695)
Splits monolithic 1-thread-per-block into three phases: (1) sequential round computation with scratch snapshots, (2) parallel per-row trace column writes, (3) cross-row dependency fixup. Phase separation improves memory coalescing. Sound.

db3bfe4 — stream migration refactoring (#2662)
Large (354 files) but systematic refactor adding explicit cudaStream_t to all CUDA launchers, removing PTDS. Lifts DeviceContext to first-class associated type. Enables proper multi-stream CUDA execution. Infrastructure change, consistent.

Other Noteworthy Commits

021a2fb — Expose AIR column names (#2771): Introduces ColumnsAir trait with default no-op across ~152 files. Pure additive, no behavioral changes. ✅

59c8412 — bigint + rv32im cleanup (#2713): Replaces assert_eq! panics with Err(StaticProgramError::InvalidInstruction). Adds safety docs to unsafe extern "C" functions. ✅

b67acb2, b24c73b — mark keccak256/sha2/algebra functions as unsafe (#2706, #2708): Correct unsafe markings for functions taking raw pointers. Safe wrappers add runtime bounds checks. ✅

44ceb46 — sha2 guest library api mismatch (#2703): Fixes API inconsistency. ✅

4ceb82f — clean up persistent boundary (#2705): Removes TouchedLabels enum, dead code removal. ✅

Outstanding Issues (Carried Forward + New)

🔴 High Priority

1. examples/u256/ still references removed openvm-ruint

Commit 947f874 removed the ruint guest library from the workspace, but examples/u256/ was not updated:

examples/u256/Cargo.toml:13 still has openvm-ruint = { git = "...", branch = "v1.14.0-openvm", package = "ruint" }
examples/u256/src/main.rs:7 still has use openvm_ruint::aliases::U256
docs/vocs/docs/pages/book/guest-libraries/ruint.mdx still documents ruint as active

Impact: The example is not in the workspace members, so it won't break workspace builds, but it will fail if built standalone. The docs page is stale.

Fix this →

🟡 Medium Priority

2. ColsRef derive macro — copy-paste bug in panic message

crates/circuits/primitives/derive/src/cols_ref/mod.rs:449-451: Inside get_const_cols_ref_fields, the panic message incorrectly says get_mut_cols_ref_fields:

panic!(
    "Unsupported field type (in get_mut_cols_ref_fields): {:?}",
    //                        ^^^ Should be "get_const_cols_ref_fields"
    f.ty
);

Line 389 in the same function correctly says "get_const_cols_ref_fields", confirming this is a copy-paste error.

Fix this →

3. Latent UB in rotr_generic<uint64_t> and rotr (SHA-2 CUDA)

extensions/sha2/circuit/cuda/include/block_hasher/variant.cuh:242-244:

template <> inline __device__ __host__ uint64_t rotr_generic<uint64_t>(uint64_t value, int n) {
    return (value >> n) | (value << (64 - n));  // UB when n=0
}

Same issue in crates/circuits/primitives/cuda/include/primitives/utils.cuh:40-42 (rotr for uint32_t). The safe (-n) & 63 pattern is already used in rotl64 at line 37 of utils.cuh. All current callers use non-zero FIPS rotation constants, so this is latent UB, not triggered today.

Fix this →

4. Unsafe pointer casts without alignment verification in ColsRef macro

crates/circuits/primitives/derive/src/cols_ref/mod.rs:341 (const) and :506 (mut): Generated code casts &[T] to &[#elem_type] via raw pointer without compile-time alignment assertions. Safety relies on implicit invariants about #[repr(C)] structs. Consider adding debug_assert! for alignment in the generated code.

5. Hardcoded DEFERRAL_AS = 4 in CUDA files

crates/vm/cuda/src/system/boundary.cu:9 and inventory.cu:23 both hardcode DEFERRAL_AS = 4 with TODO comments. If the Rust constant at crates/toolchain/instructions/src/lib.rs changes, CUDA code silently diverges.

🟢 Low Priority

6. ColsRef macro code duplication (~300 lines)

get_const_cols_ref_fields (lines 292-453) and get_mut_cols_ref_fields (lines 457-619) are ~90% identical. The copy-paste bug in issue #2 above is direct evidence this duplication causes errors.

7. ColsRef example test has no assertions

crates/circuits/primitives/derive/tests/example.rs:27: println! with no assertions. test_cols_ref.rs has good coverage — only example.rs is lacking.

8. is_columns_struct heuristic is fragile

crates/circuits/primitives/derive/src/cols_ref/mod.rs:630: Uses ends_with("Cols") for type detection. An attribute-based mechanism would be more robust.

Areas Verified Sound

SHA2-air crate: All FIPS 180-4 constants correct for SHA-256/384/512. Constraint degree ≤3. Last-row padding constraint enforced. is_enabled boolean + contiguous padding constraints on main chip. SOUNDNESS.md thorough.
Memory system: Direction constraint d = d³ correct. Initial row selection when(d*(d+1)) correct. Volatile memory + access adapters fully removed. CHUNK=8, DEFAULT_BLOCK_SIZE=4, BLOCKS_PER_CHUNK=2 verified.
constant_trace_height trait: Clean default None. All lookup chips + deferral count chips implement consistently (CPU + GPU).
MeteredCtx: Correctly separates constant vs variable trace heights. Per-AIR need_rot properly threaded through segmentation.
CUDA: CUB aliasing fixed in all 3 Poseidon2 dedup sites. rotl64 uses safe (-n)&63. Keccakf round body deduplicated. Stream migration consistent across 354 files. SHA2 3-phase trace optimization sound.
SDK proofs: Merkle path depth==0 fixed. Recursion flag equalization correct. Deferral-enabled verification restored.
Guest safety: AlignedBuf lifetime bug fixed. OOB read in xorin fixed. All raw-pointer functions marked unsafe with safe wrappers.

Summary

This is a well-executed release candidate PR with significant architectural improvements. The newest commits include several critical safety and soundness fixes that are all correctly implemented:

SHA2 when_last_row() padding constraint — prevents trace truncation attack
AlignedBuf use-after-free in SHA2 guest — classic dangling pointer
OOB read in native_xorin — reads past valid buffer length
Merkle path depth==0 off-by-one — wrong sibling in proof
Recursion flag mismatch — proof generation failure when aggregation depths diverge
Poseidon2 zero-record handling — non-empty trace from empty chip

The CUDA stream migration (~354 files) is a large infrastructure change that enables multi-stream execution, consistently applied across all extensions.

Recommended actions before merge (priority order):

Fix or remove examples/u256/ + stale ruint docs (broken after ruint removal)
Fix ColsRef panic message copy-paste bug (one-line fix)
Harden rotr_generic / rotr with safe rotation pattern (latent UB)

Gather context on latest changes since prior reviews
Review new commits (24 commits since prior reviews)
Review SHA2-air crate and soundness
Review ColsRef derive macro
Review CUDA/GPU safety fixes and optimizations
Review memory system / VM architecture changes
Review CI/benchmark/docs changes
Check outstanding issues from prior reviews
Post final review

ghost · 2026-05-18T19:03:52Z

group	app.proof_time_ms	app.cycles	leaf.proof_time_ms
fibonacci	3,770	12,000,265	925
keccak	18,360	18,655,329	3,240
sha2_bench	10,359	14,793,960	1,485
regex	1,405	4,137,067	357
ecrecover	608	123,583	243
pairing	884	1,745,757	269
kitchen_sink	1,890	2,579,903	410
fibonacci_e2e	1,771	12,000,265	409
regex_e2e	826	4,137,067	169
ecrecover_e2e	512	123,583	130
pairing_e2e	631	1,745,757	132
kitchen_sink_e2e	2,034	2,579,903	402

Note: cells_used metrics omitted because CUDA tracegen does not expose unpadded trace heights.

Commit: 0110850

Benchmark Workflow

…2785) Adds `release: types: [released]` so a non-prerelease, non-draft release publish triggers both workflows, using the release tag as the checkout ref and version. `workflow_dispatch` remains supported for manual runs. Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

ghost · 2026-05-21T00:07:26Z

group	app.proof_time_ms	app.cycles	leaf.proof_time_ms
fibonacci	3,785	12,000,265	925
keccak	18,734	18,655,329	3,302
sha2_bench	10,338	14,793,960	1,485
regex	1,393	4,137,067	352
ecrecover	606	123,583	246
pairing	898	1,745,757	262
kitchen_sink	1,896	2,579,903	413
fibonacci_e2e	1,772	12,000,265	407
regex_e2e	828	4,137,067	169
ecrecover_e2e	515	123,583	132
pairing_e2e	632	1,745,757	132
kitchen_sink_e2e	2,033	2,579,903	397

Note: cells_used metrics omitted because CUDA tracegen does not expose unpadded trace heights.

Commit: 5ddd539

Benchmark Workflow

Need these changes so axiom-crypto/openvm-eth#609 becomes possible Resolves INT-7744 --------- Co-authored-by: Allan Lin <allanl@intrinsictech.xyz>

ghost · 2026-05-21T18:37:53Z

group	app.proof_time_ms	app.cycles	leaf.proof_time_ms
fibonacci	3,776	12,000,265	915
keccak	18,607	18,655,329	3,275
sha2_bench	10,423	14,793,960	1,481
regex	1,411	4,137,067	360
ecrecover	606	123,583	256
pairing	881	1,745,757	259
kitchen_sink	1,896	2,579,903	409
fibonacci_e2e	1,636	12,000,265	411
regex_e2e	679	4,137,067	170
ecrecover_e2e	361	123,583	131
pairing_e2e	483	1,745,757	131
kitchen_sink_e2e	1,875	2,579,903	406

Note: cells_used metrics omitted because CUDA tracegen does not expose unpadded trace heights.

Commit: 5ad6253

Benchmark Workflow

Co-authored-by: Allan Lin <allanl@intrinsictech.xyz>

socket-security · 2026-05-22T14:40:11Z

Review the following changes in direct dependencies. Learn more about Socket for GitHub.

Diff	Package	Supply Chain Security	Vulnerability	Quality	Maintenance	License
	cargo/ndarray@0.16.1
	cargo/struct-reflection@0.1.0

View full report

ghost · 2026-05-22T14:51:08Z

group	app.proof_time_ms	app.cycles	leaf.proof_time_ms
fibonacci	3,750	12,000,265	913
keccak	18,516	18,655,329	3,278
sha2_bench	10,204	14,793,960	1,472
regex	1,396	4,137,067	353
ecrecover	601	123,583	249
pairing	896	1,745,757	263
kitchen_sink	1,884	2,579,903	409
fibonacci_e2e	1,634	12,000,265	413
regex_e2e	676	4,137,067	169
ecrecover_e2e	360	123,583	129
pairing_e2e	489	1,745,757	133
kitchen_sink_e2e	1,867	2,579,903	399

Note: cells_used metrics omitted because CUDA tracegen does not expose unpadded trace heights.

Commit: b8194ef

Benchmark Workflow

shuklaayush changed the title ~~Develop v2.0.0 rc.1~~ test(DON'T MERGE): develop-v2.0.0-rc.1 Mar 19, 2026

shuklaayush force-pushed the develop-v2.0.0-rc.1 branch from e9d2afb to deff2f2 Compare March 19, 2026 18:20

branch-rebase-bot Bot force-pushed the develop-v2.0.0-rc.1 branch 2 times, most recently from a6c7835 to 8d51084 Compare March 19, 2026 21:05

shuklaayush force-pushed the develop-v2.0.0-rc.1 branch from 8d51084 to 743b94e Compare March 19, 2026 22:48