diff --git a/mopro-msm/.gitignore b/mopro-msm/.gitignore index f619967..b82b9bd 100644 --- a/mopro-msm/.gitignore +++ b/mopro-msm/.gitignore @@ -24,4 +24,4 @@ src/msm/metal_msm/shader/**/*.ir src/msm/metal_msm/shader/**/*.lib # Metal shader constants file -src/msm/metal_msm/shader/**/constants.metal \ No newline at end of file +src/msm/metal_msm/shader/constants.metal \ No newline at end of file diff --git a/mopro-msm/src/msm/metal_msm/shader/bigint/bigint_add_unsafe.metal.ir b/mopro-msm/src/msm/metal_msm/shader/bigint/bigint_add_unsafe.metal.ir deleted file mode 100644 index 9863771..0000000 Binary files a/mopro-msm/src/msm/metal_msm/shader/bigint/bigint_add_unsafe.metal.ir and /dev/null differ diff --git a/mopro-msm/src/msm/metal_msm/shader/constants.metal b/mopro-msm/src/msm/metal_msm/shader/constants.metal deleted file mode 100644 index 650b981..0000000 --- a/mopro-msm/src/msm/metal_msm/shader/constants.metal +++ /dev/null @@ -1,8 +0,0 @@ -// THIS FILE IS AUTOGENERATED BY shader.rs -#define NUM_LIMBS 16 -#define NUM_LIMBS_WIDE 17 -#define LOG_LIMB_SIZE 16 -#define TWO_POW_WORD_SIZE 65536 -#define MASK 65535 -#define N0 25481 -#define NSAFE 1 diff --git a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_unsafe.rs b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_unsafe.rs index 748b1f3..c3001aa 100644 --- a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_unsafe.rs +++ b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_unsafe.rs @@ -5,29 +5,31 @@ use crate::msm::metal_msm::host::gpu::{ }; use crate::msm::metal_msm::host::shader::{compile_metal, write_constants}; use crate::msm::metal_msm::utils::limbs_conversion::{FromLimbs, ToLimbs}; -use ark_ff::{BigInt, BigInteger}; +use ark_ff::{BigInt, BigInteger, UniformRand}; +use ark_std::rand; use metal::*; #[test] #[serial_test::serial] pub fn test_bigint_add_unsafe() { - let log_limb_size = 13; - let num_limbs = 20; - - // Create two test numbers (equivalent to the previous hex values) - let a = BigInt::new([ - 0x0000000100000001, - 0x0000000000000000, - 0x1800a1101800a110, - 0x0000000d0000000d, - ]); - let b = a.clone(); // Same value as a for this test - - let mut expected = a.clone(); - let overflow = expected.add_with_carry(&b); - - // We are testing add_unsafe, so the sum should not overflow - assert!(!overflow); + // adjusted by bn254 scalar bits and mont_mul cios optimal limb size + let log_limb_size = 16; + let num_limbs = 16; + + // Create two test numbers that do not cause overflow + let mut rng = rand::thread_rng(); + let (a, b, expected) = loop { + let a = BigInt::rand(&mut rng); + let b = BigInt::rand(&mut rng); + + let mut expected = a.clone(); + let overflow = expected.add_with_carry(&b); + + // Break the loop if addition does not overflow + if !overflow { + break (a, b, expected); + } + }; let device = get_default_device(); let a_buf = create_buffer(&device, &a.to_limbs(num_limbs, log_limb_size)); @@ -87,7 +89,9 @@ pub fn test_bigint_add_unsafe() { command_buffer.wait_until_completed(); let result_limbs: Vec = read_buffer(&result_buf, num_limbs); - let result = BigInt::from_limbs(&result_limbs, log_limb_size); + let expected_limbs = expected.to_limbs(num_limbs, log_limb_size); + assert_eq!(result_limbs, expected_limbs); - assert!(result.eq(&expected)); + let result = BigInt::from_limbs(&result_limbs, log_limb_size); + assert_eq!(result, expected); } diff --git a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_wide.rs b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_wide.rs index 6e2713d..6524102 100644 --- a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_wide.rs +++ b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_add_wide.rs @@ -5,33 +5,31 @@ use crate::msm::metal_msm::host::gpu::{ }; use crate::msm::metal_msm::host::shader::{compile_metal, write_constants}; use crate::msm::metal_msm::utils::limbs_conversion::{FromLimbs, ToLimbs}; -use ark_ff::{BigInt, BigInteger}; +use ark_ff::{BigInt, BigInteger, UniformRand}; +use ark_std::rand; use metal::*; #[test] #[serial_test::serial] -pub fn test_bigint_add() { - let log_limb_size = 13; - let num_limbs = 20; - - // Create two large numbers that will overflow when added - let a = BigInt::new([ - 0xffffffffffffffff, - 0xffffffffffffffff, - 0xffffffffffffffff, - 0xffffffffffffffff, - ]); - let b = BigInt::new([ - 0x1000000000000000, - 0x0000000000000000, - 0x0000000000000000, - 0x0000000000000000, - ]); - - let mut expected = a.clone(); - - let overflow = expected.add_with_carry(&b); - assert!(overflow); +pub fn test_bigint_add_no_overflow() { + // adjusted by bn254 scalar bits and mont_mul cios optimal limb size + let log_limb_size = 16; + let num_limbs = 16; + + // Create two test numbers that do not cause overflow + let mut rng = rand::thread_rng(); + let (a, b, expected) = loop { + let a = BigInt::rand(&mut rng); + let b = BigInt::rand(&mut rng); + + let mut expected = a.clone(); + let overflow = expected.add_with_carry(&b); + + // Break the loop if addition does not overflow + if !overflow { + break (a, b, expected); + } + }; let device = get_default_device(); let a_buf = create_buffer(&device, &a.to_limbs(num_limbs, log_limb_size)); @@ -91,35 +89,35 @@ pub fn test_bigint_add() { command_buffer.commit(); command_buffer.wait_until_completed(); - let result_limbs: Vec = read_buffer(&result_buf, num_limbs + 1); - let result = BigInt::from_limbs(&result_limbs, log_limb_size); + let result_limbs: Vec = read_buffer(&result_buf, num_limbs); + let expected_limbs = expected.to_limbs(num_limbs, log_limb_size); + assert_eq!(result_limbs, expected_limbs); - assert!(result.eq(&expected)); + let result = BigInt::from_limbs(&result_limbs, log_limb_size); + assert_eq!(result, expected); } #[test] #[serial_test::serial] -pub fn test_bigint_add_no_overflow() { - let log_limb_size = 13; - let num_limbs = 20; - - // Create two numbers that won't overflow when added - let a = BigInt::new([ - 0x0000000000000000, - 0x0000000000000000, - 0x0000000000000000, - 0x0000000000000001, - ]); - let b = BigInt::new([ - 0x0000000000000000, - 0x0000000000000000, - 0x0000000000000000, - 0x0000000000000002, - ]); - - let mut expected = a.clone(); - let overflow = expected.add_with_carry(&b); - assert!(!overflow); +pub fn test_bigint_add_overflow() { + // adjusted by bn254 scalar bits and mont_mul cios optimal limb size + let log_limb_size = 16; + let num_limbs = 16; + + // Create two test numbers that cause overflow + let mut rng = rand::thread_rng(); + let (a, b, expected) = loop { + let a = BigInt::rand(&mut rng); + let b = BigInt::rand(&mut rng); + + let mut expected = a.clone(); + let overflow = expected.add_with_carry(&b); + + // Break the loop if addition overflow + if overflow { + break (a, b, expected); + } + }; let device = get_default_device(); let a_buf = create_buffer(&device, &a.to_limbs(num_limbs, log_limb_size)); @@ -179,8 +177,10 @@ pub fn test_bigint_add_no_overflow() { command_buffer.commit(); command_buffer.wait_until_completed(); - let result_limbs: Vec = read_buffer(&result_buf, num_limbs + 1); - let result = BigInt::from_limbs(&result_limbs, log_limb_size); + let result_limbs: Vec = read_buffer(&result_buf, num_limbs); + let expected_limbs = expected.to_limbs(num_limbs, log_limb_size); + assert_eq!(result_limbs, expected_limbs); - assert!(result.eq(&expected)); + let result = BigInt::from_limbs(&result_limbs, log_limb_size); + assert_eq!(result, expected); } diff --git a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_sub.rs b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_sub.rs index 8a5390e..3d0a077 100644 --- a/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_sub.rs +++ b/mopro-msm/src/msm/metal_msm/tests/bigint/bigint_sub.rs @@ -1,33 +1,41 @@ // adapted from: https://github.com/geometryxyz/msl-secp256k1 -use core::borrow; - use crate::msm::metal_msm::host::gpu::{ create_buffer, create_empty_buffer, get_default_device, read_buffer, }; use crate::msm::metal_msm::host::shader::{compile_metal, write_constants}; use crate::msm::metal_msm::utils::limbs_conversion::{FromLimbs, ToLimbs}; -use ark_ff::{BigInt, BigInteger}; +use ark_ff::{BigInt, BigInteger, UniformRand}; +use ark_std::rand; use metal::*; #[test] #[serial_test::serial] -pub fn test_bigint_sub() { - let log_limb_size = 13; - let num_limbs = 20; - - let mut a = BigInt::new([0xf09f8fb3, 0xefb88fe2, 0x808df09f, 0x8c880010]); - let b = BigInt::new([0xf09f8fb3, 0xefb88fe2, 0x808df09f, 0x8c880001]); +pub fn test_bigint_sub_no_underflow() { + // adjusted by bn254 scalar bits and mont_mul cios optimal limb size + let log_limb_size = 16; + let num_limbs = 16; + + // Create two test numbers that do not cause underflow + let mut rng = rand::thread_rng(); + let (a, b, expected) = loop { + let a = BigInt::rand(&mut rng); + let b = BigInt::rand(&mut rng); + + let mut expected = a.clone(); + let underflow = expected.sub_with_borrow(&b); + + // Break the loop if subtraction does not underflow + if !underflow { + break (a, b, expected); + } + }; let device = get_default_device(); let a_buf = create_buffer(&device, &a.to_limbs(num_limbs, log_limb_size)); let b_buf = create_buffer(&device, &b.to_limbs(num_limbs, log_limb_size)); let result_buf = create_empty_buffer(&device, num_limbs); - // perform a - b - let _borrow = a.sub_with_borrow(&b); - let expected_limbs = a.to_limbs(num_limbs, log_limb_size); - let command_queue = device.new_command_queue(); let command_buffer = command_queue.new_command_buffer(); @@ -81,46 +89,38 @@ pub fn test_bigint_sub() { command_buffer.wait_until_completed(); let result_limbs: Vec = read_buffer(&result_buf, num_limbs); - let result = BigInt::from_limbs(&result_limbs, log_limb_size); + let expected_limbs = expected.to_limbs(num_limbs, log_limb_size); + assert_eq!(result_limbs, expected_limbs); - assert!(result_limbs.eq(&expected_limbs)); - assert!(result.eq(&a)); + let result = BigInt::from_limbs(&result_limbs, log_limb_size); + assert_eq!(result, expected); } #[test] #[serial_test::serial] fn test_bigint_sub_underflow() { - let device = Device::system_default().expect("no device found"); - let num_limbs = 20; - let log_limb_size = 13; - - // Create smaller number a and larger number b - let mut a = BigInt::from_u32(100); - let b = BigInt::from_u32(200); - - let a_limbs = a.to_limbs(num_limbs, log_limb_size); - let b_limbs = b.to_limbs(num_limbs, log_limb_size); - - let a_buf = device.new_buffer_with_data( - unsafe { std::mem::transmute(a_limbs.as_ptr()) }, - (a_limbs.len() * std::mem::size_of::()) as u64, - MTLResourceOptions::StorageModeShared, - ); - - let b_buf = device.new_buffer_with_data( - unsafe { std::mem::transmute(b_limbs.as_ptr()) }, - (b_limbs.len() * std::mem::size_of::()) as u64, - MTLResourceOptions::StorageModeShared, - ); - - let result_buf = device.new_buffer( - (num_limbs * std::mem::size_of::()) as u64, - MTLResourceOptions::StorageModeShared, - ); + let num_limbs = 16; + let log_limb_size = 16; + + // Create two test numbers that cause underflow + let mut rng = rand::thread_rng(); + let (a, b, expected) = loop { + let a = BigInt::rand(&mut rng); + let b = BigInt::rand(&mut rng); + + let mut expected = a.clone(); + let underflow = expected.sub_with_borrow(&b); + + // Break the loop if subtraction does not underflow + if underflow { + break (a, b, expected); + } + }; - // Expected result is 2^256 - 100 (since we're doing a - b where b > a) - let _expected = a.sub_with_borrow(&b); - let expected_limbs = a.to_limbs(num_limbs, log_limb_size); + let device = get_default_device(); + let a_buf = create_buffer(&device, &a.to_limbs(num_limbs, log_limb_size)); + let b_buf = create_buffer(&device, &b.to_limbs(num_limbs, log_limb_size)); + let result_buf = create_empty_buffer(&device, num_limbs); let command_queue = device.new_command_queue(); let command_buffer = command_queue.new_command_buffer(); @@ -175,8 +175,9 @@ fn test_bigint_sub_underflow() { command_buffer.wait_until_completed(); let result_limbs: Vec = read_buffer(&result_buf, num_limbs); - let result = BigInt::from_limbs(&result_limbs, log_limb_size); + let expected_limbs = expected.to_limbs(num_limbs, log_limb_size); + assert_eq!(result_limbs, expected_limbs); - // assert!(result_limbs.eq(&expected_limbs)); // TODO: leading limb is incorrect - assert!(result.eq(&a)); + let result = BigInt::from_limbs(&result_limbs, log_limb_size); + assert_eq!(result, expected); }