Syscall for Fp2AddSubSyscall {
local_mem_access: rt.postprocess(),
};
match P::FIELD_TYPE {
- // All the fp2 add and sub events for a given curve are coalesced to the curve's fp2 add operation. Only check for
- // that operation.
+ // All the fp2 add and sub events for a given curve are coalesced to the curve's fp2 add
+ // operation. Only check for that operation.
// TODO: Fix this.
FieldType::Bn254 => {
let syscall_code_key = match syscall_code {
diff --git a/crates/core/executor/src/syscalls/precompiles/mod.rs b/crates/core/executor/src/syscalls/precompiles/mod.rs
index f07da94609..4b06dd3c12 100644
--- a/crates/core/executor/src/syscalls/precompiles/mod.rs
+++ b/crates/core/executor/src/syscalls/precompiles/mod.rs
@@ -2,5 +2,6 @@ pub mod edwards;
pub mod fptower;
pub mod keccak256;
pub mod sha256;
+pub mod u256x2048_mul;
pub mod uint256;
pub mod weierstrass;
diff --git a/crates/core/executor/src/syscalls/precompiles/u256x2048_mul.rs b/crates/core/executor/src/syscalls/precompiles/u256x2048_mul.rs
new file mode 100644
index 0000000000..3795055608
--- /dev/null
+++ b/crates/core/executor/src/syscalls/precompiles/u256x2048_mul.rs
@@ -0,0 +1,91 @@
+use num::{BigUint, Integer, One};
+
+use sp1_primitives::consts::{bytes_to_words_le, words_to_bytes_le_vec};
+
+use crate::{
+ events::{PrecompileEvent, U256xU2048MulEvent},
+ syscalls::{Syscall, SyscallCode, SyscallContext},
+ Register::{X12, X13},
+};
+
+const U256_NUM_WORDS: usize = 8;
+const U2048_NUM_WORDS: usize = 64;
+const U256_NUM_BYTES: usize = U256_NUM_WORDS * 4;
+const U2048_NUM_BYTES: usize = U2048_NUM_WORDS * 4;
+
+pub(crate) struct U256xU2048MulSyscall;
+
+impl Syscall for U256xU2048MulSyscall {
+ fn execute(
+ &self,
+ rt: &mut SyscallContext,
+ syscall_code: SyscallCode,
+ arg1: u32,
+ arg2: u32,
+ ) -> Option {
+ let clk = rt.clk;
+
+ let a_ptr = arg1;
+ let b_ptr = arg2;
+
+ let (lo_ptr_memory, lo_ptr) = rt.mr(X12 as u32);
+ let (hi_ptr_memory, hi_ptr) = rt.mr(X13 as u32);
+
+ let (a_memory_records, a) = rt.mr_slice(a_ptr, U256_NUM_WORDS);
+ let (b_memory_records, b) = rt.mr_slice(b_ptr, U2048_NUM_WORDS);
+ let uint256_a = BigUint::from_bytes_le(&words_to_bytes_le_vec(&a));
+ let uint2048_b = BigUint::from_bytes_le(&words_to_bytes_le_vec(&b));
+
+ let result = uint256_a * uint2048_b;
+
+ let two_to_2048 = BigUint::one() << 2048;
+
+ let (hi, lo) = result.div_rem(&two_to_2048);
+
+ let mut lo_bytes = lo.to_bytes_le();
+ lo_bytes.resize(U2048_NUM_BYTES, 0u8);
+ let lo_words = bytes_to_words_le::(&lo_bytes);
+
+ let mut hi_bytes = hi.to_bytes_le();
+ hi_bytes.resize(U256_NUM_BYTES, 0u8);
+ let hi_words = bytes_to_words_le::(&hi_bytes);
+
+ // Increment clk so that the write is not at the same cycle as the read.
+ rt.clk += 1;
+
+ let lo_memory_records = rt.mw_slice(lo_ptr, &lo_words);
+ let hi_memory_records = rt.mw_slice(hi_ptr, &hi_words);
+ let lookup_id = rt.syscall_lookup_id;
+ let shard = rt.current_shard();
+ let event = PrecompileEvent::U256xU2048Mul(U256xU2048MulEvent {
+ lookup_id,
+ shard,
+ clk,
+ a_ptr,
+ a,
+ b_ptr,
+ b,
+ lo_ptr,
+ lo: lo_words.to_vec(),
+ hi_ptr,
+ hi: hi_words.to_vec(),
+ lo_ptr_memory,
+ hi_ptr_memory,
+ a_memory_records,
+ b_memory_records,
+ lo_memory_records,
+ hi_memory_records,
+ local_mem_access: rt.postprocess(),
+ });
+
+ let sycall_event =
+ rt.rt.syscall_event(clk, syscall_code.syscall_id(), arg1, arg2, lookup_id);
+ rt.add_precompile_event(syscall_code, sycall_event, event);
+
+ None
+ }
+
+ fn num_extra_cycles(&self) -> u32 {
+ 1
+ }
+}
diff --git a/crates/core/executor/src/syscalls/verify.rs b/crates/core/executor/src/syscalls/verify.rs
index 0197199e51..b17d795d3d 100644
--- a/crates/core/executor/src/syscalls/verify.rs
+++ b/crates/core/executor/src/syscalls/verify.rs
@@ -1,3 +1,5 @@
+use crate::DeferredProofVerification;
+
use super::{Syscall, SyscallCode, SyscallContext};
pub(crate) struct VerifySyscall;
@@ -32,16 +34,22 @@ impl Syscall for VerifySyscall {
let vkey_bytes: [u32; 8] = vkey.try_into().unwrap();
let pv_digest_bytes: [u32; 8] = pv_digest.try_into().unwrap();
- ctx.rt
- .subproof_verifier
- .verify_deferred_proof(proof, proof_vk, vkey_bytes, pv_digest_bytes)
- .unwrap_or_else(|e| {
- panic!(
- "Failed to verify proof {proof_index} with digest {}: {}",
- hex::encode(bytemuck::cast_slice(&pv_digest_bytes)),
- e
- )
- });
+ // Skip deferred proof verification if the corresponding runtime flag is set.
+ match ctx.rt.deferred_proof_verification {
+ DeferredProofVerification::Enabled => {
+ ctx.rt
+ .subproof_verifier
+ .verify_deferred_proof(proof, proof_vk, vkey_bytes, pv_digest_bytes)
+ .unwrap_or_else(|e| {
+ panic!(
+ "Failed to verify proof {proof_index} with digest {}: {}",
+ hex::encode(bytemuck::cast_slice(&pv_digest_bytes)),
+ e
+ )
+ });
+ }
+ DeferredProofVerification::Disabled => {}
+ }
None
}
diff --git a/crates/core/machine/Cargo.toml b/crates/core/machine/Cargo.toml
index 3bbd29d0b9..ff60ab4838 100644
--- a/crates/core/machine/Cargo.toml
+++ b/crates/core/machine/Cargo.toml
@@ -8,11 +8,12 @@ license = { workspace = true }
repository = { workspace = true }
keywords = { workspace = true }
categories = { workspace = true }
+links = "sp1-core-machine-sys"
[dependencies]
bincode = "1.3.3"
-serde = { version = "1.0", features = ["derive", "rc"] }
-itertools = "0.13.0"
+serde = { workspace = true, features = ["derive", "rc"] }
+itertools = { workspace = true }
log = "0.4.22"
num = { version = "0.4.3" }
p3-air = { workspace = true }
@@ -27,6 +28,9 @@ p3-util = { workspace = true }
sp1-derive = { workspace = true }
sp1-primitives = { workspace = true }
+rayon = "1.10.0"
+rayon-scan = "0.1.1"
+
amcl = { package = "snowbridge-amcl", version = "1.0.2", default-features = false, features = [
"bls381",
] }
@@ -41,15 +45,15 @@ p256 = { version = "0.13.2", features = ["expose-field"] }
num_cpus = "1.16.0"
size = "0.4.1"
tempfile = "3.10.1"
-tracing = "0.1.40"
+tracing = { workspace = true }
tracing-forest = { version = "0.1.6", features = ["ansi", "smallvec"] }
-tracing-subscriber = { version = "0.3.18", features = ["std", "env-filter"] }
+tracing-subscriber = { workspace = true, features = ["std", "env-filter"] }
strum_macros = "0.26"
strum = "0.26"
web-time = "1.1.0"
thiserror = "1.0.63"
rand = "0.8.5"
-hashbrown = { version = "0.14.5", features = ["serde", "inline-more"] }
+hashbrown = { workspace = true, features = ["serde", "inline-more"] }
static_assertions = "1.1.0"
sp1-stark = { workspace = true }
@@ -63,11 +67,23 @@ num = { version = "0.4.3", features = ["rand"] }
rand = "0.8.5"
sp1-zkvm = { workspace = true }
sp1-core-executor = { workspace = true, features = ["programs"] }
+test-artifacts = { workspace = true }
+
+[build-dependencies]
+sp1-stark = { workspace = true }
+sp1-primitives = { workspace = true }
+p3-baby-bear = { workspace = true }
+cbindgen = "0.27.0"
+cc = "1.1"
+pathdiff = "0.2.1"
+glob = "0.3.1"
[features]
+default = []
programs = []
debug = []
bigint-rug = ["sp1-curves/bigint-rug"]
+sys = []
[lib]
bench = false
diff --git a/crates/core/machine/build.rs b/crates/core/machine/build.rs
new file mode 100644
index 0000000000..f2088bd6ac
--- /dev/null
+++ b/crates/core/machine/build.rs
@@ -0,0 +1,169 @@
+fn main() {
+ #[cfg(feature = "sys")]
+ sys::build_ffi();
+}
+
+#[cfg(feature = "sys")]
+mod sys {
+ use std::{
+ env, fs, os,
+ path::{Path, PathBuf},
+ };
+
+ use pathdiff::diff_paths;
+
+ /// The library name, used for the static library archive and the headers.
+ /// Should be chosen as to not conflict with other library/header names.
+ const LIB_NAME: &str = "sp1-core-machine-sys";
+
+ /// The name of all include directories involved, used to find and output header files.
+ const INCLUDE_DIRNAME: &str = "include";
+
+ /// The name of the directory to recursively search for source files in.
+ const SOURCE_DIRNAME: &str = "cpp";
+
+ /// The warning placed in the cbindgen header.
+ const AUTOGEN_WARNING: &str =
+ "/* Automatically generated by `cbindgen`. Not intended for manual editing. */";
+
+ pub fn build_ffi() {
+ // The name of the header generated by `cbindgen`.
+ let cbindgen_hpp = &format!("{LIB_NAME}-cbindgen.hpp");
+
+ // The crate directory.
+ let crate_dir = PathBuf::from(env::var("CARGO_MANIFEST_DIR").unwrap());
+
+ // The output directory, where built artifacts should be placed.
+ let out_dir = PathBuf::from(env::var("OUT_DIR").unwrap());
+
+ // The target directory that the cargo invocation is using.
+ // Headers are symlinked into `target/include` purely for IDE purposes.
+ let target_dir = {
+ let mut dir = out_dir.clone();
+ loop {
+ if dir.ends_with("target") {
+ break dir;
+ }
+ if !dir.pop() {
+ panic!("OUT_DIR does not have parent called \"target\": {:?}", out_dir);
+ }
+ }
+ };
+
+ // The directory to read headers from.
+ let source_include_dir = crate_dir.join(INCLUDE_DIRNAME);
+
+ // The directory to place headers into.
+ let target_include_dir = out_dir.join(INCLUDE_DIRNAME);
+
+ // The directory to place symlinks to headers into. Has the fixed path "target/include".
+ let target_include_dir_fixed = target_dir.join(INCLUDE_DIRNAME);
+
+ // The directory to read source files from.
+ let source_dir = crate_dir.join(SOURCE_DIRNAME);
+
+ let headers = glob::glob(source_include_dir.join("**/*.hpp").to_str().unwrap())
+ .unwrap()
+ .collect::, _>>()
+ .unwrap();
+
+ let compilation_units = glob::glob(source_dir.join("**/*.cpp").to_str().unwrap())
+ .unwrap()
+ .collect::, _>>()
+ .unwrap();
+
+ // Tell Cargo that if the given file changes, to rerun this build script.
+ println!("cargo::rerun-if-changed={INCLUDE_DIRNAME}");
+ println!("cargo::rerun-if-changed={SOURCE_DIRNAME}");
+ println!("cargo::rerun-if-changed=src");
+ println!("cargo::rerun-if-changed=Cargo.toml");
+
+ // Cargo build script metadata, used by dependents' build scripts.
+ // The root directory containing the library archive.
+ println!("cargo::metadata=root={}", out_dir.to_str().unwrap());
+
+ // The include path defining the library's API.
+ println!("cargo::metadata=include={}", target_include_dir.to_str().unwrap());
+
+ // Generate a header containing bindings to the crate.
+ match cbindgen::Builder::new()
+ .with_pragma_once(true)
+ .with_autogen_warning(AUTOGEN_WARNING)
+ .with_no_includes()
+ .with_sys_include("cstdint")
+ .with_parse_deps(true)
+ .with_parse_include(&[
+ "sp1-stark",
+ "sp1-primitives",
+ "sp1-core-machine",
+ "p3-baby-bear",
+ "sp1-core-executor",
+ ])
+ .with_parse_extra_bindings(&["sp1-stark", "sp1-primitives", "p3-baby-bear"])
+ .rename_item("BabyBear", "BabyBearP3")
+ .include_item("MemoryRecord") // Just for convenience. Not exposed, so we need to manually do this.
+ .include_item("SyscallCode") // Required for populating the CPU columns for ECALL.
+ .include_item("SepticExtension")
+ .include_item("SepticCurve")
+ .include_item("MemoryLocalCols")
+ .include_item("MEMORY_LOCAL_INITIAL_DIGEST_POS")
+ .include_item("Ghost")
+ .include_item("MemoryInitCols")
+ .include_item("MemoryInitializeFinalizeEvent")
+ .with_namespace("sp1_core_machine_sys")
+ .with_crate(crate_dir)
+ .generate()
+ {
+ Ok(bindings) => {
+ // Write the bindings to the target include directory.
+ let header_path = target_include_dir.join(cbindgen_hpp);
+ if bindings.write_to_file(&header_path) {
+ // Symlink the header to the fixed include directory.
+ rel_symlink_file(header_path, target_include_dir_fixed.join(cbindgen_hpp));
+ }
+ }
+ Err(cbindgen::Error::ParseSyntaxError { .. }) => {} // Ignore parse errors so rust-analyzer can run.
+ Err(e) => panic!("{:?}", e),
+ }
+
+ // Copy the headers to the include directory and symlink them to the fixed include directory.
+ for header in &headers {
+ // Get the path of the header relative to the source include directory.
+ let relpath = diff_paths(header, &source_include_dir).unwrap();
+
+ // Let the destination path be the same place relative to the target include directory.
+ let dst = target_include_dir.join(&relpath);
+
+ // Create the parent directory if it does not exist.
+ if let Some(parent) = dst.parent() {
+ fs::create_dir_all(parent).unwrap();
+ }
+ fs::copy(header, &dst).unwrap();
+ rel_symlink_file(dst, target_include_dir_fixed.join(relpath));
+ }
+
+ // Use the `cc` crate to build the library and statically link it to the crate.
+ let mut cc_builder = cc::Build::new();
+ cc_builder.files(&compilation_units).include(target_include_dir);
+ cc_builder.cpp(true).std("c++20");
+ cc_builder.compile(LIB_NAME)
+ }
+
+ /// Place a relative symlink pointing to `original` at `link`.
+ fn rel_symlink_file(original: P, link: Q)
+ where
+ P: AsRef,
+ Q: AsRef,
+ {
+ #[cfg(unix)]
+ use os::unix::fs::symlink;
+ #[cfg(windows)]
+ use os::windows::fs::symlink_file as symlink;
+
+ let target_dir = link.as_ref().parent().unwrap();
+ fs::create_dir_all(target_dir).unwrap();
+ let _ = fs::remove_file(&link);
+ let relpath = diff_paths(original, target_dir).unwrap();
+ symlink(relpath, link).unwrap();
+ }
+}
diff --git a/crates/core/machine/cpp/extern.cpp b/crates/core/machine/cpp/extern.cpp
new file mode 100644
index 0000000000..509f5998a7
--- /dev/null
+++ b/crates/core/machine/cpp/extern.cpp
@@ -0,0 +1,28 @@
+#include "bb31_t.hpp"
+#include "bb31_septic_extension_t.hpp"
+#include "sys.hpp"
+
+namespace sp1_core_machine_sys {
+extern void add_sub_event_to_row_babybear(
+ const AluEvent* event,
+ AddSubCols* cols
+) {
+ AddSubCols* cols_bb31 = reinterpret_cast*>(cols);
+ add_sub::event_to_row(*event, *cols_bb31);
+}
+
+extern void memory_local_event_to_row_babybear(const MemoryLocalEvent* event, SingleMemoryLocal* cols) {
+ SingleMemoryLocal* cols_bb31 = reinterpret_cast*>(cols);
+ memory_local::event_to_row(event, cols_bb31);
+}
+
+extern void memory_global_event_to_row_babybear(const MemoryInitializeFinalizeEvent* event, const bool is_receive, MemoryInitCols* cols) {
+ MemoryInitCols* cols_bb31 = reinterpret_cast*>(cols);
+ memory_global::event_to_row(event, is_receive, cols_bb31);
+}
+
+extern void syscall_event_to_row_babybear(const SyscallEvent* event, const bool is_receive, SyscallCols* cols) {
+ SyscallCols* cols_bb31 = reinterpret_cast*>(cols);
+ syscall::event_to_row(event, is_receive, cols_bb31);
+}
+} // namespace sp1_core_machine_sys
diff --git a/crates/core/machine/include/add_sub.hpp b/crates/core/machine/include/add_sub.hpp
new file mode 100644
index 0000000000..ee98c21b4f
--- /dev/null
+++ b/crates/core/machine/include/add_sub.hpp
@@ -0,0 +1,38 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::add_sub {
+template
+__SP1_HOSTDEV__ __SP1_INLINE__ uint32_t
+populate(AddOperation& op, const uint32_t a_u32, const uint32_t b_u32) {
+ array_t a = u32_to_le_bytes(a_u32);
+ array_t b = u32_to_le_bytes(b_u32);
+ bool carry = a[0] + b[0] > 0xFF;
+ op.carry[0] = F::from_bool(carry).val;
+ carry = a[1] + b[1] + carry > 0xFF;
+ op.carry[1] = F::from_bool(carry).val;
+ carry = a[2] + b[2] + carry > 0xFF;
+ op.carry[2] = F::from_bool(carry).val;
+
+ uint32_t expected = a_u32 + b_u32;
+ write_word_from_u32_v2(op.value, expected);
+ return expected;
+}
+
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, AddSubCols& cols) {
+ bool is_add = event.opcode == Opcode::ADD;
+ cols.shard = F::from_canonical_u32(event.shard);
+ cols.is_add = F::from_bool(is_add);
+ cols.is_sub = F::from_bool(!is_add);
+
+ auto operand_1 = is_add ? event.b : event.a;
+ auto operand_2 = event.c;
+
+ populate(cols.add_operation, operand_1, operand_2);
+ write_word_from_u32_v2(cols.operand_1, operand_1);
+ write_word_from_u32_v2(cols.operand_2, operand_2);
+}
+} // namespace sp1::add_sub
\ No newline at end of file
diff --git a/crates/core/machine/include/bb31_septic_extension_t.hpp b/crates/core/machine/include/bb31_septic_extension_t.hpp
new file mode 100644
index 0000000000..9737d8bb12
--- /dev/null
+++ b/crates/core/machine/include/bb31_septic_extension_t.hpp
@@ -0,0 +1,511 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "bb31_t.hpp"
+#include
+
+#ifdef __CUDA_ARCH__
+#define FUN __host__ __device__
+#endif
+#ifndef __CUDA_ARCH__
+#define FUN inline
+#endif
+
+class bb31_cipolla_t {
+ public:
+ bb31_t real;
+ bb31_t imag;
+
+ FUN bb31_cipolla_t(bb31_t real, bb31_t imag) {
+ this->real = bb31_t(real);
+ this->imag = bb31_t(imag);
+ }
+
+ FUN static bb31_cipolla_t one() {
+ return bb31_cipolla_t(bb31_t::one(), bb31_t::zero());
+ }
+
+ FUN bb31_cipolla_t mul_ext(bb31_cipolla_t other, bb31_t nonresidue) {
+ bb31_t new_real = real * other.real + nonresidue * imag * other.imag;
+ bb31_t new_imag = real * other.imag + imag * other.real;
+ return bb31_cipolla_t(new_real, new_imag);
+ }
+
+ FUN bb31_cipolla_t pow(uint32_t exponent, bb31_t nonresidue) {
+ bb31_cipolla_t result = bb31_cipolla_t::one();
+ bb31_cipolla_t base = *this;
+
+ while(exponent) {
+ if(exponent & 1) {
+ result = result.mul_ext(base, nonresidue);
+ }
+ exponent >>= 1;
+ base = base.mul_ext(base, nonresidue);
+ }
+
+ return result;
+ }
+};
+
+namespace constants {
+ #ifdef __CUDA_ARCH__
+ __constant__ constexpr const bb31_t frobenius_const[49] = {
+ bb31_t(int(1)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)),
+ bb31_t(int(954599710)), bb31_t(int(1359279693)), bb31_t(int(566669999)), bb31_t(int(1982781815)), bb31_t(int(1735718361)), bb31_t(int(1174868538)), bb31_t(int(1120871770)),
+ bb31_t(int(862825265)), bb31_t(int(597046311)), bb31_t(int(978840770)), bb31_t(int(1790138282)), bb31_t(int(1044777201)), bb31_t(int(835869808)), bb31_t(int(1342179023)),
+ bb31_t(int(596273169)), bb31_t(int(658837454)), bb31_t(int(1515468261)), bb31_t(int(367059247)), bb31_t(int(781278880)), bb31_t(int(1544222616)), bb31_t(int(155490465)),
+ bb31_t(int(557608863)), bb31_t(int(1173670028)), bb31_t(int(1749546888)), bb31_t(int(1086464137)), bb31_t(int(803900099)), bb31_t(int(1288818584)), bb31_t(int(1184677604)),
+ bb31_t(int(763416381)), bb31_t(int(1252567168)), bb31_t(int(628856225)), bb31_t(int(1771903394)), bb31_t(int(650712211)), bb31_t(int(19417363)), bb31_t(int(57990258)),
+ bb31_t(int(1734711039)), bb31_t(int(1749813853)), bb31_t(int(1227235221)), bb31_t(int(1707730636)), bb31_t(int(424560395)), bb31_t(int(1007029514)), bb31_t(int(498034669)),
+ };
+
+ __constant__ constexpr const bb31_t double_frobenius_const[49] = {
+ bb31_t(int(1)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)),
+ bb31_t(int(1013489358)), bb31_t(int(1619071628)), bb31_t(int(304593143)), bb31_t(int(1949397349)), bb31_t(int(1564307636)), bb31_t(int(327761151)), bb31_t(int(415430835)),
+ bb31_t(int(209824426)), bb31_t(int(1313900768)), bb31_t(int(38410482)), bb31_t(int(256593180)), bb31_t(int(1708830551)), bb31_t(int(1244995038)), bb31_t(int(1555324019)),
+ bb31_t(int(1475628651)), bb31_t(int(777565847)), bb31_t(int(704492386)), bb31_t(int(1218528120)), bb31_t(int(1245363405)), bb31_t(int(475884575)), bb31_t(int(649166061)),
+ bb31_t(int(550038364)), bb31_t(int(948935655)), bb31_t(int(68722023)), bb31_t(int(1251345762)), bb31_t(int(1692456177)), bb31_t(int(1177958698)), bb31_t(int(350232928)),
+ bb31_t(int(882720258)), bb31_t(int(821925756)), bb31_t(int(199955840)), bb31_t(int(812002876)), bb31_t(int(1484951277)), bb31_t(int(1063138035)), bb31_t(int(491712810)),
+ bb31_t(int(738287111)), bb31_t(int(1955364991)), bb31_t(int(552724293)), bb31_t(int(1175775744)), bb31_t(int(341623997)), bb31_t(int(1454022463)), bb31_t(int(408193320))
+ };
+
+ __constant__ constexpr const bb31_t A_EC_LOGUP[7] = {bb31_t(int(0x31415926)), bb31_t(int(0x53589793)), bb31_t(int(0x23846264)), bb31_t(int(0x33832795)), bb31_t(int(0x02884197)), bb31_t(int(0x16939937)), bb31_t(int(0x51058209))};
+
+ __constant__ constexpr const bb31_t B_EC_LOGUP[7] = {bb31_t(int(0x74944592)), bb31_t(int(0x30781640)), bb31_t(int(0x62862089)), bb31_t(int(0x9862803)), bb31_t(int(0x48253421)), bb31_t(int(0x17067982)), bb31_t(int(0x14808651))};
+
+ __constant__ constexpr const bb31_t dummy_x[7] = {bb31_t(int(0x2738281)), bb31_t(int(0x8284590)), bb31_t(int(0x4523536)), bb31_t(int(0x0287471)), bb31_t(int(0x3526624)), bb31_t(int(0x9775724)), bb31_t(int(0x7093699))};
+ __constant__ constexpr const bb31_t dummy_y[7] = {bb31_t(int(48041908)), bb31_t(int(550064556)), bb31_t(int(415267377)), bb31_t(int(1726976249)), bb31_t(int(1253299140)), bb31_t(int(209439863)), bb31_t(int(1302309485))};
+
+ __constant__ constexpr bb31_t start_x[7] = {bb31_t(int(0x1434213)), bb31_t(int(0x5623730)), bb31_t(int(0x9504880)), bb31_t(int(0x1688724)), bb31_t(int(0x2096980)), bb31_t(int(0x7856967)), bb31_t(int(0x1875376))};
+ __constant__ constexpr bb31_t start_y[7] = {bb31_t(int(885797405)), bb31_t(int(1130275556)), bb31_t(int(567836311)), bb31_t(int(52700240)), bb31_t(int(239639200)), bb31_t(int(442612155)), bb31_t(int(1839439733))};
+
+ #endif
+
+ #ifndef __CUDA_ARCH__
+ static constexpr const bb31_t frobenius_const[49] = {
+ bb31_t(int(1)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)),
+ bb31_t(int(954599710)), bb31_t(int(1359279693)), bb31_t(int(566669999)), bb31_t(int(1982781815)), bb31_t(int(1735718361)), bb31_t(int(1174868538)), bb31_t(int(1120871770)),
+ bb31_t(int(862825265)), bb31_t(int(597046311)), bb31_t(int(978840770)), bb31_t(int(1790138282)), bb31_t(int(1044777201)), bb31_t(int(835869808)), bb31_t(int(1342179023)),
+ bb31_t(int(596273169)), bb31_t(int(658837454)), bb31_t(int(1515468261)), bb31_t(int(367059247)), bb31_t(int(781278880)), bb31_t(int(1544222616)), bb31_t(int(155490465)),
+ bb31_t(int(557608863)), bb31_t(int(1173670028)), bb31_t(int(1749546888)), bb31_t(int(1086464137)), bb31_t(int(803900099)), bb31_t(int(1288818584)), bb31_t(int(1184677604)),
+ bb31_t(int(763416381)), bb31_t(int(1252567168)), bb31_t(int(628856225)), bb31_t(int(1771903394)), bb31_t(int(650712211)), bb31_t(int(19417363)), bb31_t(int(57990258)),
+ bb31_t(int(1734711039)), bb31_t(int(1749813853)), bb31_t(int(1227235221)), bb31_t(int(1707730636)), bb31_t(int(424560395)), bb31_t(int(1007029514)), bb31_t(int(498034669))
+ };
+
+ static constexpr const bb31_t double_frobenius_const[49] = {
+ bb31_t(int(1)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)), bb31_t(int(0)),
+ bb31_t(int(1013489358)), bb31_t(int(1619071628)), bb31_t(int(304593143)), bb31_t(int(1949397349)), bb31_t(int(1564307636)), bb31_t(int(327761151)), bb31_t(int(415430835)),
+ bb31_t(int(209824426)), bb31_t(int(1313900768)), bb31_t(int(38410482)), bb31_t(int(256593180)), bb31_t(int(1708830551)), bb31_t(int(1244995038)), bb31_t(int(1555324019)),
+ bb31_t(int(1475628651)), bb31_t(int(777565847)), bb31_t(int(704492386)), bb31_t(int(1218528120)), bb31_t(int(1245363405)), bb31_t(int(475884575)), bb31_t(int(649166061)),
+ bb31_t(int(550038364)), bb31_t(int(948935655)), bb31_t(int(68722023)), bb31_t(int(1251345762)), bb31_t(int(1692456177)), bb31_t(int(1177958698)), bb31_t(int(350232928)),
+ bb31_t(int(882720258)), bb31_t(int(821925756)), bb31_t(int(199955840)), bb31_t(int(812002876)), bb31_t(int(1484951277)), bb31_t(int(1063138035)), bb31_t(int(491712810)),
+ bb31_t(int(738287111)), bb31_t(int(1955364991)), bb31_t(int(552724293)), bb31_t(int(1175775744)), bb31_t(int(341623997)), bb31_t(int(1454022463)), bb31_t(int(408193320))
+ };
+
+ static constexpr const bb31_t A_EC_LOGUP[7] = {bb31_t(int(0x31415926)), bb31_t(int(0x53589793)), bb31_t(int(0x23846264)), bb31_t(int(0x33832795)), bb31_t(int(0x02884197)), bb31_t(int(0x16939937)), bb31_t(int(0x51058209))};
+ static constexpr const bb31_t B_EC_LOGUP[7] = {bb31_t(int(0x74944592)), bb31_t(int(0x30781640)), bb31_t(int(0x62862089)), bb31_t(int(0x9862803)), bb31_t(int(0x48253421)), bb31_t(int(0x17067982)), bb31_t(int(0x14808651))};
+
+ static constexpr bb31_t dummy_x[7] = {bb31_t(int(0x2738281)), bb31_t(int(0x8284590)), bb31_t(int(0x4523536)), bb31_t(int(0x0287471)), bb31_t(int(0x3526624)), bb31_t(int(0x9775724)), bb31_t(int(0x7093699))};
+ static constexpr bb31_t dummy_y[7] = {bb31_t(int(48041908)), bb31_t(int(550064556)), bb31_t(int(415267377)), bb31_t(int(1726976249)), bb31_t(int(1253299140)), bb31_t(int(209439863)), bb31_t(int(1302309485))};
+
+ static constexpr bb31_t start_x[7] = {bb31_t(int(0x1434213)), bb31_t(int(0x5623730)), bb31_t(int(0x9504880)), bb31_t(int(0x1688724)), bb31_t(int(0x2096980)), bb31_t(int(0x7856967)), bb31_t(int(0x1875376))};
+ static constexpr bb31_t start_y[7] = {bb31_t(int(885797405)), bb31_t(int(1130275556)), bb31_t(int(567836311)), bb31_t(int(52700240)), bb31_t(int(239639200)), bb31_t(int(442612155)), bb31_t(int(1839439733))};
+
+ #endif
+}
+
+class bb31_septic_extension_t {
+ // The value of BabyBear septic extension element.
+ public:
+ bb31_t value[7];
+ static constexpr const bb31_t* frobenius_const = constants::frobenius_const;
+ static constexpr const bb31_t* double_frobenius_const = constants::double_frobenius_const;
+ static constexpr const bb31_t* A_EC_LOGUP = constants::A_EC_LOGUP;
+ static constexpr const bb31_t* B_EC_LOGUP = constants::B_EC_LOGUP;
+
+ FUN bb31_septic_extension_t() {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->value[i] = bb31_t(0);
+ }
+ }
+
+ FUN bb31_septic_extension_t(bb31_t value) {
+ this->value[0] = value;
+ for (uintptr_t i = 1 ; i < 7 ; i++) {
+ this->value[i] = bb31_t(0);
+ }
+ }
+
+ FUN bb31_septic_extension_t(bb31_t value[7]) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->value[i] = value[i];
+ }
+ }
+
+ FUN bb31_septic_extension_t(const bb31_t value[7]) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->value[i] = value[i];
+ }
+ }
+
+ static FUN bb31_septic_extension_t zero() {
+ return bb31_septic_extension_t();
+ }
+
+ static FUN bb31_septic_extension_t one() {
+ return bb31_septic_extension_t(bb31_t::one());
+ }
+
+ static FUN bb31_septic_extension_t two() {
+ return bb31_septic_extension_t(bb31_t::two());
+ }
+
+ static FUN bb31_septic_extension_t from_canonical_u32(uint32_t n) {
+ return bb31_septic_extension_t(bb31_t::from_canonical_u32(n));
+ }
+
+ FUN bb31_septic_extension_t& operator+=(const bb31_t b) {
+ value[0] += b;
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator+(bb31_septic_extension_t a, const bb31_t b) {
+ return a += b;
+ }
+
+ FUN bb31_septic_extension_t& operator+=(const bb31_septic_extension_t b) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ value[i] += b.value[i];
+ }
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator+(bb31_septic_extension_t a, const bb31_septic_extension_t b) {
+ return a += b;
+ }
+
+ FUN bb31_septic_extension_t& operator-=(const bb31_t b) {
+ value[0] -= b;
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator-(bb31_septic_extension_t a, const bb31_t b) {
+ return a -= b;
+ }
+
+ FUN bb31_septic_extension_t& operator-=(const bb31_septic_extension_t b) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ value[i] -= b.value[i];
+ }
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator-(bb31_septic_extension_t a, const bb31_septic_extension_t b) {
+ return a -= b;
+ }
+
+ FUN bb31_septic_extension_t& operator*=(const bb31_t b) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ value[i] *= b;
+ }
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator*(bb31_septic_extension_t a, const bb31_t b) {
+ return a *= b;
+ }
+
+ FUN bb31_septic_extension_t& operator*=(const bb31_septic_extension_t b) {
+ {
+ bb31_t res[13] = {};
+ for(uintptr_t i = 0 ; i < 13 ; i++) {
+ res[i] = bb31_t::zero();
+ }
+ for(uintptr_t i = 0 ; i < 7 ; i++) {
+ for(uintptr_t j = 0 ; j < 7 ; j++) {
+ res[i + j] += value[i] * b.value[j];
+ }
+ }
+ for(uintptr_t i = 7 ; i < 13 ; i++) {
+ res[i - 7] += res[i] * bb31_t::from_canonical_u32(5);
+ res[i - 6] += res[i] * bb31_t::from_canonical_u32(2);
+ }
+ for(uintptr_t i = 0 ; i < 7 ; i++) {
+ value[i] = res[i];
+ }
+ }
+ return *this;
+ }
+
+ friend FUN bb31_septic_extension_t operator*(bb31_septic_extension_t a, const bb31_septic_extension_t b) {
+ return a *= b;
+ }
+
+ FUN bool operator==(const bb31_septic_extension_t rhs) const {
+ for(uintptr_t i = 0 ; i < 7 ; i++) {
+ if(value[i] != rhs.value[i]) {
+ return false;
+ }
+ }
+ return true;
+ }
+
+ FUN bb31_septic_extension_t frobenius() const {
+ bb31_t res[7] = {};
+ res[0] = value[0];
+ for(uintptr_t i = 1 ; i < 7 ; i++) {
+ res[i] = bb31_t::zero();
+ }
+ for(uintptr_t i = 1 ; i < 7 ; i++) {
+ for(uintptr_t j = 0 ; j < 7 ; j++) {
+ res[j] += value[i] * frobenius_const[7 * i + j];
+ }
+ }
+ return bb31_septic_extension_t(res);
+
+ }
+
+ FUN bb31_septic_extension_t double_frobenius() const {
+ bb31_t res[7] = {};
+ res[0] = value[0];
+ for(uintptr_t i = 1 ; i < 7 ; i++) {
+ res[i] = bb31_t::zero();
+ }
+ for(uintptr_t i = 1 ; i < 7 ; i++) {
+ for(uintptr_t j = 0 ; j < 7 ; j++) {
+ res[j] += value[i] * double_frobenius_const[7 * i + j];
+ }
+ }
+ return bb31_septic_extension_t(res);
+
+ }
+
+ FUN bb31_septic_extension_t pow_r_1() const {
+ bb31_septic_extension_t base = frobenius();
+ base *= double_frobenius();
+ bb31_septic_extension_t base_p2 = base.double_frobenius();
+ bb31_septic_extension_t base_p4 = base_p2.double_frobenius();
+ return base * base_p2 * base_p4;
+ }
+
+ FUN bb31_t pow_r() const {
+ bb31_septic_extension_t pow_r1 = pow_r_1();
+ bb31_septic_extension_t pow_r = pow_r1 * *this;
+ return pow_r.value[0];
+ }
+
+ FUN bb31_septic_extension_t reciprocal() const {
+ bb31_septic_extension_t pow_r1 = pow_r_1();
+ bb31_septic_extension_t pow_r = pow_r1 * *this;
+ return pow_r1 * pow_r.value[0].reciprocal();
+ }
+
+ friend FUN bb31_septic_extension_t operator/(bb31_septic_extension_t a, bb31_septic_extension_t b) {
+ return a * b.reciprocal();
+ }
+
+ FUN bb31_septic_extension_t& operator/=(const bb31_septic_extension_t a) {
+ return *this *= a.reciprocal();
+ }
+
+ FUN bb31_septic_extension_t sqrt(bb31_t pow_r) const {
+ if (*this == bb31_septic_extension_t::zero()) {
+ return *this;
+ }
+
+ bb31_septic_extension_t n_iter = *this;
+ bb31_septic_extension_t n_power = *this;
+ for(uintptr_t i = 1 ; i < 30 ; i++) {
+ n_iter *= n_iter;
+ if(i >= 26) {
+ n_power *= n_iter;
+ }
+ }
+
+ bb31_septic_extension_t n_frobenius = n_power.frobenius();
+ bb31_septic_extension_t denominator = n_frobenius;
+
+ n_frobenius = n_frobenius.double_frobenius();
+ denominator *= n_frobenius;
+ n_frobenius = n_frobenius.double_frobenius();
+ denominator *= n_frobenius;
+ denominator *= *this;
+
+ bb31_t base = pow_r.reciprocal();
+ bb31_t g = bb31_t::from_canonical_u32(31);
+ bb31_t a = bb31_t::one();
+ bb31_t nonresidue = bb31_t::one() - base;
+
+ while (true) {
+ bb31_t is_square = nonresidue ^ 1006632960;
+ if (is_square != bb31_t::one()) {
+ break;
+ }
+ a *= g;
+ nonresidue = a.square() - base;
+ }
+
+ bb31_cipolla_t x = bb31_cipolla_t(a, bb31_t::one());
+ x = x.pow(1006632961, nonresidue);
+
+ return denominator * x.real;
+ }
+
+ FUN bb31_septic_extension_t universal_hash() const {
+ return *this * bb31_septic_extension_t(A_EC_LOGUP) + bb31_septic_extension_t(B_EC_LOGUP);
+ }
+
+ FUN bb31_septic_extension_t curve_formula() const {
+ bb31_septic_extension_t result = *this * *this * *this;
+ result += *this;
+ result += *this;
+ result.value[5] += bb31_t::from_canonical_u32(26);
+ return result;
+ }
+
+ FUN bool is_receive() const {
+ uint32_t limb = value[6].as_canonical_u32();
+ return 1 <= limb && limb <= (bb31_t::MOD - 1) / 2;
+ }
+
+ FUN bool is_send() const {
+ uint32_t limb = value[6].as_canonical_u32();
+ return (bb31_t::MOD + 1) / 2 <= limb && limb <= (bb31_t::MOD - 1);
+ }
+
+ FUN bool is_exception() const {
+ return value[6] == bb31_t::zero();
+ }
+};
+
+
+class bb31_septic_curve_t {
+ public:
+ bb31_septic_extension_t x;
+ bb31_septic_extension_t y;
+
+ static constexpr const bb31_t* dummy_x = constants::dummy_x;
+ static constexpr const bb31_t* dummy_y = constants::dummy_y;
+ static constexpr const bb31_t* start_x = constants::start_x;
+ static constexpr const bb31_t* start_y = constants::start_y;
+
+ FUN bb31_septic_curve_t() {
+ this->x = bb31_septic_extension_t::zero();
+ this->y = bb31_septic_extension_t::zero();
+ }
+
+ FUN bb31_septic_curve_t(bb31_septic_extension_t x, bb31_septic_extension_t y) {
+ this->x = x;
+ this->y = y;
+ }
+
+ FUN bb31_septic_curve_t(bb31_t value[14]) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->x.value[i] = value[i];
+ }
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->y.value[i] = value[i + 7];
+ }
+ }
+
+ FUN bb31_septic_curve_t(bb31_t value_x[7], bb31_t value_y[7]) {
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ this->x.value[i] = value_x[i];
+ this->y.value[i] = value_y[i];
+ }
+ }
+
+ static FUN bb31_septic_curve_t dummy_point() {
+ bb31_septic_extension_t x;
+ bb31_septic_extension_t y;
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ x.value[i] = dummy_x[i];
+ y.value[i] = dummy_y[i];
+ }
+ return bb31_septic_curve_t(x, y);
+ }
+
+ static FUN bb31_septic_curve_t start_point() {
+ bb31_septic_extension_t x;
+ bb31_septic_extension_t y;
+ for (uintptr_t i = 0 ; i < 7 ; i++) {
+ x.value[i] = start_x[i];
+ y.value[i] = start_y[i];
+ }
+ return bb31_septic_curve_t(x, y);
+ }
+
+ FUN bool is_infinity() const {
+ return x == bb31_septic_extension_t::zero() && y == bb31_septic_extension_t::zero();
+ }
+
+ FUN bb31_septic_curve_t& operator+=(const bb31_septic_curve_t b) {
+ if (b.is_infinity()) {
+ return *this;
+ }
+ if (is_infinity()) {
+ x = b.x;
+ y = b.y;
+ return *this;
+ }
+ if (x == b.x) {
+ if (y == b.y) {
+ bb31_septic_extension_t slope = (x * x * bb31_t::from_canonical_u8(3) + bb31_t::two()) / (y * bb31_t::two());
+ bb31_septic_extension_t result_x = slope * slope - x - b.x;
+ bb31_septic_extension_t result_y = slope * (x - result_x) - y;
+ x = result_x;
+ y = result_y;
+ return *this;
+ }
+ else {
+ x = bb31_septic_extension_t::zero();
+ y = bb31_septic_extension_t::zero();
+ return *this;
+ }
+ }
+ else {
+ bb31_septic_extension_t slope = (b.y - y) / (b.x - x);
+ bb31_septic_extension_t result_x = slope * slope - x - b.x;
+ bb31_septic_extension_t result_y = slope * (x - result_x) - y;
+ x = result_x;
+ y = result_y;
+ return *this;
+ }
+ }
+
+ friend FUN bb31_septic_curve_t operator+(bb31_septic_curve_t a, const bb31_septic_curve_t b) {
+ return a += b;
+ }
+
+ static FUN bb31_septic_extension_t sum_checker_x(
+ const bb31_septic_curve_t& p1,
+ const bb31_septic_curve_t& p2,
+ const bb31_septic_curve_t& p3
+ ) {
+ bb31_septic_extension_t x_diff = p2.x - p1.x;
+ bb31_septic_extension_t y_diff = p2.y - p1.y;
+ return (p1.x + p2.x + p3.x) * x_diff * x_diff - y_diff * y_diff;
+ }
+};
+
+class bb31_septic_digest_t {
+ public:
+ bb31_septic_curve_t point;
+
+ FUN bb31_septic_digest_t() {
+ this->point = bb31_septic_curve_t();
+ }
+
+ FUN bb31_septic_digest_t(bb31_t value[14]) {
+ this->point = bb31_septic_curve_t(value);
+ }
+
+ FUN bb31_septic_digest_t(bb31_septic_extension_t x, bb31_septic_extension_t y) {
+ this->point = bb31_septic_curve_t(x, y);
+ }
+
+ FUN bb31_septic_digest_t(bb31_septic_curve_t point) {
+ this->point = point;
+ }
+};
+
diff --git a/crates/core/machine/include/bb31_t.hpp b/crates/core/machine/include/bb31_t.hpp
new file mode 100644
index 0000000000..387456d422
--- /dev/null
+++ b/crates/core/machine/include/bb31_t.hpp
@@ -0,0 +1,640 @@
+// Modified by Succinct Labs
+// Copyright Supranational LLC
+// Licensed under the Apache License, Version 2.0, see LICENSE for details.
+// SPDX-License-Identifier: Apache-2.0
+
+#pragma once
+
+#include
+#include
+
+#ifdef __CUDA_ARCH__
+
+#define inline __device__ __forceinline__
+#ifdef __GNUC__
+#define asm __asm__ __volatile__
+#else
+#define asm asm volatile
+#endif
+
+class bb31_t {
+ public:
+ using mem_t = bb31_t;
+ uint32_t val;
+ static const uint32_t DEGREE = 1;
+ static const uint32_t NBITS = 31;
+ static const uint32_t MOD = 0x78000001u;
+ static const uint32_t M = 0x77ffffffu;
+ static const uint32_t RR = 0x45dddde3u;
+ static const uint32_t ONE = 0x0ffffffeu;
+ static const uint32_t MONTY_BITS = 32;
+ static const uint32_t MONTY_MU = 0x88000001;
+ static const uint32_t MONTY_MASK = ((1ULL << MONTY_BITS) - 1);
+
+ static constexpr size_t __device__ bit_length() { return 31; }
+
+ inline uint32_t& operator[](size_t i) { return val; }
+
+ inline uint32_t& operator*() { return val; }
+
+ inline const uint32_t& operator[](size_t i) const { return val; }
+
+ inline uint32_t operator*() const { return val; }
+
+ inline size_t len() const { return 1; }
+
+ inline bb31_t() {}
+
+ inline bb31_t(const uint32_t a) { val = a; }
+
+ inline bb31_t(const uint32_t* p) { val = *p; }
+
+ inline constexpr bb31_t(int a) : val(((uint64_t)a << 32) % MOD) {}
+
+ static inline const bb31_t zero() { return bb31_t(0); }
+
+ static inline const bb31_t one() { return bb31_t(ONE); }
+
+ static inline const bb31_t two() { return from_canonical_u32(2); }
+
+ static inline uint32_t to_monty(uint32_t x) {
+ return (((uint64_t)x << MONTY_BITS) % MOD);
+ }
+
+ static inline uint32_t monty_reduce(uint64_t x) {
+ uint64_t t = (x * (uint64_t)MONTY_MU) & (uint64_t)MONTY_MASK;
+ uint64_t u = t * (uint64_t)MOD;
+ uint64_t x_sub_u = x - u;
+ bool over = x < u;
+ uint32_t x_sub_u_hi = (uint32_t)(x_sub_u >> MONTY_BITS);
+ uint32_t corr = over ? MOD : 0;
+ return x_sub_u_hi + corr;
+ }
+
+ static inline bb31_t from_canonical_u32(uint32_t x) {
+ return bb31_t(to_monty(x));
+ }
+
+ static inline bb31_t from_canonical_u16(uint16_t x) {
+ return from_canonical_u32((uint32_t)x);
+ }
+
+ static inline bb31_t from_canonical_u8(uint8_t x) {
+ return from_canonical_u32((uint32_t)x);
+ }
+
+ static inline bb31_t from_bool(bool x) { return bb31_t(x * one().val); }
+
+ inline uint32_t as_canonical_u32() const {
+ return monty_reduce((uint64_t)val);
+ }
+
+ inline bb31_t exp_power_of_two(size_t log_power) {
+ bb31_t ret = *this;
+ for (size_t i = 0; i < log_power; i++) {
+ ret *= ret;
+ }
+ return ret;
+ }
+
+ inline bb31_t& operator+=(const bb31_t b) {
+ val += b.val;
+ final_sub(val);
+
+ return *this;
+ }
+
+ friend inline bb31_t operator+(bb31_t a, const bb31_t b) { return a += b; }
+
+ inline bb31_t& operator<<=(uint32_t l) {
+ while (l--) {
+ val <<= 1;
+ final_sub(val);
+ }
+
+ return *this;
+ }
+
+ friend inline bb31_t operator<<(bb31_t a, uint32_t l) { return a <<= l; }
+
+ inline bb31_t& operator>>=(uint32_t r) {
+ while (r--) {
+ val += val & 1 ? MOD : 0;
+ val >>= 1;
+ }
+
+ return *this;
+ }
+
+ friend inline bb31_t operator>>(bb31_t a, uint32_t r) { return a >>= r; }
+
+ inline bb31_t& operator-=(const bb31_t b) {
+ asm("{");
+ asm(".reg.pred %brw;");
+ asm("setp.lt.u32 %brw, %0, %1;" ::"r"(val), "r"(b.val));
+ asm("sub.u32 %0, %0, %1;" : "+r"(val) : "r"(b.val));
+ asm("@%brw add.u32 %0, %0, %1;" : "+r"(val) : "r"(MOD));
+ asm("}");
+
+ return *this;
+ }
+
+ friend inline bb31_t operator-(bb31_t a, const bb31_t b) { return a -= b; }
+
+ inline bb31_t cneg(bool flag) {
+ asm("{");
+ asm(".reg.pred %flag;");
+ asm("setp.ne.u32 %flag, %0, 0;" ::"r"(val));
+ asm("@%flag setp.ne.u32 %flag, %0, 0;" ::"r"((int)flag));
+ asm("@%flag sub.u32 %0, %1, %0;" : "+r"(val) : "r"(MOD));
+ asm("}");
+
+ return *this;
+ }
+
+ static inline bb31_t cneg(bb31_t a, bool flag) { return a.cneg(flag); }
+
+ inline bb31_t operator-() const { return cneg(*this, true); }
+
+ inline bool operator==(const bb31_t rhs) const { return val == rhs.val; }
+
+ inline bool is_one() const { return val == ONE; }
+
+ inline bool is_zero() const { return val == 0; }
+
+ inline void set_to_zero() { val = 0; }
+
+ friend inline bb31_t czero(const bb31_t a, int set_z) {
+ bb31_t ret;
+
+ asm("{");
+ asm(".reg.pred %set_z;");
+ asm("setp.ne.s32 %set_z, %0, 0;" : : "r"(set_z));
+ asm("selp.u32 %0, 0, %1, %set_z;" : "=r"(ret.val) : "r"(a.val));
+ asm("}");
+
+ return ret;
+ }
+
+ static inline bb31_t csel(const bb31_t a, const bb31_t b, int sel_a) {
+ bb31_t ret;
+
+ asm("{");
+ asm(".reg.pred %sel_a;");
+ asm("setp.ne.s32 %sel_a, %0, 0;" ::"r"(sel_a));
+ asm("selp.u32 %0, %1, %2, %sel_a;"
+ : "=r"(ret.val)
+ : "r"(a.val), "r"(b.val));
+ asm("}");
+
+ return ret;
+ }
+
+ private:
+ static inline void final_sub(uint32_t& val) {
+ asm("{");
+ asm(".reg.pred %p;");
+ asm("setp.ge.u32 %p, %0, %1;" ::"r"(val), "r"(MOD));
+ asm("@%p sub.u32 %0, %0, %1;" : "+r"(val) : "r"(MOD));
+ asm("}");
+ }
+
+ inline bb31_t& mul(const bb31_t b) {
+ uint32_t tmp[2], red;
+
+ asm("mul.lo.u32 %0, %2, %3; mul.hi.u32 %1, %2, %3;"
+ : "=r"(tmp[0]), "=r"(tmp[1])
+ : "r"(val), "r"(b.val));
+ asm("mul.lo.u32 %0, %1, %2;" : "=r"(red) : "r"(tmp[0]), "r"(M));
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %4;"
+ : "+r"(tmp[0]), "=r"(val)
+ : "r"(red), "r"(MOD), "r"(tmp[1]));
+
+ final_sub(val);
+
+ return *this;
+ }
+
+ inline uint32_t mul_by_1() const {
+ uint32_t tmp[2], red;
+
+ asm("mul.lo.u32 %0, %1, %2;" : "=r"(red) : "r"(val), "r"(M));
+ asm("mad.lo.cc.u32 %0, %2, %3, %4; madc.hi.u32 %1, %2, %3, 0;"
+ : "=r"(tmp[0]), "=r"(tmp[1])
+ : "r"(red), "r"(MOD), "r"(val));
+ return tmp[1];
+ }
+
+ public:
+ friend inline bb31_t operator*(bb31_t a, const bb31_t b) { return a.mul(b); }
+
+ inline bb31_t& operator*=(const bb31_t a) { return mul(a); }
+
+ // raise to a variable power, variable in respect to threadIdx,
+ // but mind the ^ operator's precedence!
+ inline bb31_t& operator^=(uint32_t p) {
+ bb31_t sqr = *this;
+ *this = csel(val, ONE, p & 1);
+
+#pragma unroll 1
+ while (p >>= 1) {
+ sqr.mul(sqr);
+ if (p & 1)
+ mul(sqr);
+ }
+
+ return *this;
+ }
+
+ friend inline bb31_t operator^(bb31_t a, uint32_t p) {
+ return a ^= p;
+ }
+
+ inline bb31_t operator()(uint32_t p) {
+ return *this ^ p;
+ }
+
+ // raise to a constant power, e.g. x^7, to be unrolled at compile time
+ inline bb31_t& operator^=(int p) {
+ if (p < 2)
+ asm("trap;");
+
+ bb31_t sqr = *this;
+ if ((p & 1) == 0) {
+ do {
+ sqr.mul(sqr);
+ p >>= 1;
+ } while ((p & 1) == 0);
+ *this = sqr;
+ }
+ for (p >>= 1; p; p >>= 1) {
+ sqr.mul(sqr);
+ if (p & 1)
+ mul(sqr);
+ }
+
+ return *this;
+ }
+
+ friend inline bb31_t operator^(bb31_t a, int p) {
+ return a ^= p;
+ }
+
+ inline bb31_t operator()(int p) {
+ return *this ^ p;
+ }
+
+ inline bb31_t square() { return *this * *this; }
+
+ friend inline bb31_t sqr(bb31_t a) {
+ return a.sqr();
+ }
+
+ inline bb31_t& sqr() {
+ return mul(*this);
+ }
+
+ inline void to() {
+ mul(RR);
+ }
+
+ inline void from() {
+ val = mul_by_1();
+ }
+
+ template
+ static inline bb31_t dot_product(const bb31_t a[T], const bb31_t b[T]) {
+ uint32_t acc[2];
+ size_t i = 1;
+
+ asm("mul.lo.u32 %0, %2, %3; mul.hi.u32 %1, %2, %3;"
+ : "=r"(acc[0]), "=r"(acc[1])
+ : "r"(*a[0]), "r"(*b[0]));
+ if ((T & 1) == 0) {
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i]), "r"(*b[i]));
+ i++;
+ }
+ for (; i < T; i += 2) {
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i]), "r"(*b[i]));
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i + 1]), "r"(*b[i + 1]));
+ final_sub(acc[1]);
+ }
+
+ uint32_t red;
+ asm("mul.lo.u32 %0, %1, %2;" : "=r"(red) : "r"(acc[0]), "r"(M));
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(red), "r"(MOD));
+ final_sub(acc[1]);
+
+ return acc[1];
+ }
+
+ template
+ static inline bb31_t dot_product(bb31_t a0, bb31_t b0, const bb31_t a[T - 1],
+ const bb31_t* b, size_t stride_b = 1) {
+ uint32_t acc[2];
+ size_t i = 0;
+
+ asm("mul.lo.u32 %0, %2, %3; mul.hi.u32 %1, %2, %3;"
+ : "=r"(acc[0]), "=r"(acc[1])
+ : "r"(*a0), "r"(*b0));
+ if ((T & 1) == 0) {
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i]), "r"(*b[0]));
+ i++, b += stride_b;
+ }
+ for (; i < T - 1; i += 2) {
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i]), "r"(*b[0]));
+ b += stride_b;
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(*a[i + 1]), "r"(*b[0]));
+ b += stride_b;
+ final_sub(acc[1]);
+ }
+
+ uint32_t red;
+ asm("mul.lo.u32 %0, %1, %2;" : "=r"(red) : "r"(acc[0]), "r"(M));
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %1;"
+ : "+r"(acc[0]), "+r"(acc[1])
+ : "r"(red), "r"(MOD));
+ final_sub(acc[1]);
+
+ return acc[1];
+ }
+
+ private:
+ static inline bb31_t sqr_n(bb31_t s, uint32_t n) {
+#if 0
+#pragma unroll 2
+ while (n--)
+ s.sqr();
+#else // +20% [for reciprocal()]
+#pragma unroll 2
+ while (n--) {
+ uint32_t tmp[2], red;
+
+ asm("mul.lo.u32 %0, %2, %2; mul.hi.u32 %1, %2, %2;"
+ : "=r"(tmp[0]), "=r"(tmp[1])
+ : "r"(s.val));
+ asm("mul.lo.u32 %0, %1, %2;" : "=r"(red) : "r"(tmp[0]), "r"(M));
+ asm("mad.lo.cc.u32 %0, %2, %3, %0; madc.hi.u32 %1, %2, %3, %4;"
+ : "+r"(tmp[0]), "=r"(s.val)
+ : "r"(red), "r"(MOD), "r"(tmp[1]));
+
+ if (n & 1)
+ final_sub(s.val);
+ }
+#endif
+ return s;
+ }
+
+ static inline bb31_t sqr_n_mul(bb31_t s, uint32_t n, bb31_t m) {
+ s = sqr_n(s, n);
+ s.mul(m);
+
+ return s;
+ }
+
+ public:
+ inline bb31_t reciprocal() const {
+ bb31_t x11, xff, ret = *this;
+
+ x11 = sqr_n_mul(ret, 4, ret); // 0b10001
+ ret = sqr_n_mul(x11, 1, x11); // 0b110011
+ ret = sqr_n_mul(ret, 1, x11); // 0b1110111
+ xff = sqr_n_mul(ret, 1, x11); // 0b11111111
+ ret = sqr_n_mul(ret, 8, xff); // 0b111011111111111
+ ret = sqr_n_mul(ret, 8, xff); // 0b11101111111111111111111
+ ret = sqr_n_mul(ret, 8, xff); // 0b1110111111111111111111111111111
+
+ return ret;
+ }
+
+ friend inline bb31_t operator/(int one, bb31_t a) {
+ if (one != 1)
+ asm("trap;");
+ return a.reciprocal();
+ }
+
+ friend inline bb31_t operator/(bb31_t a, bb31_t b) {
+ return a * b.reciprocal();
+ }
+
+ inline bb31_t& operator/=(const bb31_t a) {
+ return *this *= a.reciprocal();
+ }
+
+ inline bb31_t heptaroot() const {
+ bb31_t x03, x18, x1b, ret = *this;
+
+ x03 = sqr_n_mul(ret, 1, ret); // 0b11
+ x18 = sqr_n(x03, 3); // 0b11000
+ x1b = x18 * x03; // 0b11011
+ ret = x18 * x1b; // 0b110011
+ ret = sqr_n_mul(ret, 6, x1b); // 0b110011011011
+ ret = sqr_n_mul(ret, 6, x1b); // 0b110011011011011011
+ ret = sqr_n_mul(ret, 6, x1b); // 0b110011011011011011011011
+ ret = sqr_n_mul(ret, 6, x1b); // 0b110011011011011011011011011011
+ ret = sqr_n_mul(ret, 1, *this); // 0b1100110110110110110110110110111
+
+ return ret;
+ }
+
+ inline void shfl_bfly(uint32_t laneMask) {
+ val = __shfl_xor_sync(0xFFFFFFFF, val, laneMask);
+ }
+};
+
+#undef inline
+#undef asm
+// # endif // __CUDA__ARCH__
+
+#else
+
+#if defined(__GNUC__) || defined(__clang__)
+#pragma GCC diagnostic push
+#pragma GCC diagnostic ignored "-Wunused-parameter"
+#endif
+
+class bb31_t {
+ private:
+ static const uint32_t M = 0x77ffffffu;
+ static const uint32_t RR = 0x45dddde3u;
+ static const uint32_t ONE = 0x0ffffffeu;
+ static const uint32_t MONTY_BITS = 32;
+ static const uint32_t MONTY_MU = 0x88000001;
+ static const uint32_t MONTY_MASK = ((1ULL << MONTY_BITS) - 1);
+
+ public:
+ using mem_t = bb31_t;
+ uint32_t val;
+ static const uint32_t DEGREE = 1;
+ static const uint32_t NBITS = 31;
+ static const uint32_t MOD = 0x78000001;
+
+ inline bb31_t() {}
+
+ inline bb31_t(uint32_t a) : val(a) {}
+
+ inline constexpr bb31_t(int a) : val(((uint64_t)a << 32) % MOD) {}
+
+ static inline const bb31_t zero() { return bb31_t(0); }
+
+ static inline const bb31_t one() { return bb31_t(ONE); }
+
+ static inline const bb31_t two() { return bb31_t(to_monty(2)); }
+
+ static inline uint32_t to_monty(uint32_t x) {
+ return (((uint64_t)x << MONTY_BITS) % MOD);
+ }
+
+ static inline uint32_t from_monty(uint32_t x) {
+ return monty_reduce((uint64_t)x);
+ }
+
+ static inline uint32_t monty_reduce(uint64_t x) {
+ uint64_t t = (x * (uint64_t)MONTY_MU) & (uint64_t)MONTY_MASK;
+ uint64_t u = t * (uint64_t)MOD;
+ uint64_t x_sub_u = x - u;
+ bool over = x < u;
+ uint32_t x_sub_u_hi = (uint32_t)(x_sub_u >> MONTY_BITS);
+ uint32_t corr = over ? MOD : 0;
+ return x_sub_u_hi + corr;
+ }
+
+ static inline bb31_t from_canonical_u32(uint32_t x) {
+ assert(x < MOD);
+ return bb31_t(to_monty(x));
+ }
+
+ static inline bb31_t from_canonical_u16(uint16_t x) {
+ return from_canonical_u32((uint32_t)x);
+ }
+
+ static inline bb31_t from_canonical_u8(uint8_t x) {
+ return from_canonical_u32((uint32_t)x);
+ }
+
+ static inline bb31_t from_bool(bool x) { return bb31_t(x * one().val); }
+
+ inline uint32_t as_canonical_u32() const { return from_monty(val); }
+
+ inline bb31_t& operator+=(bb31_t b) {
+ val += b.val;
+ if (val >= MOD)
+ val -= MOD;
+ return *this;
+ }
+
+ inline bb31_t& operator-=(bb31_t b) {
+ if (val < b.val)
+ val += MOD;
+ val -= b.val;
+ return *this;
+ }
+
+ inline bb31_t& operator*=(bb31_t b) {
+ uint64_t long_prod = (uint64_t)val * (uint64_t)b.val;
+ val = monty_reduce(long_prod);
+ return *this;
+ }
+
+ inline bb31_t square() { return *this * *this; }
+
+ friend bb31_t operator+(bb31_t a, bb31_t b) { return a += b; }
+
+ friend bb31_t operator-(bb31_t a, bb31_t b) { return a -= b; }
+
+ friend bb31_t operator*(bb31_t a, bb31_t b) { return a *= b; }
+
+ inline bb31_t& operator<<=(uint32_t l) {
+ while (l--) {
+ val <<= 1;
+ if (val >= MOD)
+ val -= MOD;
+ }
+
+ return *this;
+ }
+
+ friend inline bb31_t operator<<(bb31_t a, uint32_t l) { return a <<= l; }
+
+ inline bb31_t& operator>>=(uint32_t r) {
+ while (r--) {
+ val += val & 1 ? MOD : 0;
+ val >>= 1;
+ }
+
+ return *this;
+ }
+
+ inline bb31_t exp_power_of_2(uint32_t power_log) const {
+ bb31_t result = *this;
+ for (uint32_t i = 0; i < power_log; ++i) {
+ result = result.square();
+ }
+ return result;
+ }
+
+ inline bb31_t reciprocal() const {
+ assert(*this != zero());
+
+ bb31_t p1 = *this;
+ bb31_t p100000000 = p1.exp_power_of_2(8);
+ bb31_t p100000001 = p100000000 * p1;
+ bb31_t p10000000000000000 = p100000000.exp_power_of_2(8);
+ bb31_t p10000000100000001 = p10000000000000000 * p100000001;
+ bb31_t p10000000100000001000 = p10000000100000001.exp_power_of_2(3);
+ bb31_t p1000000010000000100000000 = p10000000100000001000.exp_power_of_2(5);
+ bb31_t p1000000010000000100000001 = p1000000010000000100000000 * p1;
+ bb31_t p1000010010000100100001001 =
+ p1000000010000000100000001 * p10000000100000001000;
+ bb31_t p10000000100000001000000010 = p1000000010000000100000001.square();
+ bb31_t p11000010110000101100001011 =
+ p10000000100000001000000010 * p1000010010000100100001001;
+ bb31_t p100000001000000010000000100 = p10000000100000001000000010.square();
+ bb31_t p111000011110000111100001111 =
+ p100000001000000010000000100 * p11000010110000101100001011;
+ bb31_t p1110000111100001111000011110000 =
+ p111000011110000111100001111.exp_power_of_2(4);
+ bb31_t p1110111111111111111111111111111 =
+ p1110000111100001111000011110000 * p111000011110000111100001111;
+
+ return p1110111111111111111111111111111;
+ }
+
+ inline bool operator==(const bb31_t rhs) const { return val == rhs.val; }
+
+ inline bb31_t &operator^=(int b) {
+ bb31_t sqr = *this;
+ if ((b & 1) == 0)
+ *this = one();
+ while (b >>= 1) {
+ sqr = sqr.square();
+ if (b & 1)
+ *this *= sqr;
+ }
+ return *this;
+ }
+
+ friend bb31_t operator^(bb31_t a, uint32_t b) { return a ^= b; }
+
+ inline bb31_t& sqr() { return *this; }
+
+ inline void set_to_zero() { val = 0; }
+
+ inline bool is_zero() const { return val == 0; }
+};
+
+#endif // __CUDA__ARCH__
\ No newline at end of file
diff --git a/crates/core/machine/include/bitwise.hpp b/crates/core/machine/include/bitwise.hpp
new file mode 100644
index 0000000000..190e5f287c
--- /dev/null
+++ b/crates/core/machine/include/bitwise.hpp
@@ -0,0 +1,19 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::bitwise {
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, BitwiseCols& cols) {
+ cols.shard = F::from_canonical_u32(event.shard);
+ write_word_from_u32(cols.a, event.a);
+ write_word_from_u32(cols.b, event.b);
+ write_word_from_u32(cols.c, event.c);
+ cols.is_xor = F::from_bool(event.opcode == Opcode::XOR);
+ cols.is_or = F::from_bool(event.opcode == Opcode::OR);
+ cols.is_and = F::from_bool(event.opcode == Opcode::AND);
+
+ // No byte lookup yet.
+}
+} // namespace sp1::bitwise
diff --git a/crates/core/machine/include/cpu.hpp b/crates/core/machine/include/cpu.hpp
new file mode 100644
index 0000000000..41f78f9ef6
--- /dev/null
+++ b/crates/core/machine/include/cpu.hpp
@@ -0,0 +1,555 @@
+#pragma once
+
+#include
+#include
+
+#include "memory.hpp"
+#include "prelude.hpp"
+#include "utils.hpp"
+
+// namespace sp1_core_machine_sys::cpu {
+
+// template
+// __SP1_HOSTDEV__ void populate_shard_clk(const CpuEventFfi& event, CpuCols& cols) {
+// // cols.shard = F::from_canonical_u32(event.shard).val;
+// // cols.clk = F::from_canonical_u32(event.clk).val;
+
+// // const uint16_t clk_16bit_limb = (uint16_t)event.clk;
+// // const uint8_t clk_8bit_limb = (uint8_t)(event.clk >> 16);
+// // cols.clk_16bit_limb = F::from_canonical_u16(clk_16bit_limb).val;
+// // cols.clk_8bit_limb = F::from_canonical_u8(clk_8bit_limb).val;
+
+// // blu_events.add_byte_lookup_event(ByteLookupEvent::new(
+// // event.shard,
+// // U16Range,
+// // event.shard as u16,
+// // 0,
+// // 0,
+// // 0,
+// // ));
+// // blu_events.add_byte_lookup_event(ByteLookupEvent::new(
+// // event.shard,
+// // U16Range,
+// // clk_16bit_limb,
+// // 0,
+// // 0,
+// // 0,
+// // ));
+// // blu_events.add_byte_lookup_event(ByteLookupEvent::new(
+// // event.shard,
+// // ByteOpcode::U8Range,
+// // 0,
+// // 0,
+// // 0,
+// // clk_8bit_limb as u8,
+// // ));
+// }
+
+// // template
+// // __SP1_HOSTDEV__ void
+// // instruction_populate(InstructionCols& self, const Instruction& instruction) {
+// // self.opcode = F::from_canonical_u32((uint32_t)instruction.opcode).val;
+// // write_word_from_u32(self.op_a, instruction.op_a);
+// // write_word_from_u32(self.op_b, instruction.op_b);
+// // write_word_from_u32(self.op_c, instruction.op_c);
+
+// // self.op_a_0 = F::from_bool(instruction.op_a == 0).val; // 0 = Register::X0
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void
+// // selectors_populate(OpcodeSelectorCols& self, const Instruction& instruction) {
+// // self.imm_b = F::from_bool(instruction.imm_b).val;
+// // self.imm_c = F::from_bool(instruction.imm_c).val;
+
+// // switch (instruction.opcode) {
+// // // Corresponds to `instruction.is_alu_instruction()` in Rust.
+// // case Opcode::ADD:
+// // case Opcode::SUB:
+// // case Opcode::XOR:
+// // case Opcode::OR:
+// // case Opcode::AND:
+// // case Opcode::SLL:
+// // case Opcode::SRL:
+// // case Opcode::SRA:
+// // case Opcode::SLT:
+// // case Opcode::SLTU:
+// // case Opcode::MUL:
+// // case Opcode::MULH:
+// // case Opcode::MULHU:
+// // case Opcode::MULHSU:
+// // case Opcode::DIV:
+// // case Opcode::DIVU:
+// // case Opcode::REM:
+// // case Opcode::REMU:
+// // self.is_alu = F::one().val;
+// // break;
+// // // Corresponds to `instruction.is_ecall_instruction()` in Rust.
+// // case Opcode::ECALL:
+// // self.is_ecall = F::one().val;
+// // break;
+// // // Cleaner version of the `instruction.is_memory_instruction()` branch from Rust.
+// // case Opcode::LB:
+// // self.is_lb = F::one().val;
+// // break;
+// // case Opcode::LBU:
+// // self.is_lbu = F::one().val;
+// // break;
+// // case Opcode::LHU:
+// // self.is_lhu = F::one().val;
+// // break;
+// // case Opcode::LH:
+// // self.is_lh = F::one().val;
+// // break;
+// // case Opcode::LW:
+// // self.is_lw = F::one().val;
+// // break;
+// // case Opcode::SB:
+// // self.is_sb = F::one().val;
+// // break;
+// // case Opcode::SH:
+// // self.is_sh = F::one().val;
+// // break;
+// // case Opcode::SW:
+// // self.is_sw = F::one().val;
+// // break;
+// // // Cleaner version of the `instruction.is_branch_instruction()` branch from Rust.
+// // case Opcode::BEQ:
+// // self.is_beq = F::one().val;
+// // break;
+// // case Opcode::BNE:
+// // self.is_bne = F::one().val;
+// // break;
+// // case Opcode::BLT:
+// // self.is_blt = F::one().val;
+// // break;
+// // case Opcode::BGE:
+// // self.is_bge = F::one().val;
+// // break;
+// // case Opcode::BLTU:
+// // self.is_bltu = F::one().val;
+// // break;
+// // case Opcode::BGEU:
+// // self.is_bgeu = F::one().val;
+// // break;
+// // // Opcodes which each have their own branch in the original Rust function.
+// // case Opcode::JAL:
+// // self.is_jal = F::one().val;
+// // break;
+// // case Opcode::JALR:
+// // self.is_jalr = F::one().val;
+// // break;
+// // case Opcode::AUIPC:
+// // self.is_auipc = F::one().val;
+// // break;
+// // case Opcode::UNIMP:
+// // self.is_unimpl = F::one().val;
+// // break;
+// // default:
+// // break;
+// // }
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void
+// // babybear_word_populate(BabyBearWordRangeChecker& self, uint32_t value) {
+// // for (uintptr_t i = 0; i < BYTE_SIZE; ++i) {
+// // self.most_sig_byte_decomp[i] = F::from_bool((value & (1 << (i + 24))) != 0).val;
+// // }
+// // self.and_most_sig_byte_decomp_3_to_5 =
+// // F::from_bool(self.most_sig_byte_decomp[3] != 0 && self.most_sig_byte_decomp[4] != 0).val;
+// // self.and_most_sig_byte_decomp_3_to_6 =
+// // F::from_bool(self.and_most_sig_byte_decomp_3_to_5 != 0 && self.most_sig_byte_decomp[5] != 0)
+// // .val;
+// // self.and_most_sig_byte_decomp_3_to_7 =
+// // F::from_bool(self.and_most_sig_byte_decomp_3_to_6 != 0 && self.most_sig_byte_decomp[6] != 0)
+// // .val;
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void populate_memory(CpuCols& cols, const CpuEventFfi& event) {
+// // // Populate addr_word and addr_aligned columns.
+// // MemoryColumns& memory_columns = cols.opcode_specific_columns.memory;
+// // // Wraps because the types involved are unsigned integers.
+// // const uint32_t memory_addr = event.b + event.c;
+// // const uint32_t aligned_addr = memory_addr - (memory_addr % (uint32_t)WORD_SIZE);
+// // write_word_from_u32(memory_columns.addr_word, memory_addr);
+// // babybear_word_populate(memory_columns.addr_word_range_checker, memory_addr);
+// // memory_columns.addr_aligned = F::from_canonical_u32(aligned_addr).val;
+
+// // // Populate the aa_least_sig_byte_decomp columns.
+// // // assert(aligned_addr % 4 == 0);
+// // const uint8_t aligned_addr_ls_byte = (uint8_t)aligned_addr;
+// // for (uintptr_t i = 0; i < 6; ++i) {
+// // memory_columns.aa_least_sig_byte_decomp[i] =
+// // F::from_bool((aligned_addr_ls_byte & (1 << (i + 2))) != 0).val;
+// // }
+// // memory_columns.addr_word_nonce = F::from_canonical_u32(event.memory_add_nonce).val;
+
+// // // // Populate memory offsets.
+// // const uint8_t addr_offset = (uint8_t)(memory_addr % (uint32_t)WORD_SIZE);
+// // memory_columns.addr_offset = F::from_canonical_u8(addr_offset).val;
+// // memory_columns.offset_is_one = F::from_bool(addr_offset == 1).val;
+// // memory_columns.offset_is_two = F::from_bool(addr_offset == 2).val;
+// // memory_columns.offset_is_three = F::from_bool(addr_offset == 3).val;
+
+// // // If it is a load instruction, set the unsigned_mem_val column.
+// // const uint32_t mem_value = memory::unwrap_value(event.memory_record);
+
+// // // // Add event to byte lookup for byte range checking each byte in the memory addr
+// // // let addr_bytes = memory_addr.to_le_bytes();
+// // // for byte_pair in addr_bytes.chunks_exact(2) {
+// // // blu_events.add_byte_lookup_event(ByteLookupEvent {
+// // // shard: event.shard,
+// // // opcode: ByteOpcode::U8Range,
+// // // a1: 0,
+// // // a2: 0,
+// // // b: byte_pair[0],
+// // // c: byte_pair[1],
+// // // });
+// // // }
+
+// // uint32_t unsigned_mem_val = mem_value;
+// // switch (event.instruction.opcode) {
+// // case Opcode::LB:
+// // case Opcode::LBU:
+// // unsigned_mem_val = (uint32_t)(uint8_t)(mem_value >> 8 * addr_offset);
+// // break;
+// // case Opcode::LH:
+// // case Opcode::LHU:
+// // unsigned_mem_val = ((addr_offset >> 1) & 0x1) == 0 ? (mem_value & 0x0000FFFF)
+// // : (mem_value & 0xFFFF0000) >> 16;
+// // break;
+// // case Opcode::LW:
+// // // The value assigned at declaration is correct.
+// // break;
+// // default:
+// // return;
+// // }
+// // // Guard above ensures instruction is a load.
+// // write_word_from_u32(cols.unsigned_mem_val, unsigned_mem_val);
+
+// // uint8_t most_sig_mem_value_byte;
+// // switch (event.instruction.opcode) {
+// // case Opcode::LB:
+
+// // most_sig_mem_value_byte = (uint8_t)unsigned_mem_val;
+// // break;
+// // case Opcode::LH:
+// // most_sig_mem_value_byte = (uint8_t)(unsigned_mem_val >> 8);
+// // break;
+// // default:
+// // // The load instruction is unsigned.
+// // // Set the `mem_value_is_pos_not_x0` composite flag.
+// // cols.mem_value_is_pos_not_x0 =
+// // F::from_bool(event.instruction.op_a != 0).val; // 0 = Register::X0
+// // return;
+// // }
+// // // Guard above ensures the load instruction is signed.
+// // for (intptr_t i = BYTE_SIZE - 1; i >= 0; --i) {
+// // memory_columns.most_sig_byte_decomp[i] =
+// // F::from_canonical_u32(most_sig_mem_value_byte >> i & 0x1).val;
+// // }
+// // bool mem_value_is_pos_not_x0 = memory_columns.most_sig_byte_decomp[7] == F::zero().val;
+// // if (!mem_value_is_pos_not_x0) {
+// // cols.mem_value_is_neg_not_x0 =
+// // F::from_bool(event.instruction.op_a != 0).val; // 0 = Register::X0
+// // cols.unsigned_mem_val_nonce = F::from_canonical_u32(event.memory_sub_nonce).val;
+// // }
+// // // Set the `mem_value_is_pos_not_x0` composite flag.
+// // cols.mem_value_is_pos_not_x0 = F::from_bool(mem_value_is_pos_not_x0).val;
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void populate_branch(CpuCols& cols, const CpuEventFfi& event) {
+// // // let branch_columns = cols.opcode_specific_columns.branch_mut();
+// // BranchCols& branch_columns = cols.opcode_specific_columns.branch;
+
+// // Opcode opcode = event.instruction.opcode;
+// // const bool use_signed_comparison = opcode == Opcode::BLT || opcode == Opcode::BGE;
+
+// // const bool a_eq_b = event.a == event.b;
+// // const bool a_lt_b =
+// // use_signed_comparison ? ((int32_t)event.a < (int32_t)event.b) : (event.a < event.b);
+// // const bool a_gt_b =
+// // use_signed_comparison ? ((int32_t)event.a > (int32_t)event.b) : (event.a > event.b);
+
+// // branch_columns.a_lt_b_nonce = F::from_canonical_u32(event.branch_lt_nonce).val;
+// // branch_columns.a_gt_b_nonce = F::from_canonical_u32(event.branch_gt_nonce).val;
+
+// // branch_columns.a_eq_b = F::from_bool(a_eq_b).val;
+// // branch_columns.a_lt_b = F::from_bool(a_lt_b).val;
+// // branch_columns.a_gt_b = F::from_bool(a_gt_b).val;
+
+// // bool branching;
+// // switch (opcode) {
+// // case Opcode::BEQ:
+// // branching = a_eq_b;
+// // break;
+// // case Opcode::BNE:
+// // branching = !a_eq_b;
+// // break;
+// // case Opcode::BLT:
+// // case Opcode::BLTU:
+// // branching = a_lt_b;
+// // break;
+// // case Opcode::BGE:
+// // case Opcode::BGEU:
+// // branching = a_eq_b || a_gt_b;
+// // break;
+// // default:
+// // // Precondition violated.
+// // assert(false);
+// // break;
+// // }
+
+// // // Unsigned arithmetic wraps.
+// // const uint32_t next_pc = event.pc + event.c;
+// // write_word_from_u32(branch_columns.pc, event.pc);
+// // write_word_from_u32(branch_columns.next_pc, next_pc);
+// // babybear_word_populate(branch_columns.pc_range_checker, event.pc);
+// // babybear_word_populate(branch_columns.next_pc_range_checker, next_pc);
+
+// // if (branching) {
+// // cols.branching = F::one().val;
+// // branch_columns.next_pc_nonce = F::from_canonical_u32(event.branch_add_nonce).val;
+// // } else {
+// // cols.not_branching = F::one().val;
+// // }
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void populate_jump(CpuCols& cols, const CpuEventFfi& event) {
+// // // let jump_columns = cols.opcode_specific_columns.jump_mut();
+// // JumpCols& jump_columns = cols.opcode_specific_columns.jump;
+
+// // switch (event.instruction.opcode) {
+// // case Opcode::JAL: {
+// // // Unsigned arithmetic wraps.
+// // uint32_t next_pc = event.pc + event.b;
+// // babybear_word_populate(jump_columns.op_a_range_checker, event.a);
+// // write_word_from_u32(jump_columns.pc, event.pc);
+// // babybear_word_populate(jump_columns.pc_range_checker, event.pc);
+// // write_word_from_u32(jump_columns.next_pc, next_pc);
+// // babybear_word_populate(jump_columns.next_pc_range_checker, next_pc);
+// // jump_columns.jal_nonce = F::from_canonical_u32(event.jump_jal_nonce).val;
+// // break;
+// // }
+// // case Opcode::JALR: {
+// // // Unsigned arithmetic wraps.
+// // uint32_t next_pc = event.b + event.c;
+// // babybear_word_populate(jump_columns.op_a_range_checker, event.a);
+// // write_word_from_u32(jump_columns.next_pc, next_pc);
+// // babybear_word_populate(jump_columns.next_pc_range_checker, next_pc);
+// // jump_columns.jalr_nonce = F::from_canonical_u32(event.jump_jalr_nonce).val;
+// // break;
+// // }
+// // default:
+// // // Precondition violated.
+// // assert(false);
+// // break;
+// // }
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void populate_auipc(CpuCols& cols, const CpuEventFfi& event) {
+// // AuipcCols& auipc_columns = cols.opcode_specific_columns.auipc;
+
+// // write_word_from_u32(auipc_columns.pc, event.pc);
+// // babybear_word_populate(auipc_columns.pc_range_checker, event.pc);
+// // auipc_columns.auipc_nonce = F::from_canonical_u32(event.auipc_nonce).val;
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ void
+// // is_zero_operation_populate_from_field_element(IsZeroOperation& self, F a) {
+// // if (a == F::zero()) {
+// // self.inverse = F::zero().val;
+// // self.result = F::one().val;
+// // } else {
+// // self.inverse = a.reciprocal().val;
+// // self.result = F::zero().val;
+// // }
+// // // F is_zero = F::one() - F(self.inverse) * a;
+// // // assert(is_zero == F(self.result));
+// // // let is_zero = one.clone() - cols.inverse * a.clone();
+// // // builder.when(is_real.clone()).assert_eq(is_zero, cols.result);
+
+// // // let prod = self.inverse * a;
+// // // debug_assert!(prod == F::one() || prod == F::zero());
+// // // (a == F::zero()) as u32
+// // }
+
+// // template
+// // __SP1_HOSTDEV__ bool populate_ecall(CpuCols& cols, const CpuEventFfi& event) {
+// // bool is_halt = false;
+
+// // // The send_to_table column is the 1st entry of the op_a_access column prev_value field.
+// // // Look at `ecall_eval` in cpu/air/mod.rs for the corresponding constraint and
+// // // explanation.
+// // EcallCols& ecall_cols = cols.opcode_specific_columns.ecall;
+
+// // cols.ecall_mul_send_to_table = cols.op_a_access.prev_value._0[1];
+
+// // F syscall_id = F(cols.op_a_access.prev_value._0[0]);
+
+// // // In the following statements, truncating to `uint8_t` is the equivalent of the
+// // // `SyscallCode::get_syscall_id` calls from the Rust code.
+
+// // // Populate `is_enter_unconstrained`.
+// // is_zero_operation_populate_from_field_element(
+// // ecall_cols.is_enter_unconstrained,
+// // syscall_id - F::from_canonical_u8((uint8_t)SyscallCode::ENTER_UNCONSTRAINED)
+// // );
+
+// // // Populate `is_hint_len`.
+// // is_zero_operation_populate_from_field_element(
+// // ecall_cols.is_hint_len,
+// // syscall_id - F::from_canonical_u8((uint8_t)SyscallCode::HINT_LEN)
+// // );
+
+// // // Populate `is_halt`.
+// // is_zero_operation_populate_from_field_element(
+// // ecall_cols.is_halt,
+// // syscall_id - F::from_canonical_u8((uint8_t)SyscallCode::HALT)
+// // );
+
+// // // Populate `is_commit`.
+// // is_zero_operation_populate_from_field_element(
+// // ecall_cols.is_commit,
+// // syscall_id - F::from_canonical_u8((uint8_t)SyscallCode::COMMIT)
+// // );
+
+// // // Populate `is_commit_deferred_proofs`.
+// // is_zero_operation_populate_from_field_element(
+// // ecall_cols.is_commit_deferred_proofs,
+// // syscall_id - F::from_canonical_u8((uint8_t)SyscallCode::COMMIT_DEFERRED_PROOFS)
+// // );
+
+// // // If the syscall is `COMMIT` or `COMMIT_DEFERRED_PROOFS`, set the index bitmap and
+// // // digest word.
+// // if (syscall_id
+// // == F::from_canonical_u8((uint8_t)SyscallCode::COMMIT
+// // ) // Comment to make my editor format nicely...
+// // || syscall_id == F::from_canonical_u8((uint8_t)SyscallCode::COMMIT_DEFERRED_PROOFS)) {
+// // uint32_t digest_idx = word_to_u32(cols.op_b_access.access.value);
+// // ecall_cols.index_bitmap[digest_idx] = F::one().val;
+// // }
+
+// // // Write the syscall nonce.
+// // ecall_cols.syscall_nonce = F::from_canonical_u32(event.syscall_nonce).val;
+
+// // is_halt = syscall_id == F::from_canonical_u32((uint8_t)SyscallCode::HALT);
+
+// // // For halt and commit deferred proofs syscalls, we need to baby bear range check one of
+// // // it's operands.
+// // if (is_halt) {
+// // write_word_from_u32(ecall_cols.operand_to_check, event.b);
+// // babybear_word_populate(ecall_cols.operand_range_check_cols, event.b);
+// // cols.ecall_range_check_operand = F::one().val;
+// // }
+
+// // if (syscall_id == F::from_canonical_u32((uint8_t)SyscallCode::COMMIT_DEFERRED_PROOFS)) {
+// // write_word_from_u32(ecall_cols.operand_to_check, event.c);
+// // babybear_word_populate(ecall_cols.operand_range_check_cols, event.c);
+// // cols.ecall_range_check_operand = F::one().val;
+// // }
+
+// // return is_halt;
+// // }
+
+// template
+// __SP1_HOSTDEV__ void event_to_row(const CpuEventFfi& event, CpuCols& cols) {
+// // // Populate shard and clk columns.
+// // populate_shard_clk(event, cols);
+
+// // // Populate the nonce.
+// // cols.nonce = F::from_canonical_u32(event.alu_nonce).val;
+
+// // // Populate basic fields.
+// // cols.pc = F::from_canonical_u32(event.pc).val;
+// // cols.next_pc = F::from_canonical_u32(event.next_pc).val;
+// // instruction_populate(cols.instruction, event.instruction);
+// // // cols.instruction.populate(event.instruction);
+// // selectors_populate(cols.selectors, event.instruction);
+// // // cols.selectors.populate(event.instruction);
+// // write_word_from_u32(cols.op_a_access.access.value, event.a);
+// // write_word_from_u32(cols.op_b_access.access.value, event.b);
+// // write_word_from_u32(cols.op_c_access.access.value, event.c);
+
+// // // // Populate memory accesses for a, b, and c.
+// // // The function guards against the record being `None`.
+// // memory::populate_read_write(cols.op_a_access, event.a_record);
+// // if (event.b_record.tag == OptionMemoryRecordEnum::Tag::Read) {
+// // memory::populate_read(cols.op_b_access, event.b_record.read._0);
+// // }
+// // if (event.c_record.tag == OptionMemoryRecordEnum::Tag::Read) {
+// // memory::populate_read(cols.op_c_access, event.c_record.read._0);
+// // }
+
+// // // // Populate range checks for a.
+// // // let a_bytes = cols
+// // // .op_a_access
+// // // .access
+// // // .val
+// // // .0
+// // // .iter()
+// // // .map(|x| x.as_canonical_u32())
+// // // .collect::>();
+// // // blu_events.add_byte_lookup_event(ByteLookupEvent {
+// // // shard: event.shard,
+// // // opcode: ByteOpcode::U8Range,
+// // // a1: 0,
+// // // a2: 0,
+// // // b: a_bytes[0] as u8,
+// // // c: a_bytes[1] as u8,
+// // // });
+// // // blu_events.add_byte_lookup_event(ByteLookupEvent {
+// // // shard: event.shard,
+// // // opcode: ByteOpcode::U8Range,
+// // // a1: 0,
+// // // a2: 0,
+// // // b: a_bytes[2] as u8,
+// // // c: a_bytes[3] as u8,
+// // // });
+
+// // // Populate memory accesses for reading from memory.
+// // // `event.memory` appears to be vestigial.
+// // // assert_eq!(event.memory_record.is_some(), event.memory.is_some());
+// // // The function guards against the record being `None`.
+// // memory::populate_read_write(
+// // cols.opcode_specific_columns.memory.memory_access,
+// // event.memory_record
+// // );
+
+// // // Populate memory, branch, jump, and auipc specific fields.
+// // const bool is_memory = opcode_utils::is_memory(event.instruction.opcode);
+// // const bool is_branch = opcode_utils::is_branch(event.instruction.opcode);
+// // const bool is_jump = opcode_utils::is_jump(event.instruction.opcode);
+// // const bool is_auipc = event.instruction.opcode == Opcode::AUIPC;
+// // const bool is_ecall = event.instruction.opcode == Opcode::ECALL;
+// // // Calculated by `populate_ecall`, if called.
+// // bool is_halt = false;
+// // // Unlike the Rust code, we guard outside the function bodies so we can reuse the booleans.
+// // if (is_memory) {
+// // populate_memory(cols, event);
+// // }
+// // if (is_branch) {
+// // populate_branch(cols, event);
+// // }
+// // if (is_jump) {
+// // populate_jump(cols, event);
+// // }
+// // if (is_auipc) {
+// // populate_auipc(cols, event);
+// // }
+// // if (is_ecall) {
+// // is_halt = populate_ecall(cols, event);
+// // }
+
+// // cols.is_sequential_instr = F::from_bool(!(is_branch || is_jump || is_halt)).val;
+
+// // // Assert that the instruction is not a no-op.
+// // cols.is_real = F::one().val;
+// }
+// } // namespace sp1::cpu
\ No newline at end of file
diff --git a/crates/core/machine/include/lt.hpp b/crates/core/machine/include/lt.hpp
new file mode 100644
index 0000000000..3c83c144f4
--- /dev/null
+++ b/crates/core/machine/include/lt.hpp
@@ -0,0 +1,100 @@
+#pragma once
+
+#include
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::lt {
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, LtCols& cols) {
+ array_t a = u32_to_le_bytes(event.a);
+ array_t b = u32_to_le_bytes(event.b);
+ array_t c = u32_to_le_bytes(event.c);
+ cols.shard = F::from_canonical_u32(event.shard).val;
+ word_from_le_bytes(cols.a, a);
+ word_from_le_bytes(cols.b, b);
+ word_from_le_bytes(cols.c, c);
+
+ // If this is SLT, mask the MSB of b & c before computing cols.bits.
+ uint8_t masked_b = b[3] & 0x7f;
+ uint8_t masked_c = c[3] & 0x7f;
+ cols.b_masked = F::from_canonical_u8(masked_b);
+ cols.c_masked = F::from_canonical_u8(masked_c);
+
+ // // Send the masked interaction.
+ // blu.add_byte_lookup_event(ByteLookupEvent {
+ // shard: event.shard,
+ // channel: event.channel,
+ // opcode: ByteOpcode::AND,
+ // a1: masked_b as u16,
+ // a2: 0,
+ // b: b[3],
+ // c: 0x7f,
+ // });
+ // blu.add_byte_lookup_event(ByteLookupEvent {
+ // shard: event.shard,
+ // channel: event.channel,
+ // opcode: ByteOpcode::AND,
+ // a1: masked_c as u16,
+ // a2: 0,
+ // b: c[3],
+ // c: 0x7f,
+ // });
+
+ array_t b_comp = b;
+ array_t c_comp = c;
+ if (event.opcode == Opcode::SLT) {
+ b_comp[3] = masked_b;
+ c_comp[3] = masked_c;
+ }
+
+ // Set the byte equality flags.
+ intptr_t i = 3;
+ while (true) {
+ uint8_t b_byte = b_comp[i];
+ uint8_t c_byte = c_comp[i];
+ if (b_byte != c_byte) {
+ cols.byte_flags[i] = F::one();
+ cols.sltu = F::from_bool(b_byte < c_byte);
+ F b_byte_f = F::from_canonical_u8(b_byte);
+ F c_byte_f = F::from_canonical_u8(c_byte);
+ cols.not_eq_inv = (b_byte_f - c_byte_f).reciprocal();
+ cols.comparison_bytes[0] = b_byte_f;
+ cols.comparison_bytes[1] = c_byte_f;
+ break;
+ }
+ if (i == 0) {
+ // The equality `b_comp == c_comp` holds.
+ cols.is_comp_eq = F::one();
+ break;
+ }
+ --i;
+ }
+
+ cols.msb_b = F::from_bool((b[3] >> 7) & 1);
+ cols.msb_c = F::from_bool((c[3] >> 7) & 1);
+ cols.is_sign_eq = F::from_bool(event.opcode != Opcode::SLT || cols.msb_b == cols.msb_c);
+
+ cols.is_slt = F::from_bool(event.opcode == Opcode::SLT);
+ cols.is_sltu = F::from_bool(event.opcode == Opcode::SLTU);
+
+ cols.bit_b = (F(cols.msb_b) * F(cols.is_slt));
+ cols.bit_c = (F(cols.msb_c) * F(cols.is_slt));
+
+ // if (F(cols.a._0[0]) != F(cols.bit_b) * (F::one() - F(cols.bit_c)) + F(cols.is_sign_eq) * F(cols.sltu))
+ // {
+ // std::exit(1);
+ // }
+
+ // blu.add_byte_lookup_event(ByteLookupEvent {
+ // shard: event.shard,
+ // channel: event.channel,
+ // opcode: ByteOpcode::LTU,
+ // a1: cols.sltu.as_canonical_u32() as u16,
+ // a2: 0,
+ // b: cols.comparison_bytes[0].as_canonical_u32() as u8,
+ // c: cols.comparison_bytes[1].as_canonical_u32() as u8,
+ // });
+}
+} // namespace sp1::lt
\ No newline at end of file
diff --git a/crates/core/machine/include/memory.hpp b/crates/core/machine/include/memory.hpp
new file mode 100644
index 0000000000..216a3a449d
--- /dev/null
+++ b/crates/core/machine/include/memory.hpp
@@ -0,0 +1,116 @@
+#pragma once
+
+#include
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+// namespace sp1_core_machine_sys::memory {
+// __SP1_HOSTDEV__ __SP1_INLINE__ uint32_t unwrap_value(const OptionMemoryRecordEnum& record) {
+// switch (record.tag) {
+// case OptionMemoryRecordEnum::Tag::Read:
+// return record.read._0.value;
+// case OptionMemoryRecordEnum::Tag::Write:
+// return record.write._0.value;
+// default:
+// // Either the tag is `None` or it is an invalid value.
+// assert(false);
+// }
+// // Unreachable.
+// return 0;
+// }
+
+// template
+// __SP1_HOSTDEV__ void populate_access(
+// MemoryAccessCols& self,
+// const MemoryRecord& current_record,
+// const MemoryRecord& prev_record
+// ) {
+// write_word_from_u32(self.value, current_record.value);
+
+// self.prev_shard = F::from_canonical_u32(prev_record.shard).val;
+// self.prev_clk = F::from_canonical_u32(prev_record.timestamp).val;
+
+// // Fill columns used for verifying current memory access time value is greater than
+// // previous's.
+// const bool use_clk_comparison = prev_record.shard == current_record.shard;
+// self.compare_clk = F::from_bool(use_clk_comparison).val;
+// const uint32_t prev_time_value = use_clk_comparison ? prev_record.timestamp : prev_record.shard;
+// const uint32_t current_time_value =
+// use_clk_comparison ? current_record.timestamp : current_record.shard;
+
+// const uint32_t diff_minus_one = current_time_value - prev_time_value - 1;
+// const uint16_t diff_16bit_limb = (uint16_t)(diff_minus_one);
+// self.diff_16bit_limb = F::from_canonical_u16(diff_16bit_limb).val;
+// const uint8_t diff_8bit_limb = (uint8_t)(diff_minus_one >> 16);
+// self.diff_8bit_limb = F::from_canonical_u8(diff_8bit_limb).val;
+
+// // let shard = current_record.shard;
+
+// // // Add a byte table lookup with the 16Range op.
+// // output.add_u16_range_check(shard, diff_16bit_limb);
+
+// // // Add a byte table lookup with the U8Range op.
+// // output.add_u8_range_check(shard, 0, diff_8bit_limb as u8);
+// }
+
+// template
+// __SP1_HOSTDEV__ void
+// populate_read(MemoryReadCols& self, const MemoryReadRecord& record) {
+// const MemoryRecord current_record = {
+// .shard = record.shard,
+// .timestamp = record.timestamp,
+// .value = record.value,
+// };
+// const MemoryRecord prev_record = {
+// .shard = record.prev_shard,
+// .timestamp = record.prev_timestamp,
+// .value = record.value,
+// };
+// populate_access(self.access, current_record, prev_record);
+// }
+
+// template
+// __SP1_HOSTDEV__ void populate_read_write(
+// MemoryReadWriteCols& self,
+// const OptionMemoryRecordEnum& record
+// ) {
+// if (record.tag == OptionMemoryRecordEnum::Tag::None) {
+// return;
+// }
+// MemoryRecord current_record;
+// MemoryRecord prev_record;
+// switch (record.tag) {
+// case OptionMemoryRecordEnum::Tag::Read:
+// current_record = {
+// .shard = record.read._0.shard,
+// .timestamp = record.read._0.timestamp,
+// .value = record.read._0.value,
+// };
+// prev_record = {
+// .shard = record.read._0.prev_shard,
+// .timestamp = record.read._0.prev_timestamp,
+// .value = record.read._0.value,
+// };
+// break;
+// case OptionMemoryRecordEnum::Tag::Write:
+// current_record = {
+// .shard = record.write._0.shard,
+// .timestamp = record.write._0.timestamp,
+// .value = record.write._0.value,
+// };
+// prev_record = {
+// .shard = record.write._0.prev_shard,
+// .timestamp = record.write._0.prev_timestamp,
+// .value = record.write._0.prev_value,
+// };
+// break;
+// default:
+// // Unreachable. `None` case guarded above.
+// assert(false);
+// break;
+// }
+// write_word_from_u32(self.prev_value, prev_record.value);
+// populate_access(self.access, current_record, prev_record);
+// }
+// } // namespace sp1::memory
\ No newline at end of file
diff --git a/crates/core/machine/include/memory_global.hpp b/crates/core/machine/include/memory_global.hpp
new file mode 100644
index 0000000000..6859686a56
--- /dev/null
+++ b/crates/core/machine/include/memory_global.hpp
@@ -0,0 +1,37 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+#include "bb31_septic_extension_t.hpp"
+#include "memory_local.hpp"
+
+namespace sp1_core_machine_sys::memory_global {
+ template
+ __SP1_HOSTDEV__ void event_to_row(const MemoryInitializeFinalizeEvent* event, const bool is_receive, MemoryInitCols* cols) {
+ MemoryRecord record;
+ if (is_receive) {
+ record.shard = event->shard;
+ record.timestamp = event->timestamp;
+ record.value = event->value;
+ } else {
+ record.shard = 0;
+ record.timestamp = 0;
+ record.value = event->value;
+ }
+ // We populate only the things in the first loop of generate_trace here. The second loop is handled in the kernel directly.
+ sp1_core_machine_sys::memory_local::populate_memory(&cols->global_interaction_cols, &record, event->addr, is_receive);
+ cols->addr = F::from_canonical_u32(event->addr);
+ for(uintptr_t i = 0 ; i < 32 ; i++) {
+ cols->addr_bits.bits[i] = F::from_canonical_u32(((event->addr) >> i) & 1);
+ }
+ cols->addr_bits.and_most_sig_byte_decomp_3_to_5 = cols->addr_bits.bits[27] * cols->addr_bits.bits[28];
+ cols->addr_bits.and_most_sig_byte_decomp_3_to_6 = cols->addr_bits.and_most_sig_byte_decomp_3_to_5 * cols->addr_bits.bits[29];
+ cols->addr_bits.and_most_sig_byte_decomp_3_to_7 = cols->addr_bits.and_most_sig_byte_decomp_3_to_6 * cols->addr_bits.bits[30];
+ cols->shard = F::from_canonical_u32(event->shard);
+ cols->timestamp = F::from_canonical_u32(event->timestamp);
+ for(uintptr_t i = 0 ; i < 32 ; i++) {
+ cols->value[i] = F::from_canonical_u32(((event->value) >> i) & 1);
+ }
+ cols->is_real = F::from_canonical_u32(event->used);
+ }
+} // namespace sp1::memory_local
\ No newline at end of file
diff --git a/crates/core/machine/include/memory_local.hpp b/crates/core/machine/include/memory_local.hpp
new file mode 100644
index 0000000000..1d799208b2
--- /dev/null
+++ b/crates/core/machine/include/memory_local.hpp
@@ -0,0 +1,82 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+#include "bb31_septic_extension_t.hpp"
+
+namespace sp1_core_machine_sys::memory_local {
+ template __SP1_HOSTDEV__ void populate_memory(GlobalInteractionOperation* cols, const MemoryRecord* record, const uint32_t& addr, bool is_receive) {
+ EF7 x_start;
+
+ {
+ x_start.value[0] = F::from_canonical_u32(record->shard + (1 << 24));
+ x_start.value[1] = F::from_canonical_u32(record->timestamp);
+ x_start.value[2] = F::from_canonical_u32(addr);
+ x_start.value[3] = F::from_canonical_u32(record->value & 255);
+ x_start.value[4] = F::from_canonical_u32((record->value >> 8) & 255);
+ x_start.value[5] = F::from_canonical_u32((record->value >> 16) & 255);
+ x_start.value[6] = F::from_canonical_u32((record->value >> 24) & 255);
+ }
+
+ #pragma unroll(1)
+ for(uint32_t offset = 0 ; offset < 256 ; offset++) {
+ EF7 x_trial = x_start.universal_hash();
+ EF7 y_sq = x_trial.curve_formula();
+ F y_sq_pow_r = y_sq.pow_r();
+ F is_square = y_sq_pow_r ^ 1006632960;
+ if(is_square == F::one()) {
+ EF7 y = y_sq.sqrt(y_sq_pow_r);
+ if (y.is_exception()) {
+ x_start += F::from_canonical_u32(1 << 16);
+ continue;
+ }
+ if (y.is_receive() != is_receive) {
+ y = EF7::zero() - y;
+ }
+ // x_trial, y
+ for(uint32_t idx = 0 ; idx < 8 ; idx++ ) {
+ cols->offset_bits[idx] = F::from_canonical_u32((offset >> idx) & 1);
+ }
+ for(uintptr_t i = 0 ; i < 7 ; i++) {
+ cols->x_coordinate._0[i] = x_trial.value[i];
+ cols->y_coordinate._0[i] = y.value[i];
+ }
+ uint32_t range_check_value;
+ if (is_receive) {
+ range_check_value = y.value[6].as_canonical_u32() - 1;
+ } else {
+ range_check_value = y.value[6].as_canonical_u32() - (F::MOD + 1) / 2;
+ }
+ F top_4_bits = F::zero();
+ for(uint32_t idx = 0 ; idx < 30 ; idx++) {
+ cols->y6_bit_decomp[idx] = F::from_canonical_u32((range_check_value >> idx) & 1);
+ if (idx >= 26) {
+ top_4_bits += cols->y6_bit_decomp[idx];
+ }
+ }
+ top_4_bits -= F::from_canonical_u32(4);
+ cols->range_check_witness = top_4_bits.reciprocal();
+ return;
+ }
+ x_start += F::from_canonical_u32(1 << 16);
+ }
+ assert(false);
+ }
+
+ template
+ __SP1_HOSTDEV__ void event_to_row(const MemoryLocalEvent* event, SingleMemoryLocal* cols) {
+ populate_memory(&cols->initial_global_interaction_cols, &event->initial_mem_access, event->addr, true);
+ populate_memory(&cols->final_global_interaction_cols, &event->final_mem_access, event->addr, false);
+ cols->addr = F::from_canonical_u32(event->addr);
+
+ cols->initial_shard = F::from_canonical_u32(event->initial_mem_access.shard);
+ cols->initial_clk = F::from_canonical_u32(event->initial_mem_access.timestamp);
+ write_word_from_u32_v2(cols->initial_value, event->initial_mem_access.value);
+
+ cols->final_shard = F::from_canonical_u32(event->final_mem_access.shard);
+ cols->final_clk = F::from_canonical_u32(event->final_mem_access.timestamp);
+ write_word_from_u32_v2(cols->final_value, event->final_mem_access.value);
+
+ cols->is_real = F::one();
+ }
+} // namespace sp1::memory_local
diff --git a/crates/core/machine/include/mul.hpp b/crates/core/machine/include/mul.hpp
new file mode 100644
index 0000000000..efa564275f
--- /dev/null
+++ b/crates/core/machine/include/mul.hpp
@@ -0,0 +1,111 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::mul {
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, MulCols& cols) {
+ // // Ensure that the opcode is MUL, MULHU, MULH, or MULHSU.
+ // assert!(
+ // event.opcode == Opcode::MUL
+ // || event.opcode == Opcode::MULHU
+ // || event.opcode == Opcode::MULH
+ // || event.opcode == Opcode::MULHSU
+ // );
+
+ const array_t a = u32_to_le_bytes(event.a);
+ const array_t b = u32_to_le_bytes(event.b);
+ const array_t c = u32_to_le_bytes(event.c);
+
+ // Handle b and c's signs.
+ {
+ uint8_t b_msb = get_msb(b);
+ cols.b_msb = F::from_canonical_u8(b_msb).val;
+ uint8_t c_msb = get_msb(c);
+ cols.c_msb = F::from_canonical_u8(c_msb).val;
+
+ // If b is signed and it is negative, sign extend b.
+ if ((event.opcode == Opcode::MULH || event.opcode == Opcode::MULHSU) && b_msb == 1) {
+ cols.b_sign_extend = F::one().val;
+ }
+
+ // If c is signed and it is negative, sign extend c.
+ if (event.opcode == Opcode::MULH && c_msb == 1) {
+ cols.c_sign_extend = F::one().val;
+ }
+
+ // // Insert the MSB lookup events.
+ // {
+ // let words = [b_word, c_word];
+ // let mut blu_events: Vec = vec![];
+ // for word in words.iter() {
+ // let most_significant_byte = word[WORD_SIZE - 1];
+ // blu_events.push(ByteLookupEvent {
+ // shard: event.shard,
+ // opcode: ByteOpcode::MSB,
+ // a1: get_msb(*word) as u16,
+ // a2: 0,
+ // b: most_significant_byte,
+ // c: 0,
+ // });
+ // }
+ // record.add_byte_lookup_events(blu_events);
+ // }
+ }
+
+ // Required for the following logic to correctly multiply.
+ static_assert(2 * WORD_SIZE == LONG_WORD_SIZE);
+
+ array_t product {};
+ for (uintptr_t i = 0; i < WORD_SIZE; ++i) {
+ for (uintptr_t j = 0; j < WORD_SIZE; ++j) {
+ product[i + j] += (uint32_t)b[i] * (uint32_t)c[j];
+ }
+ if (cols.c_sign_extend != F::zero().val) {
+ for (uintptr_t j = WORD_SIZE; j < LONG_WORD_SIZE - i; ++j) {
+ product[i + j] += (uint32_t)b[i] * (uint32_t)0xFF;
+ }
+ }
+ }
+ if (cols.b_sign_extend != F::zero().val) {
+ for (uintptr_t i = WORD_SIZE; i < LONG_WORD_SIZE; ++i) {
+ for (uintptr_t j = 0; j < LONG_WORD_SIZE - i; ++j) {
+ product[i + j] += (uint32_t)0xFF * (uint32_t)c[j];
+ }
+ }
+ }
+
+ // Calculate the correct product using the `product` array. We store the
+ // correct carry value for verification.
+ const uint32_t base = 1 << BYTE_SIZE;
+ array_t carry {};
+ for (uintptr_t i = 0; i < LONG_WORD_SIZE; ++i) {
+ carry[i] = product[i] / base;
+ product[i] %= base;
+ if (i + 1 < LONG_WORD_SIZE) {
+ product[i + 1] += carry[i];
+ }
+ cols.carry[i] = F::from_canonical_u32(carry[i]).val;
+ }
+
+ for (uintptr_t i = 0; i < LONG_WORD_SIZE; ++i) {
+ cols.product[i] = F::from_canonical_u32(product[i]).val;
+ }
+ word_from_le_bytes(cols.a, a);
+ word_from_le_bytes(cols.b, b);
+ word_from_le_bytes(cols.c, c);
+ cols.is_real = F::one().val;
+ cols.is_mul = F::from_bool(event.opcode == Opcode::MUL).val;
+ cols.is_mulh = F::from_bool(event.opcode == Opcode::MULH).val;
+ cols.is_mulhu = F::from_bool(event.opcode == Opcode::MULHU).val;
+ cols.is_mulhsu = F::from_bool(event.opcode == Opcode::MULHSU).val;
+ cols.shard = F::from_canonical_u32(event.shard).val;
+
+ // // Range check.
+ // {
+ // record.add_u16_range_checks(event.shard, &carry.map(|x| x as u16));
+ // record.add_u8_range_checks(event.shard, &product.map(|x| x as u8));
+ // }
+}
+} // namespace sp1::mul
diff --git a/crates/core/machine/include/prelude.hpp b/crates/core/machine/include/prelude.hpp
new file mode 100644
index 0000000000..1f2b0db1e1
--- /dev/null
+++ b/crates/core/machine/include/prelude.hpp
@@ -0,0 +1,23 @@
+#pragma once
+
+#include "sp1-core-machine-sys-cbindgen.hpp"
+
+#ifndef __CUDACC__
+ #define __SP1_HOSTDEV__
+ #define __SP1_INLINE__ inline
+ #include
+
+namespace sp1_core_machine_sys {
+template
+using array_t = std::array;
+} // namespace sp1
+#else
+ #define __SP1_HOSTDEV__ __host__ __device__
+ #define __SP1_INLINE__
+ #include
+
+namespace sp1_core_machine_sys {
+template
+using array_t = cuda::std::array;
+} // namespace sp1
+#endif
diff --git a/crates/core/machine/include/sll.hpp b/crates/core/machine/include/sll.hpp
new file mode 100644
index 0000000000..5f24a575a0
--- /dev/null
+++ b/crates/core/machine/include/sll.hpp
@@ -0,0 +1,66 @@
+#pragma once
+
+#include
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::sll {
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, ShiftLeftCols& cols) {
+ array_t a = u32_to_le_bytes(event.a);
+ array_t b = u32_to_le_bytes(event.b);
+ array_t c = u32_to_le_bytes(event.c);
+ cols.shard = F::from_canonical_u32(event.shard).val;
+ word_from_le_bytes(cols.a, a);
+ word_from_le_bytes(cols.b, b);
+ word_from_le_bytes(cols.c, c);
+ cols.is_real = F::one().val;
+ for (uintptr_t i = 0; i < BYTE_SIZE; ++i) {
+ cols.c_least_sig_byte[i] = F::from_canonical_u32((event.c >> i) & 1).val;
+ }
+
+ // Variables for bit shifting.
+ uintptr_t num_bits_to_shift = event.c % BYTE_SIZE;
+ for (uintptr_t i = 0; i < BYTE_SIZE; ++i) {
+ cols.shift_by_n_bits[i] = F::from_bool(num_bits_to_shift == i).val;
+ }
+
+ uint32_t bit_shift_multiplier = 1 << num_bits_to_shift;
+ cols.bit_shift_multiplier = F::from_canonical_u32(bit_shift_multiplier).val;
+
+ uint32_t carry = 0;
+ uint32_t base = 1 << BYTE_SIZE;
+
+ array_t bit_shift_result;
+ array_t bit_shift_result_carry;
+ for (uintptr_t i = 0; i < WORD_SIZE; ++i) {
+ uint32_t v = b[i] * bit_shift_multiplier + carry;
+ carry = v / base;
+ bit_shift_result[i] = (uint8_t)(v % base);
+ cols.bit_shift_result[i] = F::from_canonical_u8(bit_shift_result[i]).val;
+ bit_shift_result_carry[i] = (uint8_t)carry;
+ cols.bit_shift_result_carry[i] = F::from_canonical_u8(bit_shift_result_carry[i]).val;
+ }
+
+ // // Variables for byte shifting.
+ uintptr_t num_bytes_to_shift = (uintptr_t)(event.c & 0b11111) / BYTE_SIZE;
+ for (uintptr_t i = 0; i < WORD_SIZE; ++i) {
+ cols.shift_by_n_bytes[i] = F::from_bool(num_bytes_to_shift == i).val;
+ }
+
+ // // Range checks.
+ // {
+ // blu.add_u8_range_checks(event.shard, event.channel, &bit_shift_result);
+ // blu.add_u8_range_checks(event.shard, event.channel, &bit_shift_result_carry);
+ // }
+
+ // // Sanity check.
+ // for i in num_bytes_to_shift..WORD_SIZE {
+ // debug_assert_eq!(
+ // cols.bit_shift_result[i - num_bytes_to_shift],
+ // F::from_canonical_u8(a[i])
+ // );
+ // }
+}
+} // namespace sp1::sll
\ No newline at end of file
diff --git a/crates/core/machine/include/sr.hpp b/crates/core/machine/include/sr.hpp
new file mode 100644
index 0000000000..a5806fadb1
--- /dev/null
+++ b/crates/core/machine/include/sr.hpp
@@ -0,0 +1,106 @@
+#pragma once
+
+#include
+
+#include "prelude.hpp"
+#include "utils.hpp"
+
+namespace sp1_core_machine_sys::sr {
+template
+__SP1_HOSTDEV__ void event_to_row(const AluEvent& event, ShiftRightCols& cols) {
+ // Initialize cols with basic operands and flags derived from the current event.
+ {
+ cols.shard = F::from_canonical_u32(event.shard).val;
+ write_word_from_u32(cols.a, event.a);
+ write_word_from_u32(cols.b, event.b);
+ write_word_from_u32(cols.c, event.c);
+ cols.b_msb = F::from_canonical_u32((event.b >> 31) & 1).val;
+ cols.is_srl = F::from_bool(event.opcode == Opcode::SRL).val;
+ cols.is_sra = F::from_bool(event.opcode == Opcode::SRA).val;
+ cols.is_real = F::one().val;
+
+ for (uintptr_t i = 0; i < BYTE_SIZE; ++i) {
+ cols.c_least_sig_byte[i] = F::from_canonical_u32((event.c >> i) & 1).val;
+ }
+
+ // // Insert the MSB lookup event.
+ // let most_significant_byte = event.b.to_le_bytes()[WORD_SIZE - 1];
+ // blu.add_byte_lookup_events(vec![ByteLookupEvent {
+ // shard: event.shard,
+ // opcode: ByteOpcode::MSB,
+ // a1: ((most_significant_byte >> 7) & 1) as u16,
+ // a2: 0,
+ // b: most_significant_byte,
+ // c: 0,
+ // }]);
+ }
+
+ // Note that we take the least significant 5 bits per the RISC-V spec.
+ const uintptr_t num_bytes_to_shift = (event.c % 32) / BYTE_SIZE;
+ const uintptr_t num_bits_to_shift = (event.c % 32) % BYTE_SIZE;
+
+ // Byte shifting.
+ // Zero-initialize the array.
+ array_t byte_shift_result {};
+ {
+ for (uintptr_t i = 0; i < WORD_SIZE; ++i) {
+ cols.shift_by_n_bytes[i] = F::from_bool(num_bytes_to_shift == i).val;
+ }
+ // Sign extension is necessary only for arithmetic right shift.
+ array_t sign_extended_b = event.opcode == Opcode::SRA
+ ? u64_to_le_bytes((int64_t)(int32_t)event.b)
+ : u64_to_le_bytes((uint64_t)event.b);
+
+ for (uintptr_t i = 0; i < LONG_WORD_SIZE - num_bytes_to_shift; ++i) {
+ byte_shift_result[i] = sign_extended_b[i + num_bytes_to_shift];
+ cols.byte_shift_result[i] =
+ F::from_canonical_u8(sign_extended_b[i + num_bytes_to_shift]).val;
+ }
+ }
+
+ // Bit shifting.
+ {
+ for (uintptr_t i = 0; i < BYTE_SIZE; ++i) {
+ cols.shift_by_n_bits[i] = F::from_bool(num_bits_to_shift == i).val;
+ }
+ const uint32_t carry_multiplier = 1 << (8 - num_bits_to_shift);
+ uint32_t last_carry = 0;
+ array_t bit_shift_result;
+ array_t shr_carry_output_carry;
+ array_t shr_carry_output_shifted_byte;
+ for (intptr_t i = LONG_WORD_SIZE - 1; i >= 0; --i) {
+ auto [shift, carry] = shr_carry(byte_shift_result[i], num_bits_to_shift);
+
+ // let byte_event = ByteLookupEvent {
+ // shard: event.shard,
+ // opcode: ByteOpcode::ShrCarry,
+ // a1: shift as u16,
+ // a2: carry,
+ // b: byte_shift_result[i],
+ // c: num_bits_to_shift as u8,
+ // };
+ // blu.add_byte_lookup_event(byte_event);
+
+ shr_carry_output_carry[i] = carry;
+ cols.shr_carry_output_carry[i] = F::from_canonical_u8(carry).val;
+
+ shr_carry_output_shifted_byte[i] = shift;
+ cols.shr_carry_output_shifted_byte[i] = F::from_canonical_u8(shift).val;
+
+ uint8_t res = (uint8_t)(((uint32_t)shift + last_carry * carry_multiplier) & 0xFF);
+ bit_shift_result[i] = res;
+ cols.bit_shift_result[i] = F::from_canonical_u8(res).val;
+ last_carry = (uint32_t)carry;
+ }
+ // for (uintptr_t i = 0; i < WORD_SIZE; ++i)
+ // {
+ // assert(cols.a[i] == cols.bit_shift_result[i]);
+ // }
+ // // Range checks.
+ // blu.add_u8_range_checks(event.shard, &byte_shift_result);
+ // blu.add_u8_range_checks(event.shard, &bit_shift_result);
+ // blu.add_u8_range_checks(event.shard, &shr_carry_output_carry);
+ // blu.add_u8_range_checks(event.shard, &shr_carry_output_shifted_byte);
+ }
+}
+} // namespace sp1::sr
\ No newline at end of file
diff --git a/crates/core/machine/include/sys.hpp b/crates/core/machine/include/sys.hpp
new file mode 100644
index 0000000000..2a6f56166b
--- /dev/null
+++ b/crates/core/machine/include/sys.hpp
@@ -0,0 +1,14 @@
+#pragma once
+
+#include "add_sub.hpp"
+#include "bitwise.hpp"
+#include "cpu.hpp"
+#include "lt.hpp"
+#include "memory.hpp"
+#include "mul.hpp"
+#include "sll.hpp"
+#include "sp1-core-machine-sys-cbindgen.hpp"
+#include "sr.hpp"
+#include "memory_local.hpp"
+#include "memory_global.hpp"
+#include "syscall.hpp"
diff --git a/crates/core/machine/include/syscall.hpp b/crates/core/machine/include/syscall.hpp
new file mode 100644
index 0000000000..caa91b32f8
--- /dev/null
+++ b/crates/core/machine/include/syscall.hpp
@@ -0,0 +1,77 @@
+#pragma once
+
+#include "prelude.hpp"
+#include "utils.hpp"
+#include "bb31_septic_extension_t.hpp"
+
+namespace sp1_core_machine_sys::syscall {
+ template __SP1_HOSTDEV__ void populate_syscall(GlobalInteractionOperation* cols, const SyscallEvent* event, bool is_receive) {
+ EF7 x_start;
+
+ {
+ x_start.value[0] = F::from_canonical_u32(event->shard + 8 * (1 << 24));
+ x_start.value[1] = F::from_canonical_u32(event->clk & ((1 << 16) - 1));
+ x_start.value[2] = F::from_canonical_u32((event->clk) >> 16);
+ x_start.value[3] = F::from_canonical_u32(event->syscall_id);
+ x_start.value[4] = F::from_canonical_u32(event->arg1);
+ x_start.value[5] = F::from_canonical_u32(event->arg2);
+ x_start.value[6] = F::zero();
+ }
+
+ #pragma unroll(1)
+ for(uint32_t offset = 0 ; offset < 256 ; offset++) {
+ EF7 x_trial = x_start.universal_hash();
+ EF7 y_sq = x_trial.curve_formula();
+ F y_sq_pow_r = y_sq.pow_r();
+ F is_square = y_sq_pow_r ^ 1006632960;
+ if(is_square == F::one()) {
+ EF7 y = y_sq.sqrt(y_sq_pow_r);
+ if (y.is_exception()) {
+ x_start += F::from_canonical_u32(1 << 16);
+ continue;
+ }
+ if (y.is_receive() != is_receive) {
+ y = EF7::zero() - y;
+ }
+ // x_trial, y
+ for(uint32_t idx = 0 ; idx < 8 ; idx++ ) {
+ cols->offset_bits[idx] = F::from_canonical_u32((offset >> idx) & 1);
+ }
+ for(uintptr_t i = 0 ; i < 7 ; i++) {
+ cols->x_coordinate._0[i] = x_trial.value[i];
+ cols->y_coordinate._0[i] = y.value[i];
+ }
+ uint32_t range_check_value;
+ if (is_receive) {
+ range_check_value = y.value[6].as_canonical_u32() - 1;
+ } else {
+ range_check_value = y.value[6].as_canonical_u32() - (F::MOD + 1) / 2;
+ }
+ F top_4_bits = F::zero();
+ for(uint32_t idx = 0 ; idx < 30 ; idx++) {
+ cols->y6_bit_decomp[idx] = F::from_canonical_u32((range_check_value >> idx) & 1);
+ if (idx >= 26) {
+ top_4_bits += cols->y6_bit_decomp[idx];
+ }
+ }
+ top_4_bits -= F::from_canonical_u32(4);
+ cols->range_check_witness = top_4_bits.reciprocal();
+ return;
+ }
+ x_start += F::from_canonical_u32(1 << 16);
+ }
+ assert(false);
+ }
+
+ template
+ __SP1_HOSTDEV__ void event_to_row(const SyscallEvent* event, const bool is_receive, SyscallCols* cols) {
+ populate_syscall(&cols->global_interaction_cols, event, is_receive);
+ cols->shard = F::from_canonical_u32(event->shard);
+ cols->clk_16 = F::from_canonical_u32(event->clk & ((1 << 16) - 1));
+ cols->clk_8 = F::from_canonical_u32((event->clk) >> 16);
+ cols->syscall_id = F::from_canonical_u32(event->syscall_id);
+ cols->arg1 = F::from_canonical_u32(event->arg1);
+ cols->arg2 = F::from_canonical_u32(event->arg2);
+ cols->is_real = F::one();
+ }
+} // namespace sp1::memory_local
diff --git a/crates/core/machine/include/utils.hpp b/crates/core/machine/include/utils.hpp
new file mode 100644
index 0000000000..7f798ce314
--- /dev/null
+++ b/crates/core/machine/include/utils.hpp
@@ -0,0 +1,134 @@
+#pragma once
+
+#include
+#include
+
+#include "prelude.hpp"
+
+namespace sp1_core_machine_sys {
+
+// Compiles to a no-op with -O3 and the like.
+__SP1_HOSTDEV__ __SP1_INLINE__ array_t u32_to_le_bytes(uint32_t n) {
+ return {
+ (uint8_t)(n >> 8 * 0),
+ (uint8_t)(n >> 8 * 1),
+ (uint8_t)(n >> 8 * 2),
+ (uint8_t)(n >> 8 * 3),
+ };
+}
+
+__SP1_HOSTDEV__ __SP1_INLINE__ array_t u64_to_le_bytes(uint64_t n) {
+ return {
+ (uint8_t)(n >> 8 * 0),
+ (uint8_t)(n >> 8 * 1),
+ (uint8_t)(n >> 8 * 2),
+ (uint8_t)(n >> 8 * 3),
+ (uint8_t)(n >> 8 * 4),
+ (uint8_t)(n >> 8 * 5),
+ (uint8_t)(n >> 8 * 6),
+ (uint8_t)(n >> 8 * 7),
+ };
+}
+
+/// Shifts a byte to the right and returns both the shifted byte and the bits that carried.
+__SP1_HOSTDEV__ __SP1_INLINE__ std::tuple
+shr_carry(uint8_t input, uint8_t rotation) {
+ uint8_t c_mod = rotation & 0x7;
+ if (c_mod != 0) {
+ uint8_t res = input >> c_mod;
+ uint8_t c_mod_comp = 8 - c_mod;
+ uint8_t carry = (uint8_t)(input << c_mod_comp) >> c_mod_comp;
+ return {res, carry};
+ } else {
+ return {input, 0};
+ }
+}
+
+template