feat: integration with mopro-msm gpu acceleration

halo2_proofs/Cargo.toml

@@ -69,7 +69,7 @@ serde_derive = { version = "1", optional = true}
 bincode = { version = "1.3.3", default_features = false }
 serde = { version = "1.0.126", features = ["derive"] }
 instant = { version = "0.1" }
-
+mopro-msm = { git = "https://github.com/ElusAegis/metal-msm-gpu-acceleration.git", rev = "962e4d", optional = true, features = ["h2c"], default-features = false }
 
 # Developer tooling dependencies
 plotters = { version = "0.3.0", default-features = false, optional = true }
@@ -107,6 +107,8 @@ batch = ["rand_core/getrandom"]
 circuit-params = []
 counter = []
 icicle_gpu = ["icicle", "rustacuda"]
+macos = ["mopro-msm/macos"]
+ios = ["mopro-msm/ios"]
 mv-lookup = []
 cost-estimator = ["serde_derive"]
 derive_serde = ["halo2curves/derive_serde"]

halo2_proofs/src/arithmetic.rs

@@ -38,6 +38,25 @@ where
 ///
 /// This will use multithreading if beneficial.
 pub fn best_multiexp_cpu<C: CurveAffine>(coeffs: &[C::Scalar], bases: &[C]) -> C::Curve {
+    #[cfg(feature = "metal")]
+    if coeffs.len() >= 2_usize.pow(17) {
+        static PRINT_ONCE: Once = Once::new();
+
+        PRINT_ONCE.call_once(|| {
+            log::warn!(
+                "WARNING: Using Experimental Metal Acceleration for MSM. \
+                 Best performance improvements are observed with log row size >= 20. \
+                 Current log size: {}",
+                coeffs.len().ilog2()
+            );
+        });
+
+        use mopro_msm::metal::abstraction::limbs_conversion::h2c::{H2GAffine, H2G, H2Fr};
+        mopro_msm::metal::msm_best::<C, H2GAffine, H2G, H2Fr>(coeffs, bases)
+    } else {
+        msm_best(coeffs, bases)
+    }
+    #[cfg(not(feature = "metal"))]
     msm_best(coeffs, bases)
 }