From 5dd16307583f3553333c6b85f712bbfd281dd432 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Sun, 4 Feb 2024 19:59:50 -0600 Subject: [PATCH 01/15] add icicle to gpu --- math/Cargo.toml | 16 ++ .../curves/bls12_377/curve.rs | 2 + math/src/gpu/icicle.rs | 251 ++++++++++++++++++ math/src/gpu/mod.rs | 2 + 4 files changed, 271 insertions(+) create mode 100644 math/src/gpu/icicle.rs diff --git a/math/Cargo.toml b/math/Cargo.toml index 02df065d4..fb5517359 100644 --- a/math/Cargo.toml +++ b/math/Cargo.toml @@ -24,6 +24,13 @@ objc = { version = "0.2.7", optional = true } # cuda cudarc = { version = "0.9.7", optional = true } +# Icicle integration +icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.0.0" } +icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } +icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } +icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } +icicle-bw6-761 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } + lambdaworks-gpu = { workspace = true, optional = true } @@ -55,6 +62,15 @@ metal = [ ] cuda = ["dep:cudarc", "dep:lambdaworks-gpu"] + +#icicle = ["dep:icicle-cuda-runtime", +#"dep:icicle-cuda-runtime", +#"dep:icicle-core", +#"dep:icicle-bn254", +#"dep:icicle-bls12-381", +#"icicle-bw6-761" +#] + [[bench]] name = "criterion_elliptic_curve" harness = false diff --git a/math/src/elliptic_curve/short_weierstrass/curves/bls12_377/curve.rs b/math/src/elliptic_curve/short_weierstrass/curves/bls12_377/curve.rs index c324b35de..a9b2958b1 100644 --- a/math/src/elliptic_curve/short_weierstrass/curves/bls12_377/curve.rs +++ b/math/src/elliptic_curve/short_weierstrass/curves/bls12_377/curve.rs @@ -5,6 +5,8 @@ use crate::{ elliptic_curve::short_weierstrass::traits::IsShortWeierstrass, field::element::FieldElement, }; +pub type BLS12377FieldElement = FieldElement; + /// The description of the curve. #[derive(Clone, Debug)] pub struct BLS12377Curve; diff --git a/math/src/gpu/icicle.rs b/math/src/gpu/icicle.rs new file mode 100644 index 000000000..fa5072472 --- /dev/null +++ b/math/src/gpu/icicle.rs @@ -0,0 +1,251 @@ +use icicle_bls12_377::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; +use icicle_bls12_381::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; +use icicle_bn254::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; +use icicle_core::{ + field::Field, + msm, + traits::{FieldConfig, FieldImpl, GenerateRandom}, + Curve::{Affine, Curve, Projective}, + Field::{Field, FieldImpl, MontgomeryConvertibleField}, +}; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; + +use crate::{ + elliptic_curve::{ + short_weierstrass::{ + curves::{ + bls12_377::{ + curve::{BLS12377Curve, BLS12377FieldElement}, + field_extension::BLS12377PrimeField, + }, + bls12_381::{ + curve::{BLS12381Curve, BLS12381FieldElement, BLS12381TwistCurveFieldElement}, + twist::BLS12381TwistCurve, + }, + bn_254::{ + curve::{BN254Curve, BN254FieldElement, BN254TwistCurveFieldElement}, + twist::BN254TwistCurve, + }, + }, + point::ShortWeierstrassProjectivePoint, + }, + traits::IsEllipticCurve, + }, + errors::ByteConversionError, + field::{element::FieldElement, traits::IsField}, + traits::ByteConversion, +}; + +use core::fmt::Debug; + +/// Notes: +/// Lambdaworks supplies rust bindings generic over there internal Field and Coordinate types. +/// The best solution is to upstream a `LambdaConvertible` trait implementation that handles this conversion for us. +/// In the meantime conversions are for specific curves and field implemented as the Icicle's Field type is not abstracted +/// from the field configuration or number of underlying limbs used in its representation + +/// trait for Conversions of lambdaworks type -> Icicle type +/// NOTE: This may be removed with eliminating `LambdaConvertible` +pub trait ToIcicle: Clone + Debug { + type IcicleType; + + fn to_icicle(&self) -> Self::IcicleType; + fn from_icicle(icicle: Self::IcicleType) -> Result; +} + +impl ToIcicle for BLS12377FieldElement { + type IcicleType = icicle_bls12_377::curve::BaseField; + + fn to_icicle(&self) -> Self::IcicleType { + IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Self::from_bytes_le(icicle.to_repr().to_bytes_le()) + } +} + +impl ToIcicle for BLS12381FieldElement { + type IcicleType = icicle_bls12_381::curve::BaseField; + + fn to_icicle(&self) -> Self::IcicleType { + IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Self::from_bytes_le(icicle.to_repr().to_bytes_le()) + } +} + +impl ToIcicle for BLS12381TwistCurveFieldElement { + type IcicleType = icicle_bls12_381::curve::BaseField; + + fn to_icicle(&self) -> Self::IcicleType { + IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Self::from_bytes_le(icicle.to_repr().to_bytes_le()) + } +} + +impl ToIcicle for BN254FieldElement { + type IcicleType = icicle_bn254::curve::BaseField; + + fn to_icicle(&self) -> Self::IcicleType { + IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Self::from_bytes_le(icicle.to_repr().to_bytes_le()) + } +} + +impl ToIcicle for BN254TwistCurveFieldElement { + type IcicleType = icicle_bn254::curve::BaseField; + + fn to_icicle(&self) -> Self::IcicleType { + IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } +} + +impl ToIcicle for ShortWeierstrassProjectivePoint { + type IcicleType = icicle_bls12_377::curve::G1Projective; + + fn to_icicle(&self) -> Self::IcicleType { + Self::IcicleType { + x: self.x().to_icicle(), + y: self.y().to_icicle(), + z: self.z().to_icicle(), + } + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(icicle.x).unwrap(), + FieldElement::::from_icicle(icicle.y).unwrap(), + FieldElement::::from_icicle(icicle.z).unwrap(), + ])) + } +} + +impl ToIcicle for ShortWeierstrassProjectivePoint { + type IcicleType = icicle_bls12_3811::curve::G1Projective; + + fn to_icicle(&self) -> Self::IcicleType { + Self::IcicleType { + x: self.x().to_icicle(), + y: self.y().to_icicle(), + z: self.z().to_icicle(), + } + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(icicle.x).unwrap(), + FieldElement::::from_icicle(icicle.y).unwrap(), + FieldElement::::from_icicle(icicle.z).unwrap(), + ])) + } +} + +impl ToIcicle for ShortWeierstrassProjectivePoint { + type IcicleType = icicle_bls12_381::curve::G2Projective; + + fn to_icicle(&self) -> Self::IcicleType { + Self::IcicleType { + x: self.x().to_icicle(), + y: self.y().to_icicle(), + z: self.z().to_icicle(), + } + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(icicle.x).unwrap(), + FieldElement::::from_icicle(icicle.y).unwrap(), + FieldElement::::from_icicle(icicle.z).unwrap(), + ])) + } +} + +impl ToIcicle for ShortWeierstrassProjectivePoint { + type IcicleType = icicle_bn254::curve::G1Projective; + + fn to_icicle(&self) -> Self::IcicleType { + Self::IcicleType { + x: self.x().to_icicle(), + y: self.y().to_icicle(), + z: self.z().to_icicle(), + } + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(icicle.x).unwrap(), + FieldElement::::from_icicle(icicle.y).unwrap(), + FieldElement::::from_icicle(icicle.z).unwrap(), + ])) + } +} + +impl ToIcicle for ShortWeierstrassProjectivePoint { + type IcicleType = icicle_bn254::curve::G2Projective; + + fn to_icicle(&self) -> Self::IcicleType { + Self::IcicleType { + x: self.x().to_icicle(), + y: self.y().to_icicle(), + z: self.z().to_icicle(), + } + } + + fn from_icicle(icicle: Self::IcicleType) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(icicle.x).unwrap(), + FieldElement::::from_icicle(icicle.y).unwrap(), + FieldElement::::from_icicle(icicle.z).unwrap(), + ])) + } +} + +/// Performs msm using Icicle GPU, intitiates, allocates, and configures all gpu operations +/// TODO: determining where this setup should occur is an open question +fn msm( + scalars: &[impl ToIcicle], + points: &[impl ToIcicle], +) -> ShortWeierstrassProjectivePoint { + let scalars = HostOrDeviceSlice::Host(&scalars.iter().map(to_icicle()).collect::>()); + let point = HostOrDeviceSlice::Host(&points.iter().map(to_icicle()).collect::>()); + let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + let mut cfg = msm::get_default_msm_config(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut msm_host_result = Vec::new(); + stream.synchronize().unwrap(); + msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); +} + +/// Performs ntt using Icicle GPU, intitiates, allocates, and configures all gpu operations +fn ntt(scalars: &[impl ToIcicle], points: &[impl ToIcicle]) -> FieldElement { + let point = HostOrDeviceSlice::Host(&points.iter().map(to_icicle()).collect::>()); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + let mut cfg = msm::get_default_msm_config(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut ntt_host_result = Vec::new(); + stream.synchronize().unwrap(); + ntt_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); + + let ctx = get_default_device_context(); +} diff --git a/math/src/gpu/mod.rs b/math/src/gpu/mod.rs index ee1867d80..a0269071d 100644 --- a/math/src/gpu/mod.rs +++ b/math/src/gpu/mod.rs @@ -6,3 +6,5 @@ If you were using the `--all-features` flag please read this crate's Cargo.toml" #[cfg(feature = "cuda")] pub mod cuda; + +pub mod icicle; From 27826b29382854912465bd96422bb67f21bcb628 Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 09:33:20 +0000 Subject: [PATCH 02/15] added icicle msm --- math/Cargo.toml | 21 +-- math/src/gpu/icicle/bls12_377.rs | 219 +++++++++++++++++++++++++++++++ math/src/gpu/icicle/bls12_381.rs | 219 +++++++++++++++++++++++++++++++ math/src/gpu/icicle/bn254.rs | 218 ++++++++++++++++++++++++++++++ math/src/gpu/icicle/mod.rs | 3 + 5 files changed, 665 insertions(+), 15 deletions(-) create mode 100644 math/src/gpu/icicle/bls12_377.rs create mode 100644 math/src/gpu/icicle/bls12_381.rs create mode 100644 math/src/gpu/icicle/bn254.rs create mode 100644 math/src/gpu/icicle/mod.rs diff --git a/math/Cargo.toml b/math/Cargo.toml index fb5517359..c2f5608a0 100644 --- a/math/Cargo.toml +++ b/math/Cargo.toml @@ -25,12 +25,11 @@ objc = { version = "0.2.7", optional = true } cudarc = { version = "0.9.7", optional = true } # Icicle integration -icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.0.0" } -icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } -icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } -icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } -icicle-bw6-761 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" } - +icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } +icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } +icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } +icicle-bls12-381 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } +icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } lambdaworks-gpu = { workspace = true, optional = true } @@ -52,6 +51,7 @@ lambdaworks-serde-binary = ["dep:serde", "alloc"] lambdaworks-serde-string = ["dep:serde", "dep:serde_json", "alloc"] proptest = ["dep:proptest"] winter_compatibility = ["winter-math", "miden-core"] +icicle = ["dep:icicle-cuda-runtime", "icicle-core", "icicle-bls12-377", "icicle-bls12-381", "icicle-bn254"] # gpu metal = [ @@ -62,15 +62,6 @@ metal = [ ] cuda = ["dep:cudarc", "dep:lambdaworks-gpu"] - -#icicle = ["dep:icicle-cuda-runtime", -#"dep:icicle-cuda-runtime", -#"dep:icicle-core", -#"dep:icicle-bn254", -#"dep:icicle-bls12-381", -#"icicle-bw6-761" -#] - [[bench]] name = "criterion_elliptic_curve" harness = false diff --git a/math/src/gpu/icicle/bls12_377.rs b/math/src/gpu/icicle/bls12_377.rs new file mode 100644 index 000000000..bfa0e68f7 --- /dev/null +++ b/math/src/gpu/icicle/bls12_377.rs @@ -0,0 +1,219 @@ +use icicle_bls12_377::curve; +use icicle_core::{error::IcicleError, msm, traits::FieldImpl}; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; + +use crate::{ + elliptic_curve::short_weierstrass::{ + curves::bls12_377::{ + curve::{BLS12377Curve, BLS12377FieldElement}, + field_extension::BLS12377PrimeField, + }, + point::ShortWeierstrassProjectivePoint, + }, + errors::ByteConversionError, + field::element::FieldElement, + traits::ByteConversion, +}; + +impl BLS12377FieldElement { + fn to_icicle(&self) -> curve::BaseField { + curve::BaseField::from_bytes_le(&self.to_bytes_le()) + } + + fn to_icicle_scalar(&self) -> curve::ScalarField { + curve::ScalarField::from_bytes_le(&self.to_bytes_le()) + } + + fn from_icicle(icicle: &curve::BaseField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } + + fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } +} + +impl ShortWeierstrassProjectivePoint { + fn to_icicle(&self) -> curve::G1Affine { + let s = self.to_affine(); + curve::G1Affine { + x: s.x().to_icicle(), + y: s.y().to_icicle(), + } + } + + + fn from_icicle(icicle: &curve::G1Projective) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::::from_icicle(&icicle.z).unwrap(), + ])) + } +} + +pub fn bls12_377_g1_msm( + scalars: &[BLS12377FieldElement], + points: &[ShortWeierstrassProjectivePoint], + config: Option, +) -> Result, IcicleError> { + let mut cfg = config.unwrap_or(msm::MSMConfig::default()); + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let points = HostOrDeviceSlice::Host( + points + .iter() + .map(|point| point.to_icicle()) + .collect::>(), + ); + let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut msm_host_result = vec![curve::G1Projective::zero(); 1]; + stream.synchronize().unwrap(); + msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle(&msm_host_result[0]).unwrap(); + Ok(res) +} + +/* +fn bls12_377_g1_ntt( + scalars: &[BLS12377FieldElement], + config: Option>, + order: u64, + inverse_fft: bool, +) -> Result, IcicleError> { + let size = scalars.len(); + let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); + let dir = if inverse_fft { + ntt::NTTDir::kInverse + } else { + ntt::NTTDir::kForward + }; + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + let root_of_unity = BLS12377PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); + curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + stream.synchronize().unwrap(); + let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; + ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); + stream.destroy().unwrap(); + let res = ntt_host_results + .as_slice() + .iter() + .map(|scalar| BLS12377FieldElement::from_icicle_scalar(&scalar).unwrap()) + .collect::>(); + Ok(res) +} +*/ + +#[cfg(test)] +mod test { + use super::*; + use crate::{ + elliptic_curve::{ + short_weierstrass::curves::bls12_377::curve::BLS12377FieldElement, traits::IsEllipticCurve, + }, + field::element::FieldElement, + msm::pippenger::msm, + }; + + impl ShortWeierstrassProjectivePoint { + + fn from_icicle_affine( + icicle: &curve::G1Affine, + ) -> Result, ByteConversionError> { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::one(), + ])) + } + } + + fn point_times_5() -> ShortWeierstrassProjectivePoint { + let x = BLS12377FieldElement::from_hex_unchecked( + "3c852d5aab73fbb51e57fbf5a0a8b5d6513ec922b2611b7547bfed74cba0dcdfc3ad2eac2733a4f55d198ec82b9964", + ); + let y = BLS12377FieldElement::from_hex_unchecked( + "a71425e68e55299c64d7eada9ae9c3fb87a9626b941d17128b64685fc07d0e635f3c3a512903b4e0a43e464045967b", + ); + BLS12377Curve::create_point_from_affine(x, y).unwrap() + } + + #[test] + fn to_from_icicle() { + // convert value of 5 to icicle and back again and that icicle 5 matches + let point = point_times_5(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn to_from_icicle_generator() { + // Convert generator and see that it matches + let point = BLS12377Curve::generator(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn icicle_g1_msm() { + const LEN: usize = 20; + let eight: BLS12377FieldElement = FieldElement::from(8); + let lambda_scalars = vec![eight; LEN]; + let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); + let expected = msm( + &lambda_scalars + .clone() + .into_iter() + .map(|x| x.representative()) + .collect::>(), + &lambda_points, + ) + .unwrap(); + let res = bls12_377_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); + assert_eq!(res, expected); + } + + /* + #[test] + fn icicle_1_ntt() { + const len: usize = 20; + let eight: BLS12377FieldElement = FieldElement::from(8); + let lambda_scalars = &vec![eight; len]; + let expected = Polynomial::evaluate_fft::( + &Polynomial::new(lambda_scalars), 1, None, + ).unwrap(); + let icicle_scalars = lambda_scalars.as_slice().iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(); + println!("expected {:?}", expected); + let res = bls12_381_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); + println!(); + println!("res {:?}", res); + assert_eq!(res, expected); + } + */ +} diff --git a/math/src/gpu/icicle/bls12_381.rs b/math/src/gpu/icicle/bls12_381.rs new file mode 100644 index 000000000..282ad95bb --- /dev/null +++ b/math/src/gpu/icicle/bls12_381.rs @@ -0,0 +1,219 @@ +use icicle_bls12_381::curve; +use icicle_core::{error::IcicleError, msm, traits::FieldImpl}; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; + +use crate::{ + elliptic_curve::short_weierstrass::{ + curves::bls12_381::{ + curve::{BLS12381Curve, BLS12381FieldElement}, + field_extension::BLS12381PrimeField, + }, + point::ShortWeierstrassProjectivePoint, + }, + errors::ByteConversionError, + field::element::FieldElement, + traits::ByteConversion, +}; + +impl BLS12381FieldElement { + fn to_icicle(&self) -> curve::BaseField { + curve::BaseField::from_bytes_le(&self.to_bytes_le()) + } + + fn to_icicle_scalar(&self) -> curve::ScalarField { + curve::ScalarField::from_bytes_le(&self.to_bytes_le()) + } + + fn from_icicle(icicle: &curve::BaseField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } + + fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } +} + +impl ShortWeierstrassProjectivePoint { + fn to_icicle(&self) -> curve::G1Affine { + let s = self.to_affine(); + curve::G1Affine { + x: s.x().to_icicle(), + y: s.y().to_icicle(), + } + } + + fn from_icicle(icicle: &curve::G1Projective) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::::from_icicle(&icicle.z).unwrap(), + ])) + } +} + +pub fn bls12_381_g1_msm( + scalars: &[BLS12381FieldElement], + points: &[ShortWeierstrassProjectivePoint], + config: Option, +) -> Result, IcicleError> { + let mut cfg = config.unwrap_or(msm::MSMConfig::default()); + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let points = HostOrDeviceSlice::Host( + points + .iter() + .map(|point| point.to_icicle()) + .collect::>(), + ); + let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut msm_host_result = vec![curve::G1Projective::zero(); 1]; + stream.synchronize().unwrap(); + msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle(&msm_host_result[0]).unwrap(); + Ok(res) +} + +/* +fn bls12_381_g1_ntt( + scalars: &[BLS12381FieldElement], + config: Option>, + order: u64, + inverse_fft: bool, +) -> Result, IcicleError> { + let size = scalars.len(); + let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); + let dir = if inverse_fft { + ntt::NTTDir::kInverse + } else { + ntt::NTTDir::kForward + }; + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + let root_of_unity = BLS12381PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); + curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + stream.synchronize().unwrap(); + let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; + ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); + stream.destroy().unwrap(); + let res = ntt_host_results + .as_slice() + .iter() + .map(|scalar| BLS12381FieldElement::from_icicle_scalar(&scalar).unwrap()) + .collect::>(); + Ok(res) +} +*/ + +#[cfg(test)] +mod test { + use super::*; + use crate::{ + elliptic_curve::{ + short_weierstrass::curves::bls12_381::curve::BLS12381FieldElement, traits::IsEllipticCurve, + }, + field::element::FieldElement, + msm::pippenger::msm, + }; + + + impl ShortWeierstrassProjectivePoint { + fn from_icicle_affine( + icicle: &curve::G1Affine, + ) -> Result, ByteConversionError> { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::one(), + ])) + } + } + + + fn point_times_5() -> ShortWeierstrassProjectivePoint { + let x = BLS12381FieldElement::from_hex_unchecked( + "32bcce7e71eb50384918e0c9809f73bde357027c6bf15092dd849aa0eac274d43af4c68a65fb2cda381734af5eecd5c", + ); + let y = BLS12381FieldElement::from_hex_unchecked( + "11e48467b19458aabe7c8a42dc4b67d7390fdf1e150534caadddc7e6f729d8890b68a5ea6885a21b555186452b954d88", + ); + BLS12381Curve::create_point_from_affine(x, y).unwrap() + } + + #[test] + fn to_from_icicle() { + // convert value of 5 to icicle and back again and that icicle 5 matches + let point = point_times_5(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn to_from_icicle_generator() { + // Convert generator and see that it matches + let point = BLS12381Curve::generator(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn icicle_g1_msm() { + const LEN: usize = 20; + let eight: BLS12381FieldElement = FieldElement::from(8); + let lambda_scalars = vec![eight; LEN]; + let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); + let expected = msm( + &lambda_scalars + .clone() + .into_iter() + .map(|x| x.representative()) + .collect::>(), + &lambda_points, + ) + .unwrap(); + let res = bls12_381_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); + assert_eq!(res, expected); + } + + /* + #[test] + fn icicle_1_ntt() { + const len: usize = 20; + let eight: BLS12381FieldElement = FieldElement::from(8); + let lambda_scalars = &vec![eight; len]; + let expected = Polynomial::evaluate_fft::( + &Polynomial::new(lambda_scalars), 1, None, + ).unwrap(); + let icicle_scalars = lambda_scalars.as_slice().iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(); + println!("expected {:?}", expected); + let res = bls12_381_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); + println!(); + println!("res {:?}", res); + assert_eq!(res, expected); + } + */ +} diff --git a/math/src/gpu/icicle/bn254.rs b/math/src/gpu/icicle/bn254.rs new file mode 100644 index 000000000..76ef1390c --- /dev/null +++ b/math/src/gpu/icicle/bn254.rs @@ -0,0 +1,218 @@ +use icicle_bn254::curve; +use icicle_core::{error::IcicleError, msm, traits::FieldImpl}; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; + +use crate::{ + elliptic_curve::short_weierstrass::{ + curves::bn_254::{ + curve::{BN254Curve, BN254FieldElement}, + field_extension::BN254PrimeField, + }, + point::ShortWeierstrassProjectivePoint, + }, + errors::ByteConversionError, + field::element::FieldElement, + traits::ByteConversion, +}; + +impl BN254FieldElement { + fn to_icicle(&self) -> curve::BaseField { + curve::BaseField::from_bytes_le(&self.to_bytes_le()) + } + + fn to_icicle_scalar(&self) -> curve::ScalarField { + curve::ScalarField::from_bytes_le(&self.to_bytes_le()) + } + + fn from_icicle(icicle: &curve::BaseField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } + + fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { + Self::from_bytes_le(&icicle.to_bytes_le()) + } +} + +impl ShortWeierstrassProjectivePoint { + fn to_icicle(&self) -> curve::G1Affine { + let s = self.to_affine(); + curve::G1Affine { + x: s.x().to_icicle(), + y: s.y().to_icicle(), + } + } + + fn from_icicle(icicle: &curve::G1Projective) -> Result { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::::from_icicle(&icicle.z).unwrap(), + ])) + } +} + +pub fn bn254_g1_msm( + scalars: &[BN254FieldElement], + points: &[ShortWeierstrassProjectivePoint], + config: Option, +) -> Result, IcicleError> { + let mut cfg = config.unwrap_or(msm::MSMConfig::default()); + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let points = HostOrDeviceSlice::Host( + points + .iter() + .map(|point| point.to_icicle()) + .collect::>(), + ); + let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut msm_host_result = vec![curve::G1Projective::zero(); 1]; + stream.synchronize().unwrap(); + msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle(&msm_host_result[0]).unwrap(); + Ok(res) +} + +/* +fn bn254_g1_ntt( + scalars: &[BN254FieldElement], + config: Option>, + order: u64, + inverse_fft: bool, +) -> Result, IcicleError> { + let size = scalars.len(); + let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); + let dir = if inverse_fft { + ntt::NTTDir::kInverse + } else { + ntt::NTTDir::kForward + }; + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(), + ); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + let root_of_unity = BN254PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); + curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + stream.synchronize().unwrap(); + let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; + ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); + stream.destroy().unwrap(); + let res = ntt_host_results + .as_slice() + .iter() + .map(|scalar| BN254FieldElement::from_icicle_scalar(&scalar).unwrap()) + .collect::>(); + Ok(res) +} +*/ + +#[cfg(test)] +mod test { + use super::*; + use crate::{ + elliptic_curve::{ + short_weierstrass::curves::bn_254::curve::BN254FieldElement, traits::IsEllipticCurve, + }, + field::element::FieldElement, + msm::pippenger::msm, + }; + + + impl ShortWeierstrassProjectivePoint { + fn from_icicle_affine( + icicle: &curve::G1Affine, + ) -> Result, ByteConversionError> { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::one(), + ])) + } + } + + fn point_times_5() -> ShortWeierstrassProjectivePoint { + let x = BN254FieldElement::from_hex_unchecked( + "16ab03b69dfb4f870b0143ebf6a71b7b2e4053ca7a4421d09a913b8b834bbfa3", + ); + let y = BN254FieldElement::from_hex_unchecked( + "2512347279ba1049ef97d4ec348d838f939d2b7623e88f4826643cf3889599b2", + ); + BN254Curve::create_point_from_affine(x, y).unwrap() + } + + #[test] + fn to_from_icicle() { + // convert value of 5 to icicle and back again and that icicle 5 matches + let point = point_times_5(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn to_from_icicle_generator() { + // Convert generator and see that it matches + let point = BN254Curve::generator(); + let icicle_point = point.to_icicle(); + let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn icicle_g1_msm() { + const LEN: usize = 20; + let eight: BN254FieldElement = FieldElement::from(8); + let lambda_scalars = vec![eight; LEN]; + let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); + let expected = msm( + &lambda_scalars + .clone() + .into_iter() + .map(|x| x.representative()) + .collect::>(), + &lambda_points, + ) + .unwrap(); + let res = bn254_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); + assert_eq!(res, expected); + } + + /* + #[test] + fn icicle_1_ntt() { + const len: usize = 20; + let eight: BN254FieldElement = FieldElement::from(8); + let lambda_scalars = &vec![eight; len]; + let expected = Polynomial::evaluate_fft::( + &Polynomial::new(lambda_scalars), 1, None, + ).unwrap(); + let icicle_scalars = lambda_scalars.as_slice().iter() + .map(|scalar| scalar.to_icicle_scalar()) + .collect::>(); + println!("expected {:?}", expected); + let res = bn254_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); + println!(); + println!("res {:?}", res); + assert_eq!(res, expected); + } + */ +} diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs new file mode 100644 index 000000000..3a0698b9f --- /dev/null +++ b/math/src/gpu/icicle/mod.rs @@ -0,0 +1,3 @@ +pub mod bls12_377; +pub mod bls12_381; +pub mod bn254; From f96d491a36e54348f2218475f2c124294544060e Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 15:26:31 +0000 Subject: [PATCH 03/15] remove ntt code for ec curves --- math/src/gpu/icicle/bls12_377.rs | 40 -------------------------------- math/src/gpu/icicle/bls12_381.rs | 40 -------------------------------- math/src/gpu/icicle/bn254.rs | 40 -------------------------------- 3 files changed, 120 deletions(-) diff --git a/math/src/gpu/icicle/bls12_377.rs b/math/src/gpu/icicle/bls12_377.rs index bfa0e68f7..159926a0e 100644 --- a/math/src/gpu/icicle/bls12_377.rs +++ b/math/src/gpu/icicle/bls12_377.rs @@ -84,46 +84,6 @@ pub fn bls12_377_g1_msm( Ok(res) } -/* -fn bls12_377_g1_ntt( - scalars: &[BLS12377FieldElement], - config: Option>, - order: u64, - inverse_fft: bool, -) -> Result, IcicleError> { - let size = scalars.len(); - let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); - let dir = if inverse_fft { - ntt::NTTDir::kInverse - } else { - ntt::NTTDir::kForward - }; - let scalars = HostOrDeviceSlice::Host( - scalars - .iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(), - ); - let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); - let stream = CudaStream::create().unwrap(); - cfg.ctx.stream = &stream; - cfg.is_async = true; - let root_of_unity = BLS12377PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); - curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); - ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); - stream.synchronize().unwrap(); - let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; - ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); - stream.destroy().unwrap(); - let res = ntt_host_results - .as_slice() - .iter() - .map(|scalar| BLS12377FieldElement::from_icicle_scalar(&scalar).unwrap()) - .collect::>(); - Ok(res) -} -*/ - #[cfg(test)] mod test { use super::*; diff --git a/math/src/gpu/icicle/bls12_381.rs b/math/src/gpu/icicle/bls12_381.rs index 282ad95bb..8d9f2da7a 100644 --- a/math/src/gpu/icicle/bls12_381.rs +++ b/math/src/gpu/icicle/bls12_381.rs @@ -83,46 +83,6 @@ pub fn bls12_381_g1_msm( Ok(res) } -/* -fn bls12_381_g1_ntt( - scalars: &[BLS12381FieldElement], - config: Option>, - order: u64, - inverse_fft: bool, -) -> Result, IcicleError> { - let size = scalars.len(); - let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); - let dir = if inverse_fft { - ntt::NTTDir::kInverse - } else { - ntt::NTTDir::kForward - }; - let scalars = HostOrDeviceSlice::Host( - scalars - .iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(), - ); - let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); - let stream = CudaStream::create().unwrap(); - cfg.ctx.stream = &stream; - cfg.is_async = true; - let root_of_unity = BLS12381PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); - curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); - ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); - stream.synchronize().unwrap(); - let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; - ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); - stream.destroy().unwrap(); - let res = ntt_host_results - .as_slice() - .iter() - .map(|scalar| BLS12381FieldElement::from_icicle_scalar(&scalar).unwrap()) - .collect::>(); - Ok(res) -} -*/ - #[cfg(test)] mod test { use super::*; diff --git a/math/src/gpu/icicle/bn254.rs b/math/src/gpu/icicle/bn254.rs index 76ef1390c..2fad6f67a 100644 --- a/math/src/gpu/icicle/bn254.rs +++ b/math/src/gpu/icicle/bn254.rs @@ -83,46 +83,6 @@ pub fn bn254_g1_msm( Ok(res) } -/* -fn bn254_g1_ntt( - scalars: &[BN254FieldElement], - config: Option>, - order: u64, - inverse_fft: bool, -) -> Result, IcicleError> { - let size = scalars.len(); - let mut cfg = config.unwrap_or(ntt::NTTConfig::default()); - let dir = if inverse_fft { - ntt::NTTDir::kInverse - } else { - ntt::NTTDir::kForward - }; - let scalars = HostOrDeviceSlice::Host( - scalars - .iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(), - ); - let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); - let stream = CudaStream::create().unwrap(); - cfg.ctx.stream = &stream; - cfg.is_async = true; - let root_of_unity = BN254PrimeField::get_primitive_root_of_unity(order).unwrap().to_icicle_scalar(); - curve::ScalarCfg::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); - ntt::ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); - stream.synchronize().unwrap(); - let mut ntt_host_results = vec![curve::ScalarField::zero(); size]; - ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); - stream.destroy().unwrap(); - let res = ntt_host_results - .as_slice() - .iter() - .map(|scalar| BN254FieldElement::from_icicle_scalar(&scalar).unwrap()) - .collect::>(); - Ok(res) -} -*/ - #[cfg(test)] mod test { use super::*; From 32359a5ae697db586d62ed52c7724a16b1692a79 Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 15:27:50 +0000 Subject: [PATCH 04/15] feature gate --- math/src/gpu/mod.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/math/src/gpu/mod.rs b/math/src/gpu/mod.rs index a0269071d..c0492c052 100644 --- a/math/src/gpu/mod.rs +++ b/math/src/gpu/mod.rs @@ -7,4 +7,5 @@ If you were using the `--all-features` flag please read this crate's Cargo.toml" #[cfg(feature = "cuda")] pub mod cuda; +#[cfg(all(feature = "icicle", feature = "alloc"))] pub mod icicle; From 3a0ea3574e168a841b5ee3ee542d88bd17298d8e Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 15:50:14 +0000 Subject: [PATCH 05/15] fmt and remove unneeded test --- math/src/gpu/icicle/bls12_377.rs | 35 ++++++++------------------------ math/src/gpu/icicle/bls12_381.rs | 35 ++++++++------------------------ math/src/gpu/icicle/bn254.rs | 21 ------------------- 3 files changed, 16 insertions(+), 75 deletions(-) diff --git a/math/src/gpu/icicle/bls12_377.rs b/math/src/gpu/icicle/bls12_377.rs index 159926a0e..5741fff12 100644 --- a/math/src/gpu/icicle/bls12_377.rs +++ b/math/src/gpu/icicle/bls12_377.rs @@ -42,7 +42,6 @@ impl ShortWeierstrassProjectivePoint { } } - fn from_icicle(icicle: &curve::G1Projective) -> Result { Ok(Self::new([ FieldElement::::from_icicle(&icicle.x).unwrap(), @@ -89,14 +88,14 @@ mod test { use super::*; use crate::{ elliptic_curve::{ - short_weierstrass::curves::bls12_377::curve::BLS12377FieldElement, traits::IsEllipticCurve, + short_weierstrass::curves::bls12_377::curve::BLS12377FieldElement, + traits::IsEllipticCurve, }, field::element::FieldElement, msm::pippenger::msm, }; impl ShortWeierstrassProjectivePoint { - fn from_icicle_affine( icicle: &curve::G1Affine, ) -> Result, ByteConversionError> { @@ -123,8 +122,9 @@ mod test { // convert value of 5 to icicle and back again and that icicle 5 matches let point = point_times_5(); let icicle_point = point.to_icicle(); - let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) - .unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); assert_eq!(point, res) } @@ -133,8 +133,9 @@ mod test { // Convert generator and see that it matches let point = BLS12377Curve::generator(); let icicle_point = point.to_icicle(); - let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) - .unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); assert_eq!(point, res) } @@ -156,24 +157,4 @@ mod test { let res = bls12_377_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); assert_eq!(res, expected); } - - /* - #[test] - fn icicle_1_ntt() { - const len: usize = 20; - let eight: BLS12377FieldElement = FieldElement::from(8); - let lambda_scalars = &vec![eight; len]; - let expected = Polynomial::evaluate_fft::( - &Polynomial::new(lambda_scalars), 1, None, - ).unwrap(); - let icicle_scalars = lambda_scalars.as_slice().iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(); - println!("expected {:?}", expected); - let res = bls12_381_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); - println!(); - println!("res {:?}", res); - assert_eq!(res, expected); - } - */ } diff --git a/math/src/gpu/icicle/bls12_381.rs b/math/src/gpu/icicle/bls12_381.rs index 8d9f2da7a..9f3031eaa 100644 --- a/math/src/gpu/icicle/bls12_381.rs +++ b/math/src/gpu/icicle/bls12_381.rs @@ -88,13 +88,13 @@ mod test { use super::*; use crate::{ elliptic_curve::{ - short_weierstrass::curves::bls12_381::curve::BLS12381FieldElement, traits::IsEllipticCurve, + short_weierstrass::curves::bls12_381::curve::BLS12381FieldElement, + traits::IsEllipticCurve, }, field::element::FieldElement, msm::pippenger::msm, }; - impl ShortWeierstrassProjectivePoint { fn from_icicle_affine( icicle: &curve::G1Affine, @@ -107,7 +107,6 @@ mod test { } } - fn point_times_5() -> ShortWeierstrassProjectivePoint { let x = BLS12381FieldElement::from_hex_unchecked( "32bcce7e71eb50384918e0c9809f73bde357027c6bf15092dd849aa0eac274d43af4c68a65fb2cda381734af5eecd5c", @@ -123,8 +122,9 @@ mod test { // convert value of 5 to icicle and back again and that icicle 5 matches let point = point_times_5(); let icicle_point = point.to_icicle(); - let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) - .unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); assert_eq!(point, res) } @@ -133,8 +133,9 @@ mod test { // Convert generator and see that it matches let point = BLS12381Curve::generator(); let icicle_point = point.to_icicle(); - let res = ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) - .unwrap(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); assert_eq!(point, res) } @@ -156,24 +157,4 @@ mod test { let res = bls12_381_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); assert_eq!(res, expected); } - - /* - #[test] - fn icicle_1_ntt() { - const len: usize = 20; - let eight: BLS12381FieldElement = FieldElement::from(8); - let lambda_scalars = &vec![eight; len]; - let expected = Polynomial::evaluate_fft::( - &Polynomial::new(lambda_scalars), 1, None, - ).unwrap(); - let icicle_scalars = lambda_scalars.as_slice().iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(); - println!("expected {:?}", expected); - let res = bls12_381_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); - println!(); - println!("res {:?}", res); - assert_eq!(res, expected); - } - */ } diff --git a/math/src/gpu/icicle/bn254.rs b/math/src/gpu/icicle/bn254.rs index 2fad6f67a..c21c1a48a 100644 --- a/math/src/gpu/icicle/bn254.rs +++ b/math/src/gpu/icicle/bn254.rs @@ -94,7 +94,6 @@ mod test { msm::pippenger::msm, }; - impl ShortWeierstrassProjectivePoint { fn from_icicle_affine( icicle: &curve::G1Affine, @@ -155,24 +154,4 @@ mod test { let res = bn254_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); assert_eq!(res, expected); } - - /* - #[test] - fn icicle_1_ntt() { - const len: usize = 20; - let eight: BN254FieldElement = FieldElement::from(8); - let lambda_scalars = &vec![eight; len]; - let expected = Polynomial::evaluate_fft::( - &Polynomial::new(lambda_scalars), 1, None, - ).unwrap(); - let icicle_scalars = lambda_scalars.as_slice().iter() - .map(|scalar| scalar.to_icicle_scalar()) - .collect::>(); - println!("expected {:?}", expected); - let res = bn254_g1_ntt(&lambda_scalars, None, 4u64, false).unwrap(); - println!(); - println!("res {:?}", res); - assert_eq!(res, expected); - } - */ } From 628b9467486bc1ffe5db5a1fcf4effda22fe7e74 Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 18:39:33 +0000 Subject: [PATCH 06/15] add benches --- math/Cargo.toml | 9 +++- math/benches/criterion_icicle.rs | 71 ++++++++++++++++++++++++++++++++ 2 files changed, 78 insertions(+), 2 deletions(-) create mode 100644 math/benches/criterion_icicle.rs diff --git a/math/Cargo.toml b/math/Cargo.toml index c2f5608a0..dffd2c345 100644 --- a/math/Cargo.toml +++ b/math/Cargo.toml @@ -34,7 +34,7 @@ icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4 lambdaworks-gpu = { workspace = true, optional = true } [dev-dependencies] -rand = { version = "0.8.5", default-features = false } +rand = { version = "0.8.5" } rand_chacha = "0.3.1" criterion = "0.5.1" const-random = "0.1.15" @@ -51,7 +51,7 @@ lambdaworks-serde-binary = ["dep:serde", "alloc"] lambdaworks-serde-string = ["dep:serde", "dep:serde_json", "alloc"] proptest = ["dep:proptest"] winter_compatibility = ["winter-math", "miden-core"] -icicle = ["dep:icicle-cuda-runtime", "icicle-core", "icicle-bls12-377", "icicle-bls12-381", "icicle-bn254"] +icicle = ["dep:icicle-cuda-runtime", "dep:icicle-core", "dep:icicle-bls12-377", "dep:icicle-bls12-381", "dep:icicle-bn254"] # gpu metal = [ @@ -91,6 +91,11 @@ name = "criterion_msm" harness = false required-features = ["parallel"] +[[bench]] +name = "criterion_icicle" +harness = false +required-features = ["icicle"] + [[bench]] name = "criterion_fft" harness = false diff --git a/math/benches/criterion_icicle.rs b/math/benches/criterion_icicle.rs new file mode 100644 index 000000000..07a271220 --- /dev/null +++ b/math/benches/criterion_icicle.rs @@ -0,0 +1,71 @@ +use criterion::{black_box, criterion_group, criterion_main, Criterion}; +use lambdaworks_math::{ + cyclic_group::IsGroup, + elliptic_curve::{ + short_weierstrass::curves::{ + bls12_377::curve::BLS12377Curve, bls12_381::curve::BLS12381Curve, + bn_254::curve::BN254Curve, + }, + traits::IsEllipticCurve, + }, + field::{element::FieldElement, traits::IsField}, +}; + +use lambdaworks_math::gpu::icicle::{ + bls12_377::bls12_377_g1_msm, bls12_381::bls12_381_g1_msm, bn254::bn254_g1_msm, +}; +use rand::{rngs::StdRng, Rng, SeedableRng}; + +pub fn generate_cs_and_points( + msm_size: usize, +) -> (Vec>, Vec) +where + ::BaseType: From, +{ + // We use a seeded rng so the benchmarks are reproducible. + let mut rng = StdRng::seed_from_u64(42); + + let g = C::generator(); + + let cs: Vec<_> = (0..msm_size) + .map(|_| FieldElement::::new(rng.gen::().into())) + .collect(); + + let points: Vec<_> = (0..msm_size) + .map(|_| g.operate_with_self(rng.gen::())) + .collect(); + + (cs, points) +} + +pub fn msm_benchmarks_with_size(c: &mut Criterion, msm_size: usize) { + let mut group = c.benchmark_group(format!("MSM benchmarks with size {msm_size}")); + + let (cs, points) = generate_cs_and_points::(msm_size); + group.bench_function("BLS12_381", |bench| { + bench.iter(|| black_box(bls12_381_g1_msm(&cs, &points, None))); + }); + + let (cs, points) = generate_cs_and_points::(msm_size); + group.bench_function("BLS12_377", |bench| { + bench.iter(|| black_box(bls12_377_g1_msm(&cs, &points, None))); + }); + + let (cs, points) = generate_cs_and_points::(msm_size); + group.bench_function("BN_254", |bench| { + bench.iter(|| black_box(bn254_g1_msm(&cs, &points, None))); + }); +} + +pub fn run_benchmarks(c: &mut Criterion) { + let exponents = 1..=18; + + for exp in exponents { + let msm_size = 1 << exp; + + msm_benchmarks_with_size(c, msm_size); + } +} + +criterion_group!(icicle_msm, run_benchmarks); +criterion_main!(icicle_msm); From dc49c5b99b48594f9a319d15d5a4b5829ec75563 Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 18:41:20 +0000 Subject: [PATCH 07/15] rm unneeded method --- math/src/gpu/icicle/bls12_377.rs | 4 ---- math/src/gpu/icicle/bls12_381.rs | 4 ---- math/src/gpu/icicle/bn254.rs | 4 ---- 3 files changed, 12 deletions(-) diff --git a/math/src/gpu/icicle/bls12_377.rs b/math/src/gpu/icicle/bls12_377.rs index 5741fff12..2381e11be 100644 --- a/math/src/gpu/icicle/bls12_377.rs +++ b/math/src/gpu/icicle/bls12_377.rs @@ -27,10 +27,6 @@ impl BLS12377FieldElement { fn from_icicle(icicle: &curve::BaseField) -> Result { Self::from_bytes_le(&icicle.to_bytes_le()) } - - fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { - Self::from_bytes_le(&icicle.to_bytes_le()) - } } impl ShortWeierstrassProjectivePoint { diff --git a/math/src/gpu/icicle/bls12_381.rs b/math/src/gpu/icicle/bls12_381.rs index 9f3031eaa..bf68f2304 100644 --- a/math/src/gpu/icicle/bls12_381.rs +++ b/math/src/gpu/icicle/bls12_381.rs @@ -27,10 +27,6 @@ impl BLS12381FieldElement { fn from_icicle(icicle: &curve::BaseField) -> Result { Self::from_bytes_le(&icicle.to_bytes_le()) } - - fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { - Self::from_bytes_le(&icicle.to_bytes_le()) - } } impl ShortWeierstrassProjectivePoint { diff --git a/math/src/gpu/icicle/bn254.rs b/math/src/gpu/icicle/bn254.rs index c21c1a48a..daf91a61f 100644 --- a/math/src/gpu/icicle/bn254.rs +++ b/math/src/gpu/icicle/bn254.rs @@ -27,10 +27,6 @@ impl BN254FieldElement { fn from_icicle(icicle: &curve::BaseField) -> Result { Self::from_bytes_le(&icicle.to_bytes_le()) } - - fn from_icicle_scalar(icicle: &curve::ScalarField) -> Result { - Self::from_bytes_le(&icicle.to_bytes_le()) - } } impl ShortWeierstrassProjectivePoint { From f0ab4a75ad98e123c0327c423f8e7d0f9fee2984 Mon Sep 17 00:00:00 2001 From: Ubuntu Date: Wed, 21 Feb 2024 18:43:50 +0000 Subject: [PATCH 08/15] fix feature flag --- math/src/gpu/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/math/src/gpu/mod.rs b/math/src/gpu/mod.rs index c0492c052..a19c098b2 100644 --- a/math/src/gpu/mod.rs +++ b/math/src/gpu/mod.rs @@ -7,5 +7,5 @@ If you were using the `--all-features` flag please read this crate's Cargo.toml" #[cfg(feature = "cuda")] pub mod cuda; -#[cfg(all(feature = "icicle", feature = "alloc"))] +#[cfg(feature = "icicle")] pub mod icicle; From 11fdcc5688f4f2ad8fbb6033d655359e79514168 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Fri, 8 Mar 2024 04:00:36 +0000 Subject: [PATCH 09/15] rewrite with generic msm fail on G's type --- crypto/src/commitments/kzg.rs | 17 +-- math/src/gpu/icicle.rs | 251 ------------------------------- math/src/gpu/icicle/bls12_381.rs | 6 +- math/src/gpu/icicle/mod.rs | 180 +++++++++++++++++++++- math/src/msm/naive.rs | 3 + math/src/msm/pippenger.rs | 31 +++- provers/groth16/src/prover.rs | 10 +- provers/groth16/src/verifier.rs | 5 +- 8 files changed, 215 insertions(+), 288 deletions(-) delete mode 100644 math/src/gpu/icicle.rs diff --git a/crypto/src/commitments/kzg.rs b/crypto/src/commitments/kzg.rs index 2ecbb483c..9d8587605 100644 --- a/crypto/src/commitments/kzg.rs +++ b/crypto/src/commitments/kzg.rs @@ -5,7 +5,7 @@ use lambdaworks_math::{ cyclic_group::IsGroup, elliptic_curve::traits::IsPairing, errors::DeserializationError, - field::{element::FieldElement, traits::IsPrimeField}, + field::{element::FieldElement, traits::{IsPrimeField, IsField}}, msm::pippenger::msm, polynomial::Polynomial, traits::{AsBytes, Deserializable}, @@ -136,12 +136,12 @@ where } #[derive(Clone)] -pub struct KateZaveruchaGoldberg { +pub struct KateZaveruchaGoldberg { srs: StructuredReferenceString, phantom: PhantomData, } -impl KateZaveruchaGoldberg { +impl KateZaveruchaGoldberg { pub fn new(srs: StructuredReferenceString) -> Self { Self { srs, @@ -150,20 +150,15 @@ impl KateZaveruchaGoldberg { } } -impl>, P: IsPairing> +impl> + IsPrimeField>, P: IsPairing> IsCommitmentScheme for KateZaveruchaGoldberg { type Commitment = P::G1Point; fn commit(&self, p: &Polynomial>) -> Self::Commitment { - let coefficients: Vec<_> = p - .coefficients - .iter() - .map(|coefficient| coefficient.representative()) - .collect(); msm( - &coefficients, - &self.srs.powers_main_group[..coefficients.len()], + &p.coefficients, + &self.srs.powers_main_group[..p.coefficients.len()], ) .expect("`points` is sliced by `cs`'s length") } diff --git a/math/src/gpu/icicle.rs b/math/src/gpu/icicle.rs deleted file mode 100644 index fa5072472..000000000 --- a/math/src/gpu/icicle.rs +++ /dev/null @@ -1,251 +0,0 @@ -use icicle_bls12_377::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; -use icicle_bls12_381::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; -use icicle_bn254::{CurveCfg, G1Projective, G2CurveCfg, G2Projective, ScalarCfg}; -use icicle_core::{ - field::Field, - msm, - traits::{FieldConfig, FieldImpl, GenerateRandom}, - Curve::{Affine, Curve, Projective}, - Field::{Field, FieldImpl, MontgomeryConvertibleField}, -}; -use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; - -use crate::{ - elliptic_curve::{ - short_weierstrass::{ - curves::{ - bls12_377::{ - curve::{BLS12377Curve, BLS12377FieldElement}, - field_extension::BLS12377PrimeField, - }, - bls12_381::{ - curve::{BLS12381Curve, BLS12381FieldElement, BLS12381TwistCurveFieldElement}, - twist::BLS12381TwistCurve, - }, - bn_254::{ - curve::{BN254Curve, BN254FieldElement, BN254TwistCurveFieldElement}, - twist::BN254TwistCurve, - }, - }, - point::ShortWeierstrassProjectivePoint, - }, - traits::IsEllipticCurve, - }, - errors::ByteConversionError, - field::{element::FieldElement, traits::IsField}, - traits::ByteConversion, -}; - -use core::fmt::Debug; - -/// Notes: -/// Lambdaworks supplies rust bindings generic over there internal Field and Coordinate types. -/// The best solution is to upstream a `LambdaConvertible` trait implementation that handles this conversion for us. -/// In the meantime conversions are for specific curves and field implemented as the Icicle's Field type is not abstracted -/// from the field configuration or number of underlying limbs used in its representation - -/// trait for Conversions of lambdaworks type -> Icicle type -/// NOTE: This may be removed with eliminating `LambdaConvertible` -pub trait ToIcicle: Clone + Debug { - type IcicleType; - - fn to_icicle(&self) -> Self::IcicleType; - fn from_icicle(icicle: Self::IcicleType) -> Result; -} - -impl ToIcicle for BLS12377FieldElement { - type IcicleType = icicle_bls12_377::curve::BaseField; - - fn to_icicle(&self) -> Self::IcicleType { - IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Self::from_bytes_le(icicle.to_repr().to_bytes_le()) - } -} - -impl ToIcicle for BLS12381FieldElement { - type IcicleType = icicle_bls12_381::curve::BaseField; - - fn to_icicle(&self) -> Self::IcicleType { - IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Self::from_bytes_le(icicle.to_repr().to_bytes_le()) - } -} - -impl ToIcicle for BLS12381TwistCurveFieldElement { - type IcicleType = icicle_bls12_381::curve::BaseField; - - fn to_icicle(&self) -> Self::IcicleType { - IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Self::from_bytes_le(icicle.to_repr().to_bytes_le()) - } -} - -impl ToIcicle for BN254FieldElement { - type IcicleType = icicle_bn254::curve::BaseField; - - fn to_icicle(&self) -> Self::IcicleType { - IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Self::from_bytes_le(icicle.to_repr().to_bytes_le()) - } -} - -impl ToIcicle for BN254TwistCurveFieldElement { - type IcicleType = icicle_bn254::curve::BaseField; - - fn to_icicle(&self) -> Self::IcicleType { - IcicleType::from_bytes_le(self.to_representative().to_bytes_le()) - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Self::from_bytes_le(&icicle.to_bytes_le()) - } -} - -impl ToIcicle for ShortWeierstrassProjectivePoint { - type IcicleType = icicle_bls12_377::curve::G1Projective; - - fn to_icicle(&self) -> Self::IcicleType { - Self::IcicleType { - x: self.x().to_icicle(), - y: self.y().to_icicle(), - z: self.z().to_icicle(), - } - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Ok(Self::new([ - FieldElement::::from_icicle(icicle.x).unwrap(), - FieldElement::::from_icicle(icicle.y).unwrap(), - FieldElement::::from_icicle(icicle.z).unwrap(), - ])) - } -} - -impl ToIcicle for ShortWeierstrassProjectivePoint { - type IcicleType = icicle_bls12_3811::curve::G1Projective; - - fn to_icicle(&self) -> Self::IcicleType { - Self::IcicleType { - x: self.x().to_icicle(), - y: self.y().to_icicle(), - z: self.z().to_icicle(), - } - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Ok(Self::new([ - FieldElement::::from_icicle(icicle.x).unwrap(), - FieldElement::::from_icicle(icicle.y).unwrap(), - FieldElement::::from_icicle(icicle.z).unwrap(), - ])) - } -} - -impl ToIcicle for ShortWeierstrassProjectivePoint { - type IcicleType = icicle_bls12_381::curve::G2Projective; - - fn to_icicle(&self) -> Self::IcicleType { - Self::IcicleType { - x: self.x().to_icicle(), - y: self.y().to_icicle(), - z: self.z().to_icicle(), - } - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Ok(Self::new([ - FieldElement::::from_icicle(icicle.x).unwrap(), - FieldElement::::from_icicle(icicle.y).unwrap(), - FieldElement::::from_icicle(icicle.z).unwrap(), - ])) - } -} - -impl ToIcicle for ShortWeierstrassProjectivePoint { - type IcicleType = icicle_bn254::curve::G1Projective; - - fn to_icicle(&self) -> Self::IcicleType { - Self::IcicleType { - x: self.x().to_icicle(), - y: self.y().to_icicle(), - z: self.z().to_icicle(), - } - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Ok(Self::new([ - FieldElement::::from_icicle(icicle.x).unwrap(), - FieldElement::::from_icicle(icicle.y).unwrap(), - FieldElement::::from_icicle(icicle.z).unwrap(), - ])) - } -} - -impl ToIcicle for ShortWeierstrassProjectivePoint { - type IcicleType = icicle_bn254::curve::G2Projective; - - fn to_icicle(&self) -> Self::IcicleType { - Self::IcicleType { - x: self.x().to_icicle(), - y: self.y().to_icicle(), - z: self.z().to_icicle(), - } - } - - fn from_icicle(icicle: Self::IcicleType) -> Result { - Ok(Self::new([ - FieldElement::::from_icicle(icicle.x).unwrap(), - FieldElement::::from_icicle(icicle.y).unwrap(), - FieldElement::::from_icicle(icicle.z).unwrap(), - ])) - } -} - -/// Performs msm using Icicle GPU, intitiates, allocates, and configures all gpu operations -/// TODO: determining where this setup should occur is an open question -fn msm( - scalars: &[impl ToIcicle], - points: &[impl ToIcicle], -) -> ShortWeierstrassProjectivePoint { - let scalars = HostOrDeviceSlice::Host(&scalars.iter().map(to_icicle()).collect::>()); - let point = HostOrDeviceSlice::Host(&points.iter().map(to_icicle()).collect::>()); - let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); - let stream = CudaStream::create().unwrap(); - let mut cfg = msm::get_default_msm_config(); - cfg.ctx.stream = &stream; - cfg.is_async = true; - msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); - let mut msm_host_result = Vec::new(); - stream.synchronize().unwrap(); - msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); - stream.destroy().unwrap(); -} - -/// Performs ntt using Icicle GPU, intitiates, allocates, and configures all gpu operations -fn ntt(scalars: &[impl ToIcicle], points: &[impl ToIcicle]) -> FieldElement { - let point = HostOrDeviceSlice::Host(&points.iter().map(to_icicle()).collect::>()); - let mut ntt_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); - let stream = CudaStream::create().unwrap(); - let mut cfg = msm::get_default_msm_config(); - cfg.ctx.stream = &stream; - cfg.is_async = true; - msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); - let mut ntt_host_result = Vec::new(); - stream.synchronize().unwrap(); - ntt_results.copy_to_host(&mut msm_host_result[..]).unwrap(); - stream.destroy().unwrap(); - - let ctx = get_default_device_context(); -} diff --git a/math/src/gpu/icicle/bls12_381.rs b/math/src/gpu/icicle/bls12_381.rs index bf68f2304..ce1f123b8 100644 --- a/math/src/gpu/icicle/bls12_381.rs +++ b/math/src/gpu/icicle/bls12_381.rs @@ -142,11 +142,7 @@ mod test { let lambda_scalars = vec![eight; LEN]; let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); let expected = msm( - &lambda_scalars - .clone() - .into_iter() - .map(|x| x.representative()) - .collect::>(), + &lambda_scalars, &lambda_points, ) .unwrap(); diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index 3a0698b9f..5a85dec91 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -1,3 +1,177 @@ -pub mod bls12_377; -pub mod bls12_381; -pub mod bn254; +//pub mod bls12_377; +//pub mod bls12_381; +//pub mod bn254; + +use icicle_bls12_381::curve::CurveCfg as IcicleBLS12381Curve; +use icicle_bls12_377::curve::CurveCfg as IcicleBLS12377Curve; +use icicle_bn254::curve::CurveCfg as IcicleBN254Curve; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; +use icicle_core::{error::IcicleError, msm, curve::{Curve, Affine, Projective}, traits::FieldImpl}; +use crate::{ + elliptic_curve::{short_weierstrass::{ + curves::{ + bls12_381::curve::BLS12381Curve, + bls12_377::curve::BLS12377Curve, + bn_254::curve::BN254Curve + }, + traits::IsShortWeierstrass, point::ShortWeierstrassProjectivePoint}, traits::IsEllipticCurve, + }, + field::{element::FieldElement, traits::IsField}, + unsigned_integer::element::UnsignedInteger, + cyclic_group::IsGroup, + errors::ByteConversionError, + traits::ByteConversion +}; + +use std::fmt::Debug; + +impl Icicle for BLS12381Curve {} +impl Icicle for BLS12377Curve {} +impl Icicle for BN254Curve {} + +pub trait Icicle +where + FieldElement: ByteConversion +{ + /// Used for searching this field's implementation in other languages, e.g in MSL + /// for executing parallel operations with the Metal API. + fn field_name() -> &'static str { + "" + } + + fn to_icicle_field(element: &FieldElement) -> I::BaseField { + I::BaseField::from_bytes_le(&element.to_bytes_le()) + } + + fn to_icicle_scalar(element: &FieldElement) -> I::ScalarField { + I::ScalarField::from_bytes_le(&element.to_bytes_le()) + } + + fn from_icicle_field(icicle: &I::BaseField) -> Result, ByteConversionError> { + FieldElement::::from_bytes_le(&icicle.to_bytes_le()) + } + + fn to_icicle_affine(point: &ShortWeierstrassProjectivePoint) -> Affine { + let s = ShortWeierstrassProjectivePoint::::to_affine(point); + Affine:: { + x: Self::to_icicle_field(s.x()), + y: Self::to_icicle_field(s.y()), + } + } + + fn from_icicle_projective(icicle: &Projective) -> Result, ByteConversionError> { + Ok(ShortWeierstrassProjectivePoint::::new([ + Self::from_icicle_field(&icicle.x).unwrap(), + Self::from_icicle_field(&icicle.y).unwrap(), + Self::from_icicle_field(&icicle.z).unwrap(), + ])) + } + +} + +pub fn icicle_msm>( + scalars: &[FieldElement], + points: &[ShortWeierstrassProjectivePoint] + ) -> Result, IcicleError> +where + C: Icicle, + FieldElement<::BaseField>: ByteConversion +{ + let mut cfg = msm::MSMConfig::default(); + let scalars = HostOrDeviceSlice::Host( + scalars + .iter() + .map(|scalar| C::to_icicle_scalar(&scalar)) + .collect::>(), + ); + let points = HostOrDeviceSlice::Host( + points + .iter() + .map(|point| C::to_icicle_affine(&point)) + .collect::>(), + ); + let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); + let mut msm_host_result = vec![Projective::::zero(); 1]; + stream.synchronize().unwrap(); + msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); + stream.destroy().unwrap(); + let res = + C::from_icicle_projective(&msm_host_result[0]).unwrap(); + Ok(res) +} + +#[cfg(test)] +mod test { + use super::*; + use crate::{ + elliptic_curve::{ + short_weierstrass::curves::bls12_381::curve::BLS12381FieldElement, + traits::IsEllipticCurve, + }, + field::element::FieldElement, + msm::pippenger::msm, + }; + + impl ShortWeierstrassProjectivePoint { + fn from_icicle_affine( + icicle: &curve::G1Affine, + ) -> Result, ByteConversionError> { + Ok(Self::new([ + FieldElement::::from_icicle(&icicle.x).unwrap(), + FieldElement::::from_icicle(&icicle.y).unwrap(), + FieldElement::one(), + ])) + } + } + + fn point_times_5() -> ShortWeierstrassProjectivePoint { + let x = BLS12381FieldElement::from_hex_unchecked( + "32bcce7e71eb50384918e0c9809f73bde357027c6bf15092dd849aa0eac274d43af4c68a65fb2cda381734af5eecd5c", + ); + let y = BLS12381FieldElement::from_hex_unchecked( + "11e48467b19458aabe7c8a42dc4b67d7390fdf1e150534caadddc7e6f729d8890b68a5ea6885a21b555186452b954d88", + ); + BLS12381Curve::create_point_from_affine(x, y).unwrap() + } + + #[test] + fn to_from_icicle() { + // convert value of 5 to icicle and back again and that icicle 5 matches + let point = point_times_5(); + let icicle_point = point.to_icicle(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn to_from_icicle_generator() { + // Convert generator and see that it matches + let point = BLS12381Curve::generator(); + let icicle_point = point.to_icicle(); + let res = + ShortWeierstrassProjectivePoint::::from_icicle_affine(&icicle_point) + .unwrap(); + assert_eq!(point, res) + } + + #[test] + fn icicle_g1_msm() { + const LEN: usize = 20; + let eight: BLS12381FieldElement = FieldElement::from(8); + let lambda_scalars = vec![eight; LEN]; + let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); + let expected = msm( + &lambda_scalars, + &lambda_points, + ) + .unwrap(); + let res = bls12_381_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); + assert_eq!(res, expected); + } +} diff --git a/math/src/msm/naive.rs b/math/src/msm/naive.rs index fec773073..985d08079 100644 --- a/math/src/msm/naive.rs +++ b/math/src/msm/naive.rs @@ -2,10 +2,13 @@ use core::fmt::Display; use crate::cyclic_group::IsGroup; use crate::unsigned_integer::traits::IsUnsignedInteger; +#[cfg(feature = "icicle")] +use icicle_core::error::IcicleError; #[derive(Debug)] pub enum MSMError { LengthMismatch(usize, usize), + Icicle(IcicleError) } impl Display for MSMError { diff --git a/math/src/msm/pippenger.rs b/math/src/msm/pippenger.rs index 1ccfce9c3..0f813e9a7 100644 --- a/math/src/msm/pippenger.rs +++ b/math/src/msm/pippenger.rs @@ -1,4 +1,4 @@ -use crate::{cyclic_group::IsGroup, unsigned_integer::element::UnsignedInteger}; +use crate::{cyclic_group::IsGroup, unsigned_integer::element::UnsignedInteger, field::{traits::IsField, element::FieldElement}, gpu::icicle::icicle_msm}; use super::naive::MSMError; @@ -15,8 +15,8 @@ use alloc::vec; /// If `points` and `cs` are empty, then `msm` returns the zero element of the group. /// /// Panics if `cs` and `points` have different lengths. -pub fn msm( - cs: &[UnsignedInteger], +pub fn msm>>( + cs: &[FieldElement], points: &[G], ) -> Result where @@ -26,9 +26,30 @@ where return Err(MSMError::LengthMismatch(cs.len(), points.len())); } - let window_size = optimum_window_size(cs.len()); + #[cfg(feature = "icicle")] + { + icicle_msm(cs, points).map_err(|e| MSMError::Icicle(e)) + /* + if !F::field_name().is_empty() { + icicle_msm(cs, points) + } else { + println!( + "Icicle msm failed for field {}. Program will fallback to CPU.", + core::any::type_name::() + ); + let window_size = optimum_window_size(cs.len()); + let cs = cs.into_iter().map(|cs| *cs.calue()).collect::>(); + Ok(msm_with(cs, points, window_size)) + } + */ + } - Ok(msm_with(cs, points, window_size)) + #[cfg(not(feature = "icicle"))] + { + let window_size = optimum_window_size(cs.len()); + let cs = cs.into_iter().map(|cs| *cs.value()).collect::>(); + Ok(msm_with(cs, points, window_size)) + } } fn optimum_window_size(data_length: usize) -> usize { diff --git a/provers/groth16/src/prover.rs b/provers/groth16/src/prover.rs index b0dbeb579..a007b3a67 100644 --- a/provers/groth16/src/prover.rs +++ b/provers/groth16/src/prover.rs @@ -67,15 +67,7 @@ pub struct Prover; impl Prover { pub fn prove(w: &[FrElement], qap: &QuadraticArithmeticProgram, pk: &ProvingKey) -> Proof { let h_coefficients = qap - .calculate_h_coefficients(w) - .iter() - .map(|elem| elem.representative()) - .collect::>(); - - let w = w - .iter() - .map(|elem| elem.representative()) - .collect::>(); + .calculate_h_coefficients(w); // Sample randomness for hiding let r = sample_fr_elem(); diff --git a/provers/groth16/src/verifier.rs b/provers/groth16/src/verifier.rs index b83c4b215..f3343c45f 100644 --- a/provers/groth16/src/verifier.rs +++ b/provers/groth16/src/verifier.rs @@ -7,10 +7,7 @@ use crate::setup::VerifyingKey; pub fn verify(vk: &VerifyingKey, proof: &Proof, pub_inputs: &[FrElement]) -> bool { // [γ^{-1} * (β*l(τ) + α*r(τ) + o(τ))]_1 let k_tau_assigned_verifier_g1 = msm( - &pub_inputs - .iter() - .map(|elem| elem.representative()) - .collect::>(), + &pub_inputs, &vk.verifier_k_tau_g1, ) .unwrap(); From cdce3ae0848a68337e8e5c18f69e985578c6a027 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Tue, 12 Mar 2024 04:26:25 +0000 Subject: [PATCH 10/15] msm implementation working --- crypto/src/commitments/kzg.rs | 18 ++- math/Cargo.toml | 2 +- math/src/gpu/icicle/mod.rs | 206 +++++++++++++++++++++++--------- math/src/msm/naive.rs | 5 +- math/src/msm/pippenger.rs | 28 +++-- provers/groth16/src/prover.rs | 9 +- provers/groth16/src/verifier.rs | 6 +- 7 files changed, 192 insertions(+), 82 deletions(-) diff --git a/crypto/src/commitments/kzg.rs b/crypto/src/commitments/kzg.rs index 9d8587605..c65d64d8c 100644 --- a/crypto/src/commitments/kzg.rs +++ b/crypto/src/commitments/kzg.rs @@ -1,11 +1,16 @@ use super::traits::IsCommitmentScheme; use alloc::{borrow::ToOwned, vec::Vec}; use core::{marker::PhantomData, mem}; +use lambdaworks_math::gpu::icicle::GpuMSMPoint; +use lambdaworks_math::traits::ByteConversion; use lambdaworks_math::{ cyclic_group::IsGroup, elliptic_curve::traits::IsPairing, errors::DeserializationError, - field::{element::FieldElement, traits::{IsPrimeField, IsField}}, + field::{ + element::FieldElement, + traits::{IsField, IsPrimeField}, + }, msm::pippenger::msm, polynomial::Polynomial, traits::{AsBytes, Deserializable}, @@ -150,8 +155,15 @@ impl KateZaveruchaGoldberg { } } -impl> + IsPrimeField>, P: IsPairing> - IsCommitmentScheme for KateZaveruchaGoldberg +impl< + const N: usize, + F: IsField> + + IsPrimeField>, + P: IsPairing, + > IsCommitmentScheme for KateZaveruchaGoldberg +where + FieldElement: ByteConversion, + P::G1Point: GpuMSMPoint, { type Commitment = P::G1Point; diff --git a/math/Cargo.toml b/math/Cargo.toml index dffd2c345..847afb37b 100644 --- a/math/Cargo.toml +++ b/math/Cargo.toml @@ -27,7 +27,7 @@ cudarc = { version = "0.9.7", optional = true } # Icicle integration icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } -icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } +icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true} icicle-bls12-381 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0", optional = true } diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index 5a85dec91..b44db162c 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -2,92 +2,193 @@ //pub mod bls12_381; //pub mod bn254; -use icicle_bls12_381::curve::CurveCfg as IcicleBLS12381Curve; -use icicle_bls12_377::curve::CurveCfg as IcicleBLS12377Curve; -use icicle_bn254::curve::CurveCfg as IcicleBN254Curve; -use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; -use icicle_core::{error::IcicleError, msm, curve::{Curve, Affine, Projective}, traits::FieldImpl}; use crate::{ - elliptic_curve::{short_weierstrass::{ - curves::{ - bls12_381::curve::BLS12381Curve, - bls12_377::curve::BLS12377Curve, - bn_254::curve::BN254Curve + cyclic_group::IsGroup, + elliptic_curve::{ + short_weierstrass::{ + curves::{ + bls12_377::curve::BLS12377Curve, + bls12_381::{curve::BLS12381Curve, twist::BLS12381TwistCurve}, + bn_254::curve::BN254Curve, + }, + point::ShortWeierstrassProjectivePoint, }, - traits::IsShortWeierstrass, point::ShortWeierstrassProjectivePoint}, traits::IsEllipticCurve, + traits::IsEllipticCurve, }, - field::{element::FieldElement, traits::IsField}, - unsigned_integer::element::UnsignedInteger, - cyclic_group::IsGroup, errors::ByteConversionError, - traits::ByteConversion + field::{element::FieldElement, traits::IsField}, + msm::naive::MSMError, + traits::ByteConversion, +}; +use icicle_bls12_377::curve::CurveCfg as IcicleBLS12377Curve; +use icicle_bls12_381::curve::CurveCfg as IcicleBLS12381Curve; +use icicle_bn254::curve::CurveCfg as IcicleBN254Curve; +use icicle_core::{ + curve::{Affine, Curve, Projective}, + msm, + traits::FieldImpl, }; +use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; use std::fmt::Debug; -impl Icicle for BLS12381Curve {} -impl Icicle for BLS12377Curve {} -impl Icicle for BN254Curve {} +impl GpuMSMPoint for ShortWeierstrassProjectivePoint { + type LambdaCurve = BLS12381Curve; + type GpuCurve = IcicleBLS12381Curve; -pub trait Icicle -where - FieldElement: ByteConversion -{ - /// Used for searching this field's implementation in other languages, e.g in MSL - /// for executing parallel operations with the Metal API. - fn field_name() -> &'static str { + fn curve_name() -> &'static str { + "BLS12381" + } + + fn to_icicle_affine(point: &Self) -> Affine { + let s = Self::to_affine(point); + Affine:: { + x: Self::to_icicle_field(s.x()), + y: Self::to_icicle_field(s.y()), + } + } + + fn from_icicle_projective( + icicle: &Projective, + ) -> Result { + Ok(Self::new([ + Self::from_icicle_field(&icicle.x).unwrap(), + Self::from_icicle_field(&icicle.y).unwrap(), + Self::from_icicle_field(&icicle.z).unwrap(), + ])) + } +} + +//NOTE THIS IS A PLACEHOLDER COMPILING ICICLE G2 TOOK TO LONG! +impl GpuMSMPoint for ShortWeierstrassProjectivePoint { + type LambdaCurve = BLS12381Curve; + type GpuCurve = IcicleBLS12381Curve; + + fn curve_name() -> &'static str { "" } - fn to_icicle_field(element: &FieldElement) -> I::BaseField { - I::BaseField::from_bytes_le(&element.to_bytes_le()) + fn to_icicle_affine(point: &Self) -> Affine { + let s = Self::to_affine(point); + Affine:: { + x: Self::to_icicle_field(s.x()), + y: Self::to_icicle_field(s.y()), + } } - fn to_icicle_scalar(element: &FieldElement) -> I::ScalarField { - I::ScalarField::from_bytes_le(&element.to_bytes_le()) + fn from_icicle_projective( + icicle: &Projective, + ) -> Result { + Ok(Self::new([ + Self::from_icicle_field(&icicle.x).unwrap(), + Self::from_icicle_field(&icicle.y).unwrap(), + Self::from_icicle_field(&icicle.z).unwrap(), + ])) + } +} + +impl GpuMSMPoint for ShortWeierstrassProjectivePoint { + type LambdaCurve = BLS12377Curve; + type GpuCurve = IcicleBLS12377Curve; + fn curve_name() -> &'static str { + "BLS12377" + } + + fn to_icicle_affine(point: &Self) -> Affine { + let s = Self::to_affine(point); + Affine:: { + x: Self::to_icicle_field(s.x()), + y: Self::to_icicle_field(s.y()), + } + } + + fn from_icicle_projective( + icicle: &Projective, + ) -> Result { + Ok(Self::new([ + Self::from_icicle_field(&icicle.x).unwrap(), + Self::from_icicle_field(&icicle.y).unwrap(), + Self::from_icicle_field(&icicle.z).unwrap(), + ])) } +} - fn from_icicle_field(icicle: &I::BaseField) -> Result, ByteConversionError> { - FieldElement::::from_bytes_le(&icicle.to_bytes_le()) +impl GpuMSMPoint for ShortWeierstrassProjectivePoint { + type LambdaCurve = BN254Curve; + type GpuCurve = IcicleBN254Curve; + fn curve_name() -> &'static str { + "BN254" } - fn to_icicle_affine(point: &ShortWeierstrassProjectivePoint) -> Affine { - let s = ShortWeierstrassProjectivePoint::::to_affine(point); - Affine:: { + fn to_icicle_affine(point: &Self) -> Affine { + let s = Self::to_affine(point); + Affine:: { x: Self::to_icicle_field(s.x()), y: Self::to_icicle_field(s.y()), } } - fn from_icicle_projective(icicle: &Projective) -> Result, ByteConversionError> { - Ok(ShortWeierstrassProjectivePoint::::new([ + fn from_icicle_projective( + icicle: &Projective, + ) -> Result { + Ok(Self::new([ Self::from_icicle_field(&icicle.x).unwrap(), Self::from_icicle_field(&icicle.y).unwrap(), Self::from_icicle_field(&icicle.z).unwrap(), ])) } +} +pub trait GpuMSMPoint: IsGroup { + type LambdaCurve: IsEllipticCurve + Clone + Debug; + type GpuCurve: Curve + msm::MSM; + //type FE: ByteConversion; + /// Used for searching this field's implementation in other languages, e.g in MSL + /// for executing parallel operations with the Metal API. + fn curve_name() -> &'static str { + "" + } + + fn to_icicle_field(element: &FE) -> ::BaseField { + ::BaseField::from_bytes_le(&element.to_bytes_le()) + } + + fn to_icicle_scalar( + element: &FE, + ) -> ::ScalarField { + ::ScalarField::from_bytes_le(&element.to_bytes_le()) + } + + fn from_icicle_field( + icicle: &::BaseField, + ) -> Result { + FE::from_bytes_le(&icicle.to_bytes_le()) + } + + fn to_icicle_affine(point: &Self) -> Affine; + + fn from_icicle_projective( + icicle: &Projective, + ) -> Result; } -pub fn icicle_msm>( - scalars: &[FieldElement], - points: &[ShortWeierstrassProjectivePoint] - ) -> Result, IcicleError> -where - C: Icicle, - FieldElement<::BaseField>: ByteConversion +pub fn icicle_msm( + cs: &[FieldElement], + points: &[G], +) -> Result +where + FieldElement: ByteConversion, { let mut cfg = msm::MSMConfig::default(); let scalars = HostOrDeviceSlice::Host( - scalars - .iter() - .map(|scalar| C::to_icicle_scalar(&scalar)) + cs.iter() + .map(|scalar| G::to_icicle_scalar(scalar)) .collect::>(), ); let points = HostOrDeviceSlice::Host( points .iter() - .map(|point| C::to_icicle_affine(&point)) + .map(|point| G::to_icicle_affine(point)) .collect::>(), ); let mut msm_results = HostOrDeviceSlice::cuda_malloc(1).unwrap(); @@ -95,12 +196,11 @@ where cfg.ctx.stream = &stream; cfg.is_async = true; msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap(); - let mut msm_host_result = vec![Projective::::zero(); 1]; + let mut msm_host_result = [Projective::::zero(); 1]; stream.synchronize().unwrap(); msm_results.copy_to_host(&mut msm_host_result[..]).unwrap(); stream.destroy().unwrap(); - let res = - C::from_icicle_projective(&msm_host_result[0]).unwrap(); + let res = G::from_icicle_projective(&msm_host_result[0]).unwrap(); Ok(res) } @@ -166,11 +266,7 @@ mod test { let eight: BLS12381FieldElement = FieldElement::from(8); let lambda_scalars = vec![eight; LEN]; let lambda_points = (0..LEN).map(|_| point_times_5()).collect::>(); - let expected = msm( - &lambda_scalars, - &lambda_points, - ) - .unwrap(); + let expected = msm(&lambda_scalars, &lambda_points).unwrap(); let res = bls12_381_g1_msm(&lambda_scalars, &lambda_points, None).unwrap(); assert_eq!(res, expected); } diff --git a/math/src/msm/naive.rs b/math/src/msm/naive.rs index 985d08079..50a1bad9b 100644 --- a/math/src/msm/naive.rs +++ b/math/src/msm/naive.rs @@ -8,13 +8,16 @@ use icicle_core::error::IcicleError; #[derive(Debug)] pub enum MSMError { LengthMismatch(usize, usize), - Icicle(IcicleError) + #[cfg(feature = "icicle")] + Icicle(IcicleError), } impl Display for MSMError { fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result { match self { MSMError::LengthMismatch(cs, points) => write!(f, "`cs` and `points` must be of the same length to compute `msm`. Got: {cs} and {points}"), + #[cfg(feature = "icicle")] + MSMError::Icicle(e) => write!(f, "Icicle GPU backend failure. {:?}", e), } } } diff --git a/math/src/msm/pippenger.rs b/math/src/msm/pippenger.rs index 0f813e9a7..99c7dc759 100644 --- a/math/src/msm/pippenger.rs +++ b/math/src/msm/pippenger.rs @@ -1,7 +1,13 @@ -use crate::{cyclic_group::IsGroup, unsigned_integer::element::UnsignedInteger, field::{traits::IsField, element::FieldElement}, gpu::icicle::icicle_msm}; +use crate::{ + cyclic_group::IsGroup, + field::{element::FieldElement, traits::IsField}, + unsigned_integer::element::UnsignedInteger, +}; use super::naive::MSMError; - +#[cfg(feature = "icicle")] +use crate::gpu::icicle::{icicle_msm, GpuMSMPoint}; +use crate::traits::ByteConversion; use alloc::vec; /// This function computes the multiscalar multiplication (MSM). @@ -15,12 +21,13 @@ use alloc::vec; /// If `points` and `cs` are empty, then `msm` returns the zero element of the group. /// /// Panics if `cs` and `points` have different lengths. -pub fn msm>>( +pub fn msm>, G>( cs: &[FieldElement], points: &[G], ) -> Result where - G: IsGroup, + G: IsGroup + GpuMSMPoint, + FieldElement: ByteConversion, { if cs.len() != points.len() { return Err(MSMError::LengthMismatch(cs.len(), points.len())); @@ -28,9 +35,7 @@ where #[cfg(feature = "icicle")] { - icicle_msm(cs, points).map_err(|e| MSMError::Icicle(e)) - /* - if !F::field_name().is_empty() { + if !G::curve_name().is_empty() { icicle_msm(cs, points) } else { println!( @@ -38,17 +43,16 @@ where core::any::type_name::() ); let window_size = optimum_window_size(cs.len()); - let cs = cs.into_iter().map(|cs| *cs.calue()).collect::>(); - Ok(msm_with(cs, points, window_size)) + let cs = cs.iter().map(|cs| *cs.value()).collect::>(); + Ok(msm_with(&cs, points, window_size)) } - */ } #[cfg(not(feature = "icicle"))] { let window_size = optimum_window_size(cs.len()); - let cs = cs.into_iter().map(|cs| *cs.value()).collect::>(); - Ok(msm_with(cs, points, window_size)) + let cs = cs.iter().map(|cs| *cs.value()).collect::>(); + Ok(msm_with(&cs, points, window_size)) } } diff --git a/provers/groth16/src/prover.rs b/provers/groth16/src/prover.rs index a007b3a67..559bfb5f5 100644 --- a/provers/groth16/src/prover.rs +++ b/provers/groth16/src/prover.rs @@ -66,21 +66,20 @@ impl Proof { pub struct Prover; impl Prover { pub fn prove(w: &[FrElement], qap: &QuadraticArithmeticProgram, pk: &ProvingKey) -> Proof { - let h_coefficients = qap - .calculate_h_coefficients(w); + let h_coefficients = qap.calculate_h_coefficients(w); // Sample randomness for hiding let r = sample_fr_elem(); let s = sample_fr_elem(); // [π_1]_1 - let pi1 = msm(&w, &pk.l_tau_g1) + let pi1 = msm(w, &pk.l_tau_g1) .unwrap() .operate_with(&pk.alpha_g1) .operate_with(&pk.delta_g1.operate_with_self(r.representative())); // [π_2]_2 - let pi2 = msm(&w, &pk.r_tau_g2) + let pi2 = msm(w, &pk.r_tau_g2) .unwrap() .operate_with(&pk.beta_g2) .operate_with(&pk.delta_g2.operate_with_self(s.representative())); @@ -100,7 +99,7 @@ impl Prover { .unwrap(); // [π_2]_1 - let pi2_g1 = msm(&w, &pk.r_tau_g1) + let pi2_g1 = msm(w, &pk.r_tau_g1) .unwrap() .operate_with(&pk.beta_g1) .operate_with(&pk.delta_g1.operate_with_self(s.representative())); diff --git a/provers/groth16/src/verifier.rs b/provers/groth16/src/verifier.rs index f3343c45f..0c0728475 100644 --- a/provers/groth16/src/verifier.rs +++ b/provers/groth16/src/verifier.rs @@ -6,11 +6,7 @@ use crate::setup::VerifyingKey; pub fn verify(vk: &VerifyingKey, proof: &Proof, pub_inputs: &[FrElement]) -> bool { // [γ^{-1} * (β*l(τ) + α*r(τ) + o(τ))]_1 - let k_tau_assigned_verifier_g1 = msm( - &pub_inputs, - &vk.verifier_k_tau_g1, - ) - .unwrap(); + let k_tau_assigned_verifier_g1 = msm(pub_inputs, &vk.verifier_k_tau_g1).unwrap(); Pairing::compute(&proof.pi3, &vk.delta_g2).unwrap() * vk.alpha_g1_times_beta_g2.clone() From c8a0baccc642276a62e3d1a7fa97cb20206315bf Mon Sep 17 00:00:00 2001 From: PatStiles Date: Tue, 12 Mar 2024 17:22:13 +0000 Subject: [PATCH 11/15] rm unneeded modules --- math/src/gpu/icicle/mod.rs | 4 ---- 1 file changed, 4 deletions(-) diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index b44db162c..7ae949446 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -1,7 +1,3 @@ -//pub mod bls12_377; -//pub mod bls12_381; -//pub mod bn254; - use crate::{ cyclic_group::IsGroup, elliptic_curve::{ From 277a672626db9cb54af3258a5fe9bfdfcd69e369 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Wed, 13 Mar 2024 19:54:28 +0000 Subject: [PATCH 12/15] nnt compiles --- math/src/fft/errors.rs | 7 ++ math/src/fft/polynomial.rs | 27 +++++++- math/src/gpu/icicle/mod.rs | 129 ++++++++++++++++++++++++++++++++++--- 3 files changed, 153 insertions(+), 10 deletions(-) diff --git a/math/src/fft/errors.rs b/math/src/fft/errors.rs index 4bbec91c2..e27e7d74d 100644 --- a/math/src/fft/errors.rs +++ b/math/src/fft/errors.rs @@ -8,6 +8,9 @@ use lambdaworks_gpu::metal::abstractions::errors::MetalError; #[cfg(feature = "cuda")] use lambdaworks_gpu::cuda::abstractions::errors::CudaError; +#[cfg(feature = "icicle")] +use icicle_core::error::IcicleError; + #[derive(Debug)] pub enum FFTError { RootOfUnityError(u64), @@ -17,6 +20,8 @@ pub enum FFTError { MetalError(MetalError), #[cfg(feature = "cuda")] CudaError(CudaError), + #[cfg(feature = "icicle")] + IcicleError(IcicleError), } impl Display for FFTError { @@ -37,6 +42,8 @@ impl Display for FFTError { FFTError::CudaError(_) => { write!(f, "A CUDA related error has ocurred") } + #[cfg(feature = "icicle")] + FFTError::IcicleError(e) => write!(f, "Icicle GPU backend failure. {:?}", e), } } } diff --git a/math/src/fft/polynomial.rs b/math/src/fft/polynomial.rs index a27811d31..70ad7cabc 100644 --- a/math/src/fft/polynomial.rs +++ b/math/src/fft/polynomial.rs @@ -7,6 +7,7 @@ use crate::{ traits::{IsFFTField, RootsConfig}, }, polynomial::Polynomial, + traits::ByteConversion, }; use alloc::{vec, vec::Vec}; @@ -14,6 +15,13 @@ use alloc::{vec, vec::Vec}; use crate::fft::gpu::cuda::polynomial::{evaluate_fft_cuda, interpolate_fft_cuda}; #[cfg(feature = "metal")] use crate::fft::gpu::metal::polynomial::{evaluate_fft_metal, interpolate_fft_metal}; +#[cfg(feature = "icicle")] +use crate::gpu::icicle::{evaluate_fft_icicle, IcicleFFT}; +#[cfg(feature = "icicle")] +use icicle_core::{ + ntt::NTT, + traits::FieldImpl, +}; use super::cpu::{ops, roots_of_unity}; @@ -26,7 +34,8 @@ impl Polynomial> { poly: &Polynomial>, blowup_factor: usize, domain_size: Option, - ) -> Result>, FFTError> { + ) -> Result>, FFTError> + { let domain_size = domain_size.unwrap_or(0); let len = core::cmp::max(poly.coeff_len(), domain_size).next_power_of_two() * blowup_factor; @@ -51,6 +60,21 @@ impl Polynomial> { } } + /* + #[cfg(feature = "icicle")] + { + if !F::field_name().is_empty() { + Ok(evaluate_fft_icicle::(&coeffs)?) + } else { + println!( + "GPU evaluation failed for field {}. Program will fallback to CPU.", + core::any::type_name::() + ); + evaluate_fft_cpu::(&coeffs) + } + } + */ + #[cfg(feature = "cuda")] { // TODO: support multiple fields with CUDA @@ -134,6 +158,7 @@ pub fn compose_fft( where F: IsFFTField + IsSubFieldOf, E: IsField, + FieldElement: ByteConversion { let poly_2_evaluations = Polynomial::evaluate_fft::(poly_2, 1, None).unwrap(); diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index 7ae949446..1eb113d53 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -4,24 +4,32 @@ use crate::{ short_weierstrass::{ curves::{ bls12_377::curve::BLS12377Curve, - bls12_381::{curve::BLS12381Curve, twist::BLS12381TwistCurve}, - bn_254::curve::BN254Curve, + bls12_381::{curve::BLS12381Curve, twist::BLS12381TwistCurve, field_extension::BLS12381PrimeField}, + bn_254::{curve::BN254Curve, field_extension::BN254PrimeField} }, point::ShortWeierstrassProjectivePoint, }, traits::IsEllipticCurve, }, errors::ByteConversionError, - field::{element::FieldElement, traits::IsField}, + field::{element::FieldElement, traits::{IsField, IsSubFieldOf, IsFFTField}}, msm::naive::MSMError, + fft::errors::FFTError, traits::ByteConversion, }; use icicle_bls12_377::curve::CurveCfg as IcicleBLS12377Curve; -use icicle_bls12_381::curve::CurveCfg as IcicleBLS12381Curve; -use icicle_bn254::curve::CurveCfg as IcicleBN254Curve; +use icicle_bls12_381::curve::{ + CurveCfg as IcicleBLS12381Curve, + ScalarCfg as IcicleBLS12381ScalarCfg +}; +use icicle_bn254::curve::{ + CurveCfg as IcicleBN254Curve, + ScalarCfg as IcicleBN254ScalarCfg +}; use icicle_core::{ curve::{Affine, Curve, Projective}, msm, + ntt::{ntt, NTT, NTTConfig, NTTDir}, traits::FieldImpl, }; use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; @@ -145,6 +153,12 @@ pub trait GpuMSMPoint: IsGroup { "" } + fn to_icicle_affine(point: &Self) -> Affine; + + fn from_icicle_projective( + icicle: &Projective, + ) -> Result; + fn to_icicle_field(element: &FE) -> ::BaseField { ::BaseField::from_bytes_le(&element.to_bytes_le()) } @@ -160,12 +174,32 @@ pub trait GpuMSMPoint: IsGroup { ) -> Result { FE::from_bytes_le(&icicle.to_bytes_le()) } +} - fn to_icicle_affine(point: &Self) -> Affine; +pub trait IcicleFFT: IsField +where + FieldElement: ByteConversion, + ::Config: NTT +{ + type ScalarField: FieldImpl; - fn from_icicle_projective( - icicle: &Projective, - ) -> Result; + fn to_icicle_scalar(element: &FieldElement) -> Self::ScalarField { + Self::ScalarField::from_bytes_le(&element.to_bytes_le()) + } + + fn from_icicle_scalar( + icicle: &Self::ScalarField, + ) -> Result, ByteConversionError> { + FieldElement::::from_bytes_le(&icicle.to_bytes_le()) + } +} + +impl IcicleFFT for BLS12381PrimeField { + type ScalarField = ::ScalarField; +} + +impl IcicleFFT for BN254PrimeField { + type ScalarField = ::ScalarField; } pub fn icicle_msm( @@ -200,6 +234,83 @@ where Ok(res) } +pub fn evaluate_fft_icicle( + coeffs: &Vec>, +) -> Result>, FFTError> +where + E: IsSubFieldOf, + FieldElement: ByteConversion, + <::ScalarField as FieldImpl>::Config: NTT<::ScalarField>, + E: IsField + IcicleFFT + IsFFTField, +{ + let size = coeffs.len(); + let mut cfg = NTTConfig::default(); + let order = coeffs.len() as u64; + let dir = NTTDir::kForward; + let scalars = HostOrDeviceSlice::Host( + coeffs + .iter() + .map(|scalar| E::to_icicle_scalar(&scalar)) + .collect::>(), + ); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + let root_of_unity = E::to_icicle_scalar(&E::get_primitive_root_of_unity(order).unwrap()); + ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + stream.synchronize().unwrap(); + let mut ntt_host_results = vec![E::ScalarField::zero(); size]; + ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); + stream.destroy().unwrap(); + let res = ntt_host_results + .as_slice() + .iter() + .map(|scalar| E::from_icicle_scalar(&scalar).unwrap()) + .collect::>(); + Ok(res) +} + +pub fn interpolate_fft_icicle( + coeffs: &Vec>, +) -> Result>, FFTError> +where + F: IsSubFieldOf, + FieldElement: ByteConversion, + <::ScalarField as FieldImpl>::Config: NTT<::ScalarField>, + E: IsField + IcicleFFT + IsFFTField, +{ + let size = coeffs.len(); + let mut cfg = NTTConfig::default(); + let order = coeffs.len() as u64; + let dir = NTTDir::kInverse; + let scalars = HostOrDeviceSlice::Host( + coeffs + .iter() + .map(|scalar| E::to_icicle_scalar(scalar)) + .collect::>(), + ); + let mut ntt_results = HostOrDeviceSlice::cuda_malloc(size).unwrap(); + let stream = CudaStream::create().unwrap(); + cfg.ctx.stream = &stream; + cfg.is_async = true; + let root_of_unity = E::to_icicle_scalar(&E::get_primitive_root_of_unity(order).unwrap()); + ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + stream.synchronize().unwrap(); + let mut ntt_host_results = vec![E::ScalarField::zero(); size]; + ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); + stream.destroy().unwrap(); + let res = ntt_host_results + .as_slice() + .iter() + .map(|scalar| E::from_icicle_scalar(&scalar).unwrap()) + .collect::>(); + Ok(res) +} + + #[cfg(test)] mod test { use super::*; From 4367b56a7d0bbdc19ba77186f4faebca05f24772 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Wed, 13 Mar 2024 19:59:08 +0000 Subject: [PATCH 13/15] nit --- math/src/gpu/icicle/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index 1eb113d53..a44c8c6e4 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -238,7 +238,7 @@ pub fn evaluate_fft_icicle( coeffs: &Vec>, ) -> Result>, FFTError> where - E: IsSubFieldOf, + F: IsSubFieldOf, FieldElement: ByteConversion, <::ScalarField as FieldImpl>::Config: NTT<::ScalarField>, E: IsField + IcicleFFT + IsFFTField, From c5fe38bcd19b0c74bf9cf16215e6f00350c06189 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Wed, 13 Mar 2024 20:29:58 +0000 Subject: [PATCH 14/15] some more fixes --- math/src/fft/polynomial.rs | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/math/src/fft/polynomial.rs b/math/src/fft/polynomial.rs index 70ad7cabc..3ebdfc2c1 100644 --- a/math/src/fft/polynomial.rs +++ b/math/src/fft/polynomial.rs @@ -25,7 +25,11 @@ use icicle_core::{ use super::cpu::{ops, roots_of_unity}; -impl Polynomial> { +impl Polynomial> +where + FieldElement: ByteConversion, + <::ScalarField as FieldImpl>::Config: NTT<::ScalarField> +{ /// Returns `N` evaluations of this polynomial using FFT over a domain in a subfield F of E (so the results /// are P(w^i), with w being a primitive root of unity). /// `N = max(self.coeff_len(), domain_size).next_power_of_two() * blowup_factor`. @@ -60,7 +64,6 @@ impl Polynomial> { } } - /* #[cfg(feature = "icicle")] { if !F::field_name().is_empty() { @@ -73,7 +76,6 @@ impl Polynomial> { evaluate_fft_cpu::(&coeffs) } } - */ #[cfg(feature = "cuda")] { @@ -85,7 +87,7 @@ impl Polynomial> { } } - #[cfg(all(not(feature = "metal"), not(feature = "cuda")))] + #[cfg(all(not(feature = "metal"), not(feature = "cuda"), not(feature = "icicle")))] { evaluate_fft_cpu::(&coeffs) } @@ -156,9 +158,10 @@ pub fn compose_fft( poly_2: &Polynomial>, ) -> Polynomial> where - F: IsFFTField + IsSubFieldOf, + F: IsFFTField + IsSubFieldOf + IcicleFFT, E: IsField, - FieldElement: ByteConversion + FieldElement: ByteConversion, + <::ScalarField as FieldImpl>::Config: NTT<::ScalarField> { let poly_2_evaluations = Polynomial::evaluate_fft::(poly_2, 1, None).unwrap(); From 86ca47353902a65484151dd9190b012e4eeb76d8 Mon Sep 17 00:00:00 2001 From: PatStiles Date: Fri, 5 Apr 2024 06:09:41 +0000 Subject: [PATCH 15/15] ntt compiles in all provers --- math/src/fft/polynomial.rs | 58 ++++++-------- math/src/gpu/icicle/mod.rs | 78 ++++++++++++------- math/src/msm/pippenger.rs | 3 +- provers/plonk/Cargo.toml | 1 + provers/plonk/src/prover.rs | 5 +- provers/plonk/src/setup.rs | 20 ++++- provers/plonk/src/verifier.rs | 8 +- provers/stark/Cargo.toml | 1 + provers/stark/src/constraints/evaluator.rs | 10 ++- provers/stark/src/constraints/transition.rs | 8 +- provers/stark/src/debug.rs | 10 ++- provers/stark/src/examples/dummy_air.rs | 8 +- .../src/examples/fibonacci_2_cols_shifted.rs | 12 ++- .../stark/src/examples/fibonacci_2_columns.rs | 15 +++- provers/stark/src/examples/fibonacci_rap.rs | 9 ++- provers/stark/src/examples/quadratic_air.rs | 12 ++- .../stark/src/examples/simple_fibonacci.rs | 12 ++- .../src/examples/simple_periodic_cols.rs | 12 ++- provers/stark/src/fri/mod.rs | 10 ++- provers/stark/src/prover.rs | 34 +++++--- provers/stark/src/trace.rs | 5 +- provers/stark/src/traits.rs | 14 +++- provers/stark/src/verifier.rs | 12 ++- 23 files changed, 239 insertions(+), 118 deletions(-) diff --git a/math/src/fft/polynomial.rs b/math/src/fft/polynomial.rs index 3ebdfc2c1..cbb341a3c 100644 --- a/math/src/fft/polynomial.rs +++ b/math/src/fft/polynomial.rs @@ -8,6 +8,7 @@ use crate::{ }, polynomial::Polynomial, traits::ByteConversion, + gpu::icicle::{IcicleFFT, GpuMSMPoint} }; use alloc::{vec, vec::Vec}; @@ -16,19 +17,13 @@ use crate::fft::gpu::cuda::polynomial::{evaluate_fft_cuda, interpolate_fft_cuda} #[cfg(feature = "metal")] use crate::fft::gpu::metal::polynomial::{evaluate_fft_metal, interpolate_fft_metal}; #[cfg(feature = "icicle")] -use crate::gpu::icicle::{evaluate_fft_icicle, IcicleFFT}; -#[cfg(feature = "icicle")] -use icicle_core::{ - ntt::NTT, - traits::FieldImpl, -}; +use crate::gpu::icicle::{evaluate_fft_icicle, interpolate_fft_icicle}; use super::cpu::{ops, roots_of_unity}; -impl Polynomial> -where +impl Polynomial> +where FieldElement: ByteConversion, - <::ScalarField as FieldImpl>::Config: NTT<::ScalarField> { /// Returns `N` evaluations of this polynomial using FFT over a domain in a subfield F of E (so the results /// are P(w^i), with w being a primitive root of unity). @@ -38,7 +33,9 @@ where poly: &Polynomial>, blowup_factor: usize, domain_size: Option, - ) -> Result>, FFTError> + ) -> Result>, FFTError> + where + FieldElement: ByteConversion, { let domain_size = domain_size.unwrap_or(0); let len = core::cmp::max(poly.coeff_len(), domain_size).next_power_of_two() * blowup_factor; @@ -77,17 +74,7 @@ where } } - #[cfg(feature = "cuda")] - { - // TODO: support multiple fields with CUDA - if F::field_name() == "stark256" { - Ok(evaluate_fft_cuda(&coeffs)?) - } else { - evaluate_fft_cpu::(&coeffs) - } - } - - #[cfg(all(not(feature = "metal"), not(feature = "cuda"), not(feature = "icicle")))] + #[cfg(all(not(feature = "metal"), not(feature = "icicle")))] { evaluate_fft_cpu::(&coeffs) } @@ -102,7 +89,10 @@ where blowup_factor: usize, domain_size: Option, offset: &FieldElement, - ) -> Result>, FFTError> { + ) -> Result>, FFTError> + where + FieldElement: ByteConversion, + { let scaled = poly.scale(offset); Polynomial::evaluate_fft::(&scaled, blowup_factor, domain_size) } @@ -126,16 +116,20 @@ where } } - #[cfg(feature = "cuda")] + #[cfg(feature = "icicle")] { if !F::field_name().is_empty() { - Ok(interpolate_fft_cuda(fft_evals)?) + Ok(interpolate_fft_icicle::(fft_evals)?) } else { - interpolate_fft_cpu::(fft_evals) + println!( + "GPU evaluation failed for field {}. Program will fallback to CPU.", + core::any::type_name::() + ); + interpolate_fft_cpu::(&fft_evals) } } - #[cfg(all(not(feature = "metal"), not(feature = "cuda")))] + #[cfg(all(not(feature = "metal"), not(feature = "icicle")))] { interpolate_fft_cpu::(fft_evals) } @@ -153,15 +147,13 @@ where } } -pub fn compose_fft( - poly_1: &Polynomial>, - poly_2: &Polynomial>, -) -> Polynomial> +pub fn compose_fft( + poly_1: &Polynomial>, + poly_2: &Polynomial>, +) -> Polynomial> where - F: IsFFTField + IsSubFieldOf + IcicleFFT, - E: IsField, + F: IsFFTField + IcicleFFT, FieldElement: ByteConversion, - <::ScalarField as FieldImpl>::Config: NTT<::ScalarField> { let poly_2_evaluations = Polynomial::evaluate_fft::(poly_2, 1, None).unwrap(); diff --git a/math/src/gpu/icicle/mod.rs b/math/src/gpu/icicle/mod.rs index a44c8c6e4..170bce439 100644 --- a/math/src/gpu/icicle/mod.rs +++ b/math/src/gpu/icicle/mod.rs @@ -4,32 +4,36 @@ use crate::{ short_weierstrass::{ curves::{ bls12_377::curve::BLS12377Curve, - bls12_381::{curve::BLS12381Curve, twist::BLS12381TwistCurve, field_extension::BLS12381PrimeField}, - bn_254::{curve::BN254Curve, field_extension::BN254PrimeField} + bls12_381::{ + curve::BLS12381Curve, default_types::FrField, + field_extension::BLS12381PrimeField, twist::BLS12381TwistCurve, + }, + bn_254::{curve::BN254Curve, field_extension::BN254PrimeField}, }, point::ShortWeierstrassProjectivePoint, }, traits::IsEllipticCurve, }, errors::ByteConversionError, - field::{element::FieldElement, traits::{IsField, IsSubFieldOf, IsFFTField}}, - msm::naive::MSMError, fft::errors::FFTError, + field::{ + element::FieldElement, + fields::fft_friendly::stark_252_prime_field::Stark252PrimeField, + traits::{IsFFTField, IsField, IsSubFieldOf}, + }, + msm::naive::MSMError, + polynomial::Polynomial, traits::ByteConversion, }; use icicle_bls12_377::curve::CurveCfg as IcicleBLS12377Curve; use icicle_bls12_381::curve::{ - CurveCfg as IcicleBLS12381Curve, - ScalarCfg as IcicleBLS12381ScalarCfg -}; -use icicle_bn254::curve::{ - CurveCfg as IcicleBN254Curve, - ScalarCfg as IcicleBN254ScalarCfg + CurveCfg as IcicleBLS12381Curve, ScalarCfg as IcicleBLS12381ScalarCfg, }; +use icicle_bn254::curve::{CurveCfg as IcicleBN254Curve, ScalarCfg as IcicleBN254ScalarCfg}; use icicle_core::{ curve::{Affine, Curve, Projective}, msm, - ntt::{ntt, NTT, NTTConfig, NTTDir}, + ntt::{NTTConfig, NTTDir, NTT}, traits::FieldImpl, }; use icicle_cuda_runtime::{memory::HostOrDeviceSlice, stream::CudaStream}; @@ -176,12 +180,12 @@ pub trait GpuMSMPoint: IsGroup { } } -pub trait IcicleFFT: IsField +pub trait IcicleFFT: IsField where FieldElement: ByteConversion, - ::Config: NTT { type ScalarField: FieldImpl; + type Config: NTT<::ScalarField>; fn to_icicle_scalar(element: &FieldElement) -> Self::ScalarField { Self::ScalarField::from_bytes_le(&element.to_bytes_le()) @@ -196,10 +200,23 @@ where impl IcicleFFT for BLS12381PrimeField { type ScalarField = ::ScalarField; + type Config = IcicleBLS12381ScalarCfg; +} + +impl IcicleFFT for FrField { + type ScalarField = ::ScalarField; + type Config = IcicleBLS12381ScalarCfg; +} + +// DUMMY IMPLEMENTATION OF STARK252 -> Fails when Icicle feature flag enabled +impl IcicleFFT for Stark252PrimeField { + type ScalarField = ::ScalarField; + type Config = IcicleBLS12381ScalarCfg; } impl IcicleFFT for BN254PrimeField { type ScalarField = ::ScalarField; + type Config = IcicleBN254ScalarCfg; } pub fn icicle_msm( @@ -236,12 +253,11 @@ where pub fn evaluate_fft_icicle( coeffs: &Vec>, -) -> Result>, FFTError> +) -> Result>, FFTError> where - F: IsSubFieldOf, + F: IsFFTField + IsSubFieldOf, FieldElement: ByteConversion, - <::ScalarField as FieldImpl>::Config: NTT<::ScalarField>, - E: IsField + IcicleFFT + IsFFTField, + E: IsField + IcicleFFT, { let size = coeffs.len(); let mut cfg = NTTConfig::default(); @@ -257,9 +273,11 @@ where let stream = CudaStream::create().unwrap(); cfg.ctx.stream = &stream; cfg.is_async = true; - let root_of_unity = E::to_icicle_scalar(&E::get_primitive_root_of_unity(order).unwrap()); - ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); - ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + let root_of_unity = E::to_icicle_scalar( + &(F::get_primitive_root_of_unity(order).unwrap() * FieldElement::::one()), + ); + ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ::Config::ntt_unchecked(&scalars, dir, &cfg, &mut ntt_results).unwrap(); stream.synchronize().unwrap(); let mut ntt_host_results = vec![E::ScalarField::zero(); size]; ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); @@ -273,13 +291,12 @@ where } pub fn interpolate_fft_icicle( - coeffs: &Vec>, -) -> Result>, FFTError> + coeffs: &[FieldElement], +) -> Result>, FFTError> where - F: IsSubFieldOf, + F: IsFFTField + IsSubFieldOf, FieldElement: ByteConversion, - <::ScalarField as FieldImpl>::Config: NTT<::ScalarField>, - E: IsField + IcicleFFT + IsFFTField, + E: IsField + IcicleFFT, { let size = coeffs.len(); let mut cfg = NTTConfig::default(); @@ -295,9 +312,11 @@ where let stream = CudaStream::create().unwrap(); cfg.ctx.stream = &stream; cfg.is_async = true; - let root_of_unity = E::to_icicle_scalar(&E::get_primitive_root_of_unity(order).unwrap()); - ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); - ntt(&scalars, dir, &cfg, &mut ntt_results).unwrap(); + let root_of_unity = E::to_icicle_scalar( + &(F::get_primitive_root_of_unity(order).unwrap() * FieldElement::::one()), + ); + ::Config::initialize_domain(root_of_unity, &cfg.ctx).unwrap(); + ::Config::ntt_unchecked(&scalars, dir, &cfg, &mut ntt_results).unwrap(); stream.synchronize().unwrap(); let mut ntt_host_results = vec![E::ScalarField::zero(); size]; ntt_results.copy_to_host(&mut ntt_host_results[..]).unwrap(); @@ -307,10 +326,9 @@ where .iter() .map(|scalar| E::from_icicle_scalar(&scalar).unwrap()) .collect::>(); - Ok(res) + Ok(Polynomial::new(&res)) } - #[cfg(test)] mod test { use super::*; diff --git a/math/src/msm/pippenger.rs b/math/src/msm/pippenger.rs index 99c7dc759..8e7c1ba59 100644 --- a/math/src/msm/pippenger.rs +++ b/math/src/msm/pippenger.rs @@ -2,11 +2,12 @@ use crate::{ cyclic_group::IsGroup, field::{element::FieldElement, traits::IsField}, unsigned_integer::element::UnsignedInteger, + gpu::icicle::GpuMSMPoint }; use super::naive::MSMError; #[cfg(feature = "icicle")] -use crate::gpu::icicle::{icicle_msm, GpuMSMPoint}; +use crate::gpu::icicle::icicle_msm; use crate::traits::ByteConversion; use alloc::vec; diff --git a/provers/plonk/Cargo.toml b/provers/plonk/Cargo.toml index 5f3784d48..e66804397 100644 --- a/provers/plonk/Cargo.toml +++ b/provers/plonk/Cargo.toml @@ -9,6 +9,7 @@ edition = "2021" lambdaworks-math.workspace = true lambdaworks-crypto.workspace = true serde = { version = "1.0", features = ["derive"] } +icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0" } serde_json = "1.0" sha3 = { version = "0.10", default-features = false } sha2 = { version = "0.10", default-features = false } diff --git a/provers/plonk/src/prover.rs b/provers/plonk/src/prover.rs index 3281c48db..aa2654e0c 100644 --- a/provers/plonk/src/prover.rs +++ b/provers/plonk/src/prover.rs @@ -15,6 +15,9 @@ use lambdaworks_math::{ }; use lambdaworks_math::{field::traits::IsField, traits::ByteConversion}; +//TODO: feature gate +use lambdaworks_math::gpu::icicle::IcicleFFT; + /// Plonk proof. /// The challenges are denoted /// Round 2: β,γ, @@ -279,7 +282,7 @@ struct Round5Result { impl Prover where - F: IsField + IsFFTField, + F: IsField + IsFFTField + IcicleFFT, CS: IsCommitmentScheme, FieldElement: ByteConversion, CS::Commitment: AsBytes, diff --git a/provers/plonk/src/setup.rs b/provers/plonk/src/setup.rs index 01f08588f..c8398415d 100644 --- a/provers/plonk/src/setup.rs +++ b/provers/plonk/src/setup.rs @@ -11,6 +11,9 @@ use lambdaworks_math::field::{element::FieldElement, traits::IsField}; use lambdaworks_math::polynomial::Polynomial; use lambdaworks_math::traits::{AsBytes, ByteConversion}; +//TODO: feature gate +use lambdaworks_math::gpu::icicle::IcicleFFT; + // TODO: implement getters pub struct Witness { pub a: Vec>, @@ -34,7 +37,10 @@ impl Witness { // TODO: implement getters #[derive(Clone)] -pub struct CommonPreprocessedInput { +pub struct CommonPreprocessedInput +where + FieldElement: ByteConversion, +{ pub n: usize, /// Number of constraints pub domain: Vec>, @@ -56,7 +62,10 @@ pub struct CommonPreprocessedInput { pub s3_lagrange: Vec>, } -impl CommonPreprocessedInput { +impl CommonPreprocessedInput +where + FieldElement: ByteConversion, +{ pub fn from_constraint_system( system: &ConstraintSystem, order_r_minus_1_root_unity: &FieldElement, @@ -113,10 +122,13 @@ pub struct VerificationKey { pub s3_1: G1Point, } -pub fn setup>( +pub fn setup>( common_input: &CommonPreprocessedInput, commitment_scheme: &CS, -) -> VerificationKey { +) -> VerificationKey +where + FieldElement: ByteConversion, +{ VerificationKey { qm_1: commitment_scheme.commit(&common_input.qm), ql_1: commitment_scheme.commit(&common_input.ql), diff --git a/provers/plonk/src/verifier.rs b/provers/plonk/src/verifier.rs index 8393148dd..cf3bff2e6 100644 --- a/provers/plonk/src/verifier.rs +++ b/provers/plonk/src/verifier.rs @@ -9,12 +9,18 @@ use std::marker::PhantomData; use crate::prover::Proof; use crate::setup::{new_strong_fiat_shamir_transcript, CommonPreprocessedInput, VerificationKey}; +//TODO: feature gate +use lambdaworks_math::gpu::icicle::IcicleFFT; + pub struct Verifier> { commitment_scheme: CS, phantom: PhantomData, } -impl> Verifier { +impl> Verifier +where + FieldElement: ByteConversion, +{ pub fn new(commitment_scheme: CS) -> Self { Self { commitment_scheme, diff --git a/provers/stark/Cargo.toml b/provers/stark/Cargo.toml index 2c181a729..a8245319c 100644 --- a/provers/stark/Cargo.toml +++ b/provers/stark/Cargo.toml @@ -11,6 +11,7 @@ crate-type = ["cdylib", "rlib"] [dependencies] lambdaworks-math = { workspace = true , features = ["std", "lambdaworks-serde-binary"] } lambdaworks-crypto = { workspace = true, features = ["std", "serde"] } +icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.4.0" } miden-core = { git="https://github.com/lambdaclass/miden-vm", optional=true} rand = "0.8.5" diff --git a/provers/stark/src/constraints/evaluator.rs b/provers/stark/src/constraints/evaluator.rs index 140a104ab..4a8032b79 100644 --- a/provers/stark/src/constraints/evaluator.rs +++ b/provers/stark/src/constraints/evaluator.rs @@ -8,7 +8,12 @@ use crate::{frame::Frame, prover::evaluate_polynomial_on_lde_domain}; use itertools::Itertools; #[cfg(all(debug_assertions, not(feature = "parallel")))] use lambdaworks_math::polynomial::Polynomial; -use lambdaworks_math::{fft::errors::FFTError, field::element::FieldElement, traits::AsBytes}; +use lambdaworks_math::{ + fft::errors::FFTError, + field::element::FieldElement, + gpu::icicle::IcicleFFT, + traits::{AsBytes, ByteConversion}, +}; #[cfg(feature = "parallel")] use rayon::{ iter::IndexedParallelIterator, @@ -39,9 +44,10 @@ impl ConstraintEvaluator { rap_challenges: &[FieldElement], ) -> Vec> where - FieldElement: AsBytes + Send + Sync, + FieldElement: AsBytes + Send + Sync + ByteConversion, FieldElement: AsBytes + Send + Sync, A: Send + Sync, + A::Field: IcicleFFT, { let boundary_constraints = &self.boundary_constraints; let number_of_b_constraints = boundary_constraints.constraints.len(); diff --git a/provers/stark/src/constraints/transition.rs b/provers/stark/src/constraints/transition.rs index 52fe0ba70..09ed5a63a 100644 --- a/provers/stark/src/constraints/transition.rs +++ b/provers/stark/src/constraints/transition.rs @@ -4,7 +4,9 @@ use crate::prover::evaluate_polynomial_on_lde_domain; use itertools::Itertools; use lambdaworks_math::field::element::FieldElement; use lambdaworks_math::field::traits::{IsFFTField, IsField, IsSubFieldOf}; +use lambdaworks_math::gpu::icicle::IcicleFFT; use lambdaworks_math::polynomial::Polynomial; +use lambdaworks_math::traits::ByteConversion; use num_integer::Integer; use std::ops::Div; /// TransitionConstraint represents the behaviour that a transition constraint @@ -105,7 +107,11 @@ where } /// Compute evaluations of the constraints zerofier over a LDE domain. - fn zerofier_evaluations_on_extended_domain(&self, domain: &Domain) -> Vec> { + fn zerofier_evaluations_on_extended_domain(&self, domain: &Domain) -> Vec> + where + F: IcicleFFT, + FieldElement: ByteConversion, + { let blowup_factor = domain.blowup_factor; let trace_length = domain.trace_roots_of_unity.len(); let trace_primitive_root = &domain.trace_primitive_root; diff --git a/provers/stark/src/debug.rs b/provers/stark/src/debug.rs index fc62257a4..c4e2aa80e 100644 --- a/provers/stark/src/debug.rs +++ b/provers/stark/src/debug.rs @@ -6,7 +6,9 @@ use lambdaworks_math::{ element::FieldElement, traits::{IsFFTField, IsField}, }, + gpu::icicle::IcicleFFT, polynomial::Polynomial, + traits::ByteConversion, }; use log::{error, info}; @@ -17,7 +19,13 @@ pub fn validate_trace( aux_trace_polys: &[Polynomial>], domain: &Domain, rap_challenges: &[FieldElement], -) -> bool { +) -> bool +where + ::Field: IcicleFFT, + FieldElement<::Field>: ByteConversion, + FieldElement<::FieldExtension>: ByteConversion, + ::FieldExtension: IsFFTField + IcicleFFT, +{ info!("Starting constraints validation over trace..."); let mut ret = true; diff --git a/provers/stark/src/examples/dummy_air.rs b/provers/stark/src/examples/dummy_air.rs index 1206ed153..8102addb9 100644 --- a/provers/stark/src/examples/dummy_air.rs +++ b/provers/stark/src/examples/dummy_air.rs @@ -15,6 +15,8 @@ use lambdaworks_math::field::{ element::FieldElement, fields::fft_friendly::stark_252_prime_field::Stark252PrimeField, traits::IsFFTField, }; +use lambdaworks_math::gpu::icicle::IcicleFFT; +use lambdaworks_math::traits::ByteConversion; type StarkField = Stark252PrimeField; @@ -32,7 +34,8 @@ impl FibConstraint { impl TransitionConstraint for FibConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -81,7 +84,8 @@ impl BitConstraint { impl TransitionConstraint for BitConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 2 diff --git a/provers/stark/src/examples/fibonacci_2_cols_shifted.rs b/provers/stark/src/examples/fibonacci_2_cols_shifted.rs index 6d3eb76d1..6315b36dd 100644 --- a/provers/stark/src/examples/fibonacci_2_cols_shifted.rs +++ b/provers/stark/src/examples/fibonacci_2_cols_shifted.rs @@ -11,7 +11,8 @@ use crate::{ }; use lambdaworks_math::{ field::{element::FieldElement, traits::IsFFTField}, - traits::AsBytes, + gpu::icicle::IcicleFFT, + traits::{AsBytes, ByteConversion}, }; use std::marker::PhantomData; @@ -30,7 +31,8 @@ impl ShiftedFibTransition1 { impl TransitionConstraint for ShiftedFibTransition1 where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -78,7 +80,8 @@ impl ShiftedFibTransition2 { impl TransitionConstraint for ShiftedFibTransition2 where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -149,7 +152,8 @@ where /// for all `i`. Also, `Col0_0` is constrained to be `1`. impl AIR for Fibonacci2ColsShifted where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, + FieldElement: ByteConversion, { type Field = F; type FieldExtension = F; diff --git a/provers/stark/src/examples/fibonacci_2_columns.rs b/provers/stark/src/examples/fibonacci_2_columns.rs index 0cd2a789a..9846cecb4 100644 --- a/provers/stark/src/examples/fibonacci_2_columns.rs +++ b/provers/stark/src/examples/fibonacci_2_columns.rs @@ -12,7 +12,11 @@ use crate::{ trace::TraceTable, traits::AIR, }; -use lambdaworks_math::field::{element::FieldElement, traits::IsFFTField}; +use lambdaworks_math::{ + field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, + traits::ByteConversion, +}; #[derive(Clone)] struct FibTransition1 { @@ -29,7 +33,8 @@ impl FibTransition1 { impl TransitionConstraint for FibTransition1 where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -79,7 +84,8 @@ impl FibTransition2 { impl TransitionConstraint for FibTransition2 where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -128,7 +134,8 @@ where /// stacked in row-major order. impl AIR for Fibonacci2ColsAIR where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, + FieldElement: ByteConversion, { type Field = F; type FieldExtension = F; diff --git a/provers/stark/src/examples/fibonacci_rap.rs b/provers/stark/src/examples/fibonacci_rap.rs index 18084f86f..d2a1cd7d3 100644 --- a/provers/stark/src/examples/fibonacci_rap.rs +++ b/provers/stark/src/examples/fibonacci_rap.rs @@ -14,6 +14,7 @@ use crate::{ use lambdaworks_crypto::fiat_shamir::is_transcript::IsTranscript; use lambdaworks_math::{ field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, helpers::resize_to_next_power_of_two, traits::ByteConversion, }; @@ -33,7 +34,8 @@ impl FibConstraint { impl TransitionConstraint for FibConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -85,7 +87,8 @@ impl PermutationConstraint { impl TransitionConstraint for PermutationConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 2 @@ -145,7 +148,7 @@ where impl AIR for FibonacciRAP where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, FieldElement: ByteConversion, { type Field = F; diff --git a/provers/stark/src/examples/quadratic_air.rs b/provers/stark/src/examples/quadratic_air.rs index 3d07d78e3..f64b35400 100644 --- a/provers/stark/src/examples/quadratic_air.rs +++ b/provers/stark/src/examples/quadratic_air.rs @@ -11,7 +11,11 @@ use crate::{ trace::TraceTable, traits::AIR, }; -use lambdaworks_math::field::{element::FieldElement, traits::IsFFTField}; +use lambdaworks_math::{ + field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, + traits::ByteConversion, +}; #[derive(Clone)] struct QuadraticConstraint { @@ -28,7 +32,8 @@ impl QuadraticConstraint { impl TransitionConstraint for QuadraticConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 2 @@ -81,7 +86,8 @@ where impl AIR for QuadraticAIR where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, + FieldElement: ByteConversion, { type Field = F; type FieldExtension = F; diff --git a/provers/stark/src/examples/simple_fibonacci.rs b/provers/stark/src/examples/simple_fibonacci.rs index 204aa938c..72b490ea0 100644 --- a/provers/stark/src/examples/simple_fibonacci.rs +++ b/provers/stark/src/examples/simple_fibonacci.rs @@ -9,7 +9,11 @@ use crate::{ trace::TraceTable, traits::AIR, }; -use lambdaworks_math::field::{element::FieldElement, traits::IsFFTField}; +use lambdaworks_math::{ + field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, + traits::ByteConversion, +}; use std::marker::PhantomData; #[derive(Clone)] @@ -27,7 +31,8 @@ impl FibConstraint { impl TransitionConstraint for FibConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -83,7 +88,8 @@ where impl AIR for FibonacciAIR where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, + FieldElement: ByteConversion, { type Field = F; type FieldExtension = F; diff --git a/provers/stark/src/examples/simple_periodic_cols.rs b/provers/stark/src/examples/simple_periodic_cols.rs index 369ff2dd0..23c8b05b5 100644 --- a/provers/stark/src/examples/simple_periodic_cols.rs +++ b/provers/stark/src/examples/simple_periodic_cols.rs @@ -11,7 +11,11 @@ use crate::{ trace::TraceTable, traits::AIR, }; -use lambdaworks_math::field::{element::FieldElement, traits::IsFFTField}; +use lambdaworks_math::{ + field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, + traits::ByteConversion, +}; pub struct PeriodicConstraint { phantom: PhantomData, @@ -31,7 +35,8 @@ impl Default for PeriodicConstraint { impl TransitionConstraint for PeriodicConstraint where - F: IsFFTField + Send + Sync, + F: IsFFTField + Send + Sync + IcicleFFT, + FieldElement: ByteConversion, { fn degree(&self) -> usize { 1 @@ -101,7 +106,8 @@ where impl AIR for SimplePeriodicAIR where - F: IsFFTField + Send + Sync + 'static, + F: IsFFTField + Send + Sync + 'static + IcicleFFT, + FieldElement: ByteConversion, { type Field = F; type FieldExtension = F; diff --git a/provers/stark/src/fri/mod.rs b/provers/stark/src/fri/mod.rs index a841a8aea..4598b92aa 100644 --- a/provers/stark/src/fri/mod.rs +++ b/provers/stark/src/fri/mod.rs @@ -10,7 +10,9 @@ use lambdaworks_math::{ }; pub use lambdaworks_math::{ field::{element::FieldElement, fields::u64_prime_field::U64PrimeField}, + gpu::icicle::IcicleFFT, polynomial::Polynomial, + traits::ByteConversion, }; use crate::config::{BatchedMerkleTree, BatchedMerkleTreeBackend}; @@ -19,7 +21,7 @@ use self::fri_commitment::FriLayer; use self::fri_decommit::FriDecommitment; use self::fri_functions::fold_polynomial; -pub fn commit_phase, E: IsField>( +pub fn commit_phase, E: IsField + IsFFTField + IcicleFFT>( number_layers: usize, p_0: Polynomial>, transcript: &mut impl IsTranscript, @@ -31,7 +33,7 @@ pub fn commit_phase, E: IsField>( ) where FieldElement: AsBytes + Sync + Send, - FieldElement: AsBytes + Sync + Send, + FieldElement: AsBytes + Sync + Send + ByteConversion, { let mut domain_size = domain_size; @@ -112,14 +114,14 @@ where } } -pub fn new_fri_layer, E: IsField>( +pub fn new_fri_layer, E: IsField + IsFFTField + IcicleFFT>( poly: &Polynomial>, coset_offset: &FieldElement, domain_size: usize, ) -> crate::fri::fri_commitment::FriLayer> where FieldElement: AsBytes + Sync + Send, - FieldElement: AsBytes + Sync + Send, + FieldElement: AsBytes + Sync + Send + ByteConversion, { let mut evaluation = Polynomial::evaluate_offset_fft(poly, 1, Some(domain_size), coset_offset).unwrap(); // TODO: return error diff --git a/provers/stark/src/prover.rs b/provers/stark/src/prover.rs index 7a5a4210e..746e9a41c 100644 --- a/provers/stark/src/prover.rs +++ b/provers/stark/src/prover.rs @@ -10,7 +10,9 @@ use lambdaworks_math::field::traits::{IsField, IsSubFieldOf}; use lambdaworks_math::traits::AsBytes; use lambdaworks_math::{ field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, polynomial::Polynomial, + traits::ByteConversion, }; use log::info; @@ -154,7 +156,8 @@ pub fn evaluate_polynomial_on_lde_domain( ) -> Result>, FFTError> where F: IsFFTField + IsSubFieldOf, - E: IsField, + E: IsField + IsFFTField + IcicleFFT, + FieldElement: ByteConversion, { let evaluations = Polynomial::evaluate_offset_fft(p, blowup_factor, Some(domain_size), offset)?; let step = evaluations.len() / (domain_size * blowup_factor); @@ -203,9 +206,9 @@ pub trait IsStarkProver { ) where FieldElement: AsBytes + Send + Sync, - FieldElement: AsBytes + Send + Sync, + FieldElement: AsBytes + Send + Sync + ByteConversion, FieldElement: AsBytes + Send + Sync, - E: IsSubFieldOf, + E: IsSubFieldOf + IcicleFFT + IsFFTField, A::Field: IsSubFieldOf, { // Interpolate columns of `trace`. @@ -243,8 +246,8 @@ pub trait IsStarkProver { ) -> Vec>> where FieldElement: Send + Sync, - FieldElement: Send + Sync, - E: IsSubFieldOf, + FieldElement: Send + Sync + ByteConversion, + E: IsSubFieldOf + IsFFTField + IcicleFFT, A::Field: IsSubFieldOf, { #[cfg(not(feature = "parallel"))] @@ -273,8 +276,10 @@ pub trait IsStarkProver { transcript: &mut impl IsTranscript, ) -> Result, ProvingError> where - FieldElement: AsBytes + Send + Sync, - FieldElement: AsBytes + Send + Sync, + A::Field: IcicleFFT, + A::FieldExtension: IsFFTField + IcicleFFT, + FieldElement: AsBytes + Send + Sync + ByteConversion, + FieldElement: AsBytes + Send + Sync + ByteConversion, { let (trace_polys, evaluations, main_merkle_tree, main_merkle_root) = Self::interpolate_and_commit::(main_trace, domain, transcript); @@ -358,8 +363,10 @@ pub trait IsStarkProver { ) -> Round2 where A: Send + Sync, - FieldElement: AsBytes + Send + Sync, - FieldElement: AsBytes + Send + Sync, + A::Field: IcicleFFT, + A::FieldExtension: IsFFTField + IcicleFFT, + FieldElement: AsBytes + Send + Sync + ByteConversion, + FieldElement: AsBytes + Send + Sync + ByteConversion, { // Compute the evaluations of the composition polynomial on the LDE domain. let evaluator = ConstraintEvaluator::new(air, &round_1_result.rap_challenges); @@ -464,8 +471,9 @@ pub trait IsStarkProver { transcript: &mut impl IsTranscript, ) -> Round4 where + A::FieldExtension: IsFFTField + IcicleFFT, FieldElement: AsBytes + Send + Sync, - FieldElement: AsBytes + Send + Sync, + FieldElement: AsBytes + Send + Sync + ByteConversion, { let coset_offset_u64 = air.context().proof_options.coset_offset; let coset_offset = FieldElement::::from(coset_offset_u64); @@ -797,8 +805,10 @@ pub trait IsStarkProver { ) -> Result, ProvingError> where A: Send + Sync, - FieldElement: AsBytes + Send + Sync, - FieldElement: AsBytes + Send + Sync, + A::Field: IcicleFFT, + A::FieldExtension: IsFFTField + IcicleFFT, + FieldElement: AsBytes + Send + Sync + ByteConversion, + FieldElement: AsBytes + Send + Sync + ByteConversion, { info!("Started proof generation..."); #[cfg(feature = "instruments")] diff --git a/provers/stark/src/trace.rs b/provers/stark/src/trace.rs index b9e3de0cf..9093fa78f 100644 --- a/provers/stark/src/trace.rs +++ b/provers/stark/src/trace.rs @@ -4,7 +4,9 @@ use lambdaworks_math::fft::errors::FFTError; use lambdaworks_math::field::traits::{IsField, IsSubFieldOf}; use lambdaworks_math::{ field::{element::FieldElement, traits::IsFFTField}, + gpu::icicle::IcicleFFT, polynomial::Polynomial, + traits::ByteConversion, }; #[cfg(feature = "parallel")] use rayon::prelude::{IntoParallelRefIterator, ParallelIterator}; @@ -137,7 +139,8 @@ impl TraceTable { pub fn compute_trace_polys(&self) -> Vec>> where S: IsFFTField + IsSubFieldOf, - FieldElement: Send + Sync, + FieldElement: Send + Sync + ByteConversion, + F: IcicleFFT + IsFFTField, { let columns = self.columns(); #[cfg(feature = "parallel")] diff --git a/provers/stark/src/traits.rs b/provers/stark/src/traits.rs index d33343bb2..f78c1e425 100644 --- a/provers/stark/src/traits.rs +++ b/provers/stark/src/traits.rs @@ -6,7 +6,9 @@ use lambdaworks_math::{ element::FieldElement, traits::{IsFFTField, IsField, IsSubFieldOf}, }, + gpu::icicle::IcicleFFT, polynomial::Polynomial, + traits::ByteConversion, }; use crate::{constraints::transition::TransitionConstraint, domain::Domain}; @@ -124,7 +126,11 @@ pub trait AIR { vec![] } - fn get_periodic_column_polynomials(&self) -> Vec>> { + fn get_periodic_column_polynomials(&self) -> Vec>> + where + ::Field: IcicleFFT, + FieldElement<::Field>: ByteConversion, + { let mut result = Vec::new(); for periodic_column in self.get_periodic_column_values() { let values: Vec<_> = periodic_column @@ -150,7 +156,11 @@ pub trait AIR { fn transition_zerofier_evaluations( &self, domain: &Domain, - ) -> Vec>> { + ) -> Vec>> + where + Self::Field: IcicleFFT, + FieldElement: ByteConversion, + { let mut evals = vec![Vec::new(); self.num_transition_constraints()]; let mut zerofier_groups: HashMap>> = diff --git a/provers/stark/src/verifier.rs b/provers/stark/src/verifier.rs index ea25dbe79..d4afa89d9 100644 --- a/provers/stark/src/verifier.rs +++ b/provers/stark/src/verifier.rs @@ -14,7 +14,8 @@ use lambdaworks_math::{ element::FieldElement, traits::{IsFFTField, IsField, IsSubFieldOf}, }, - traits::AsBytes, + gpu::icicle::IcicleFFT, + traits::{AsBytes, ByteConversion}, }; #[cfg(not(feature = "test_fiat_shamir"))] use log::error; @@ -216,7 +217,11 @@ pub trait IsStarkVerifier { proof: &StarkProof, domain: &Domain, challenges: &Challenges, - ) -> bool { + ) -> bool + where + A::Field: IcicleFFT, + FieldElement<::Field>: ByteConversion, + { let boundary_constraints = air.boundary_constraints(&challenges.rap_challenges); let trace_length = air.trace_length(); @@ -710,8 +715,9 @@ pub trait IsStarkVerifier { mut transcript: impl IsTranscript, ) -> bool where - FieldElement: AsBytes + Sync + Send, + FieldElement: AsBytes + Sync + Send + ByteConversion, FieldElement: AsBytes + Sync + Send, + A::Field: IcicleFFT, { // Verify there are enough queries if proof.query_list.len() < proof_options.fri_number_of_queries {