From 3a104c716bc8a37e1789b508590d5f83146bbcc8 Mon Sep 17 00:00:00 2001 From: kunxian xia Date: Tue, 10 Feb 2026 11:55:02 +0800 Subject: [PATCH 1/2] clippy --- .../native/circuit/src/sumcheck/chip.rs | 7 +++--- .../native/circuit/src/sumcheck/execution.rs | 2 +- extensions/native/compiler/src/ir/sumcheck.rs | 1 + extensions/native/recursion/tests/sumcheck.rs | 25 ++++++++----------- 4 files changed, 15 insertions(+), 20 deletions(-) diff --git a/extensions/native/circuit/src/sumcheck/chip.rs b/extensions/native/circuit/src/sumcheck/chip.rs index 2ef609bba8..d4fbf2524d 100644 --- a/extensions/native/circuit/src/sumcheck/chip.rs +++ b/extensions/native/circuit/src/sumcheck/chip.rs @@ -10,9 +10,7 @@ use openvm_circuit::{ native_adapter::util::{memory_read_native, tracing_write_native_inplace}, }, }; -use openvm_instructions::{ - instruction::Instruction, program::DEFAULT_PC_STEP, LocalOpcode, NATIVE_AS, -}; +use openvm_instructions::{instruction::Instruction, program::DEFAULT_PC_STEP, LocalOpcode}; use openvm_native_compiler::SumcheckOpcode::SUMCHECK_LAYER_EVAL; use openvm_stark_backend::p3_field::PrimeField32; @@ -227,7 +225,8 @@ where let mut eval_acc = elem_to_ext(F::from_canonical_u32(0)); let mut alpha_acc = elem_to_ext(F::from_canonical_u32(1)); - // all rows share same register values, ctx, challenges, max_round, hint_space_ptrs (optional) + // all rows share same register values, ctx, challenges, max_round, hint_space_ptrs + // (optional) for row in rows.iter_mut() { // c1, c2 are same during the entire execution row.challenges[EXT_DEG..3 * EXT_DEG].copy_from_slice(&challenges[EXT_DEG..3 * EXT_DEG]); diff --git a/extensions/native/circuit/src/sumcheck/execution.rs b/extensions/native/circuit/src/sumcheck/execution.rs index 484b306392..92bd3ed1aa 100644 --- a/extensions/native/circuit/src/sumcheck/execution.rs +++ b/extensions/native/circuit/src/sumcheck/execution.rs @@ -33,7 +33,7 @@ impl NativeSumcheckExecutor { #[inline(always)] fn pre_compute_impl( &self, - pc: u32, + _pc: u32, inst: &Instruction, data: &mut NativeSumcheckPreCompute, ) -> Result<(), StaticProgramError> { diff --git a/extensions/native/compiler/src/ir/sumcheck.rs b/extensions/native/compiler/src/ir/sumcheck.rs index 15e8fe0031..5c15eee50f 100644 --- a/extensions/native/compiler/src/ir/sumcheck.rs +++ b/extensions/native/compiler/src/ir/sumcheck.rs @@ -26,6 +26,7 @@ impl Builder { /// /// 2. for computing expected eval of next layer, output[1+i] = eq(0,r)*p[i][0] + eq(1,r) * /// p[i][1]. + #[allow(clippy::too_many_arguments)] pub fn sumcheck_layer_eval( &mut self, input_ctx: &Array>, // Context variables diff --git a/extensions/native/recursion/tests/sumcheck.rs b/extensions/native/recursion/tests/sumcheck.rs index 1e3ec31166..df6194468f 100644 --- a/extensions/native/recursion/tests/sumcheck.rs +++ b/extensions/native/recursion/tests/sumcheck.rs @@ -1,4 +1,4 @@ -use std::iter::{once, repeat_n}; +use std::iter::once; use openvm_circuit::{arch::instructions::program::Program, utils::air_test_impl}; #[cfg(feature = "cuda")] @@ -35,12 +35,10 @@ fn test_sumcheck_layer_eval_with_hint_ids() { let num_logup_specs = 8; let prod_evals: Vec = (0..(num_prod_specs * num_layers * 2)) - .into_iter() .map(|_| new_rand_ext(&mut rng)) .collect(); let logup_evals: Vec = (0..(num_logup_specs * num_layers * 4)) - .into_iter() .map(|_| new_rand_ext(&mut rng)) .collect(); @@ -73,13 +71,10 @@ fn test_sumcheck_layer_eval_with_hint_ids() { standard_fri_params_with_100_bits_conjectured_security(1) }; - let mut input_stream: Vec> = vec![]; - input_stream.push( - prod_evals - .into_iter() - .flat_map(|e| >::as_base_slice(&e).to_vec()) - .collect(), - ); + let mut input_stream: Vec> = vec![prod_evals + .into_iter() + .flat_map(|e| >::as_base_slice(&e).to_vec()) + .collect()]; input_stream.push( logup_evals .into_iter() @@ -137,7 +132,7 @@ fn build_test_program( ) { let mode = 1; // current_layer - let mut ctx_u32s = vec![ + let ctx_u32s = vec![ round, num_prod_specs, num_logup_specs, @@ -175,16 +170,16 @@ fn build_test_program( let num_prod_evals = num_prod_specs * num_layers * 2; let prod_spec_evals: Array> = builder.dyn_array(num_prod_evals); - for idx in 0..num_prod_evals { - let e: Ext = builder.constant(prod_evals[idx]); + for (idx, prod_eval) in prod_evals.into_iter().enumerate() { + let e: Ext = builder.constant(prod_eval); builder.set(&prod_spec_evals, idx, e); } let num_logup_evals = num_logup_specs * num_layers * 4; let logup_spec_evals: Array> = builder.dyn_array(num_logup_evals); - for idx in 0..num_logup_evals { - let e: Ext = builder.constant(logup_evals[idx]); + for (idx, logup_eval) in logup_evals.into_iter().enumerate() { + let e: Ext = builder.constant(logup_eval); builder.set(&logup_spec_evals, idx, e); } From 5d385444b94c2bcb78ee4c62cb3cf4c43555d4a6 Mon Sep 17 00:00:00 2001 From: kunxian xia Date: Tue, 10 Feb 2026 12:00:06 +0800 Subject: [PATCH 2/2] fix gpu bug --- .../native/circuit/cuda/include/native/sumcheck.cuh | 4 ++-- extensions/native/circuit/cuda/src/sumcheck.cu | 8 +++++++- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/extensions/native/circuit/cuda/include/native/sumcheck.cuh b/extensions/native/circuit/cuda/include/native/sumcheck.cuh index a67b134aaf..a7d6eee536 100644 --- a/extensions/native/circuit/cuda/include/native/sumcheck.cuh +++ b/extensions/native/circuit/cuda/include/native/sumcheck.cuh @@ -17,18 +17,18 @@ template struct HeaderSpecificCols { template struct ProdSpecificCols { T data_ptr; T p[EXT_DEG * 2]; - MemoryReadAuxCols read_records[1]; T p_evals[EXT_DEG]; MemoryWriteAuxCols write_record; + MemoryWriteAuxCols ps_record; T eval_rlc[EXT_DEG]; }; template struct LogupSpecificCols { T data_ptr; T pq[EXT_DEG * 4]; - MemoryReadAuxCols read_records[1]; T p_evals[EXT_DEG]; T q_evals[EXT_DEG]; + MemoryWriteAuxCols pqs_record; MemoryWriteAuxCols write_records[2]; T eval_rlc[EXT_DEG]; }; diff --git a/extensions/native/circuit/cuda/src/sumcheck.cu b/extensions/native/circuit/cuda/src/sumcheck.cu index f89b93f72d..139a56473f 100644 --- a/extensions/native/circuit/cuda/src/sumcheck.cu +++ b/extensions/native/circuit/cuda/src/sumcheck.cu @@ -6,12 +6,18 @@ using namespace native; +constexpr uint32_t header_read_records_len() { + return sizeof(((HeaderSpecificCols *)nullptr)->read_records) + / sizeof(MemoryReadAuxCols); +} + __device__ void fill_sumcheck_specific(RowSlice row, MemoryAuxColsFactory &mem_helper) { RowSlice specific = row.slice_from(COL_INDEX(NativeSumcheckCols, specific)); uint32_t start_timestamp = row[COL_INDEX(NativeSumcheckCols, start_timestamp)].asUInt32(); if (row[COL_INDEX(NativeSumcheckCols, header_row)] == Fp::one()) { - for (uint32_t i = 0; i < 8; ++i) { + constexpr uint32_t header_records = header_read_records_len(); + for (uint32_t i = 0; i < header_records; ++i) { mem_fill_base( mem_helper, start_timestamp + i,