feat: split AddSub into AddSub and AddI chips

crates/rvr/rvr-openvm-lift/src/opcode.rs

@@ -12,9 +12,10 @@ use openvm_instructions::{
     instruction::Instruction, riscv::RV64_REGISTER_NUM_LIMBS, LocalOpcode, SysPhantom, SystemOpcode,
 };
 use openvm_riscv_transpiler::{
-    BaseAluOpcode, BaseAluWOpcode, BranchEqualOpcode, BranchLessThanOpcode, DivRemOpcode,
-    DivRemWOpcode, LessThanOpcode, MulHOpcode, MulOpcode, MulWOpcode, Rv64AuipcOpcode,
-    Rv64JalLuiOpcode, Rv64JalrOpcode, Rv64LoadStoreOpcode, ShiftOpcode, ShiftWOpcode,
+    AddIOpcode, BaseAluOpcode, BaseAluWOpcode, BranchEqualOpcode, BranchLessThanOpcode,
+    DivRemOpcode, DivRemWOpcode, LessThanOpcode, MulHOpcode, MulOpcode, MulWOpcode,
+    Rv64AuipcOpcode, Rv64JalLuiOpcode, Rv64JalrOpcode, Rv64LoadStoreOpcode, ShiftOpcode,
+    ShiftWOpcode,
 };
 use openvm_stark_backend::p3_field::PrimeField32;
 use rvr_openvm_ext_ffi_common::{AS_MEMORY, AS_PUBLIC_VALUES, AS_REGISTER};
@@ -84,6 +85,11 @@ pub fn lift_instruction<F: PrimeField32>(
         return Some(lift_alu_w(insn, pc, e, AluOp::Sub));
     }
 
+    // AddI: ADDI — always immediate (e == 0 from the transpiler)
+    if opcode == AddIOpcode::ADDI.global_opcode_usize() {
+        return Some(lift_alu(insn, pc, e, AluOp::Add));
+    }
+
     // Shift: SLL=0x205, SRL=0x206, SRA=0x207, SLLW=0x275, SRLW=0x276, SRAW=0x277
     if opcode == ShiftOpcode::SLL.global_opcode_usize() {
         return Some(lift_shift(insn, pc, e, AluOp::Sll));

extensions/riscv/circuit/cuda/include/riscv/adapters/alu_u16_addi.cuh

@@ -0,0 +1,60 @@
+#pragma once
+
+#include "primitives/execution.h"
+#include "primitives/histogram.cuh"
+#include "primitives/trace_access.h"
+#include "primitives/constants.h"
+#include "system/memory/controller.cuh"
+#include "system/memory/offline_checker.cuh"
+
+using namespace riscv;
+
+// Adapter columns for ADDI — immediate-only variant of Rv64BaseAluU16AdapterCols.
+// rs2 is always an immediate, so there is no rs2_as / rs2_imm_sign and only one reads_aux.
+template <typename T> struct Rv64AddIAdapterCols {
+    ExecutionState<T> from_state;
+    T rd_ptr;
+    T rs1_ptr;
+    MemoryReadAuxCols<T> reads_aux;
+    MemoryWriteAuxCols<T, BLOCK_FE_WIDTH> writes_aux;
+};
+
+struct Rv64AddIAdapterRecord {
+    uint32_t from_pc;
+    uint32_t from_timestamp;
+    uint32_t rd_ptr;
+    uint32_t rs1_ptr;
+    MemoryReadAuxRecord reads_aux;
+    MemoryWriteU16AuxRecord<BLOCK_FE_WIDTH> writes_aux;
+};
+
+struct Rv64AddIAdapter {
+    MemoryAuxColsFactory mem_helper;
+
+    __device__ Rv64AddIAdapter(VariableRangeChecker rc, uint32_t timestamp_max_bits)
+        : mem_helper(rc, timestamp_max_bits) {}
+
+    __device__ void fill_trace_row(RowSlice row, Rv64AddIAdapterRecord record) {
+        COL_WRITE_VALUE(row, Rv64AddIAdapterCols, from_state.pc, record.from_pc);
+        COL_WRITE_VALUE(row, Rv64AddIAdapterCols, from_state.timestamp, record.from_timestamp);
+        COL_WRITE_VALUE(row, Rv64AddIAdapterCols, rd_ptr, record.rd_ptr);
+        COL_WRITE_VALUE(row, Rv64AddIAdapterCols, rs1_ptr, record.rs1_ptr);
+
+        // rs1 register read at timestamp slot 0.
+        mem_helper.fill(
+            row.slice_from(COL_INDEX(Rv64AddIAdapterCols, reads_aux)),
+            record.reads_aux.prev_timestamp,
+            record.from_timestamp
+        );
+
+        // rd write at timestamp slot 1.
+        Fp prev[BLOCK_FE_WIDTH];
+        copy_u16_cells(prev, record.writes_aux.prev_data);
+        COL_WRITE_ARRAY(row, Rv64AddIAdapterCols, writes_aux.prev_data, prev);
+        mem_helper.fill(
+            row.slice_from(COL_INDEX(Rv64AddIAdapterCols, writes_aux)),
+            record.writes_aux.prev_timestamp,
+            record.from_timestamp + 1
+        );
+    }
+};

extensions/riscv/circuit/cuda/include/riscv/cores/addi.cuh

@@ -0,0 +1,61 @@
+#pragma once
+
+#include "primitives/constants.h"
+#include "primitives/histogram.cuh"
+#include "primitives/trace_access.h"
+
+using namespace riscv;
+
+template <size_t NUM_LIMBS> struct AddICoreRecord {
+    uint16_t rs1[NUM_LIMBS];
+    uint16_t imm_low11;
+    uint16_t imm_sign;
+};
+
+template <typename T, size_t NUM_LIMBS> struct AddICoreCols {
+    T rd[NUM_LIMBS];
+    T rs1[NUM_LIMBS];
+    T imm_low11;
+    T imm_sign;
+    T is_valid;
+};
+
+template <size_t NUM_LIMBS, size_t LIMB_BITS> struct AddICore {
+    VariableRangeChecker range_checker;
+
+    template <typename T> using Cols = AddICoreCols<T, NUM_LIMBS>;
+
+    __device__ AddICore(VariableRangeChecker rc) : range_checker(rc) {}
+
+    __device__ void fill_trace_row(RowSlice row, AddICoreRecord<NUM_LIMBS> record) {
+        uint16_t rd[NUM_LIMBS];
+
+        // First limb: rs1[0] + imm_low11 + imm_sign * 0xF800
+        uint32_t overflow = static_cast<uint32_t>(record.rs1[0])
+            + static_cast<uint32_t>(record.imm_low11)
+            + static_cast<uint32_t>(record.imm_sign) * 0xF800u;
+        uint32_t carry = overflow >> LIMB_BITS;
+        rd[0] = static_cast<uint16_t>(overflow & ((1u << LIMB_BITS) - 1));
+
+        // Remaining limbs: rs1[i] + sign_u16 + carry
+        uint32_t sign_u16 = static_cast<uint32_t>(record.imm_sign) * 0xFFFFu;
+#pragma unroll
+        for (size_t i = 1; i < NUM_LIMBS; i++) {
+            overflow = static_cast<uint32_t>(record.rs1[i]) + sign_u16 + carry;
+            carry = overflow >> LIMB_BITS;
+            rd[i] = static_cast<uint16_t>(overflow & ((1u << LIMB_BITS) - 1));
+        }
+
+        COL_WRITE_ARRAY(row, Cols, rs1, record.rs1);
+        COL_WRITE_ARRAY(row, Cols, rd, rd);
+        COL_WRITE_VALUE(row, Cols, imm_low11, record.imm_low11);
+        COL_WRITE_VALUE(row, Cols, imm_sign, record.imm_sign);
+        COL_WRITE_VALUE(row, Cols, is_valid, 1u);
+
+        range_checker.add_count(record.imm_low11, 11);
+#pragma unroll
+        for (size_t i = 0; i < NUM_LIMBS; i++) {
+            range_checker.add_count(rd[i], LIMB_BITS);
+        }
+    }
+};

extensions/riscv/circuit/cuda/src/addi.cu

@@ -0,0 +1,70 @@
+#include "launcher.cuh"
+#include "primitives/buffer_view.cuh"
+#include "primitives/constants.h"
+#include "primitives/histogram.cuh"
+#include "primitives/trace_access.h"
+#include "riscv/adapters/alu_u16_addi.cuh"
+#include "riscv/cores/addi.cuh"
+#include "system/memory/params.cuh"
+
+using namespace riscv;
+
+// Concrete type aliases for RV64
+using Rv64AddICoreRecord = AddICoreRecord<BLOCK_FE_WIDTH>;
+using Rv64AddICore = AddICore<BLOCK_FE_WIDTH, U16_BITS>;
+template <typename T> using Rv64AddICoreCols = AddICoreCols<T, BLOCK_FE_WIDTH>;
+
+template <typename T> struct Rv64AddICols {
+    Rv64AddIAdapterCols<T> adapter;
+    Rv64AddICoreCols<T> core;
+};
+
+struct Rv64AddIRecord {
+    Rv64AddIAdapterRecord adapter;
+    Rv64AddICoreRecord core;
+};
+
+__global__ void addi_tracegen(
+    Fp *d_trace,
+    size_t height,
+    DeviceBufferConstView<Rv64AddIRecord> d_records,
+    uint32_t *d_range_checker_ptr,
+    uint32_t range_checker_num_bins,
+    uint32_t timestamp_max_bits
+) {
+    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    RowSlice row(d_trace + idx, height);
+    if (idx < d_records.len()) {
+        auto const &rec = d_records[idx];
+
+        auto adapter = Rv64AddIAdapter(
+            VariableRangeChecker(d_range_checker_ptr, range_checker_num_bins),
+            timestamp_max_bits
+        );
+        adapter.fill_trace_row(row, rec.adapter);
+
+        auto core =
+            Rv64AddICore(VariableRangeChecker(d_range_checker_ptr, range_checker_num_bins));
+        core.fill_trace_row(row.slice_from(COL_INDEX(Rv64AddICols, core)), rec.core);
+    } else {
+        row.fill_zero(0, sizeof(Rv64AddICols<uint8_t>));
+    }
+}
+
+extern "C" int _addi_tracegen(
+    Fp *d_trace,
+    size_t height,
+    size_t width,
+    DeviceBufferConstView<Rv64AddIRecord> d_records,
+    uint32_t *d_range_checker,
+    uint32_t range_checker_num_bins,
+    uint32_t timestamp_max_bits,
+    cudaStream_t stream
+) {
+    assert(width == sizeof(Rv64AddICols<uint8_t>));
+    auto [grid, block] = kernel_launch_params(height);
+    addi_tracegen<<<grid, block, 0, stream>>>(
+        d_trace, height, d_records, d_range_checker, range_checker_num_bins, timestamp_max_bits
+    );
+    return CHECK_KERNEL();
+}