Merge pull request kroma-network#372 from kroma-network/perf/avoid-in…

…-place-operator-where-ineffective perf: avoid in place operator where ineffective
ashjeong · Apr 4, 2024 · 3424346 · 3424346
2 parents c3c6276 + 26cd9fe
commit 3424346
Show file tree

Hide file tree

Showing 18 changed files with 158 additions and 296 deletions.
diff --git a/tachyon/device/tracking_allocator.cc b/tachyon/device/tracking_allocator.cc
@@ -57,7 +57,7 @@ void* TrackingAllocator::AllocateRaw(
     size_t allocated_bytes = allocator_->AllocatedSizeSlow(ptr);
     allocated_bytes = std::max(num_bytes, allocated_bytes);
     absl::MutexLock lock(&mu_);
-    next_allocation_id_ += 1;
+    ++next_allocation_id_;
     Chunk chunk = {num_bytes, allocated_bytes, next_allocation_id_};
     in_use_.emplace(std::make_pair(ptr, chunk));
     allocated_ += allocated_bytes;

diff --git a/tachyon/math/elliptic_curves/bn/bn_curve.h b/tachyon/math/elliptic_curves/bn/bn_curve.h
@@ -82,9 +82,8 @@ class BNCurve : public PairingFriendlyCurve<Config> {
     // Follows, e.g., Beuchat et al page 9, by computing result as follows:
     //   f^((q⁶ - 1) * (q² + 1)) = (conj(f) * f⁻¹)^(q² + 1)
 
-    // f1 = f.CyclotomicInverseInPlace() = f^(q⁶)
-    Fp12 f1 = f;
-    f1.CyclotomicInverseInPlace();
+    // f1 = f.CyclotomicInverse() = f^(q⁶)
+    Fp12 f1 = f.CyclotomicInverse();
 
     // f2 = f⁻¹
     Fp12 f2 = f.Inverse();

diff --git a/tachyon/math/elliptic_curves/msm/algorithms/cuzk/cuzk.h b/tachyon/math/elliptic_curves/msm/algorithms/cuzk/cuzk.h
@@ -300,7 +300,7 @@ class CUZK : public PippengerBase<AffinePoint<GpuCurve>>,
       unsigned int start = row_ptrs[i];
       unsigned int end = row_ptrs[i + 1];
       if ((end - start > gz) || (end - start < z)) continue;
-      n_other_total += 1;
+      ++n_other_total;
     }
     if (n_other_total == 0) {
       return true;
@@ -322,7 +322,7 @@ class CUZK : public PippengerBase<AffinePoint<GpuCurve>>,
       unsigned int n = (end - start) / z;
       row_ptrs2[n_other_total_idx + 1] = row_ptrs2[n_other_total_idx] + n;
 
-      n_other_total_idx += 1;
+      ++n_other_total_idx;
     }
 
     auto intermediate_datas =
@@ -358,7 +358,7 @@ class CUZK : public PippengerBase<AffinePoint<GpuCurve>>,
       }
 
       stream_id = (stream_id + 1) % grid_size_;
-      n_other_total_idx += 1;
+      ++n_other_total_idx;
     }
 
     for (unsigned int i = 0; i < grid_size_; ++i) {

diff --git a/tachyon/math/elliptic_curves/msm/kernels/bellman/bellman_msm_kernels.cu.h b/tachyon/math/elliptic_curves/msm/kernels/bellman/bellman_msm_kernels.cu.h
@@ -214,7 +214,7 @@ __global__ void AggregateBucketsKernel(
     auto base =
         Load<AffinePoint<Curve>, CacheOperator::kNone>(&bases[base_index]);
     if (sign) {
-      base = base.NegInPlace();
+      base.NegInPlace();
     }
     bucket = base.ToXYZZ();
   } else {

diff --git a/tachyon/math/elliptic_curves/short_weierstrass/affine_point_impl.h b/tachyon/math/elliptic_curves/short_weierstrass/affine_point_impl.h
@@ -19,8 +19,7 @@ constexpr ProjectivePoint<Curve> CLASS::DoubleProjective() const {
 
   // https://hyperelliptic.org/EFD/g1p/auto-shortw-projective.html#doubling-mdbl-2007-bl
   // XX = X1²
-  BaseField xx = x_;
-  xx.SquareInPlace();
+  BaseField xx = x_.Square();
 
   // w = a + 3 * XX
   BaseField w = xx;
@@ -32,42 +31,34 @@ constexpr ProjectivePoint<Curve> CLASS::DoubleProjective() const {
   }
 
   // Y1Y1 = Y1²
-  BaseField y1y1 = y_;
-  y1y1.SquareInPlace();
+  BaseField y1y1 = y_.Square();
 
   // R = 2 * Y1Y1
-  BaseField r = std::move(y1y1);
-  r.DoubleInPlace();
+  BaseField r = y1y1.Double();
 
   // sss = 4 * Y1 * R
-  BaseField sss = y_;
-  sss *= r;
+  BaseField sss = y_ * r;
   sss.DoubleInPlace().DoubleInPlace();
 
   // RR = R²
-  BaseField rr = r;
-  r.SquareInPlace();
+  BaseField rr = r.Square();
 
   // B = (X1 + R)² - XX - RR
-  BaseField b = std::move(r);
-  b += x_;
+  BaseField b = x_ + r;
   b.SquareInPlace();
   b -= xx;
   b -= rr;
 
   // h = w² - 2 * B
-  BaseField h = w;
-  h.SquareInPlace();
+  BaseField h = w.Square();
   h -= b.Double();
 
   // X3 = 2 * h * Y1
-  BaseField x = std::move(y_);
-  x *= h;
+  BaseField x = h * y_;
   x.DoubleInPlace();
 
   // Y3 = w * (B - h) - 2 * RR
-  BaseField y = std::move(b);
-  y -= h;
+  BaseField y = b - h;
   y *= w;
   y -= rr.Double();
 
@@ -85,24 +76,19 @@ constexpr PointXYZZ<Curve> CLASS::DoubleXYZZ() const {
 
   // https://hyperelliptic.org/EFD/g1p/auto-shortw-xyzz.html#doubling-mdbl-2008-s-1
   // U = 2 * Y1
-  BaseField u = y_;
-  u.DoubleInPlace();
+  BaseField u = y_.Double();
 
   // V = U²
-  BaseField v = u;
-  v.SquareInPlace();
+  BaseField v = u.Square();
 
   // W = U * V
-  BaseField w = u;
-  w *= v;
+  BaseField w = u * v;
 
   // S = X1 * V
-  BaseField s = x_;
-  s *= v;
+  BaseField s = x_ * v;
 
   // M = 3 * X1² + a
-  BaseField m = x_;
-  m.SquareInPlace();
+  BaseField m = x_.Square();
   m += m.Double();
   if constexpr (!Curve::Config::kAIsZero) {
     // TODO(chokobole): Implement constexpr version of Curve::Config::AddByA()
@@ -111,8 +97,7 @@ constexpr PointXYZZ<Curve> CLASS::DoubleXYZZ() const {
   }
 
   // X3 = M² - 2 * S
-  BaseField x = m;
-  x.SquareInPlace();
+  BaseField x = m.Square();
   x -= s.Double();
 
   // Y3 = M * (S - X3) - W * Y1