From 75192edcfb217c3d4ca09863cfc6d32f168ce5b3 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Wed, 19 Jun 2024 22:58:07 +0900 Subject: [PATCH 01/15] refac(math): remove unused headers --- tachyon/math/finite_fields/BUILD.bazel | 2 -- tachyon/math/finite_fields/cubic_extension_field.h | 1 - tachyon/math/finite_fields/quadratic_extension_field.h | 1 - 3 files changed, 4 deletions(-) diff --git a/tachyon/math/finite_fields/BUILD.bazel b/tachyon/math/finite_fields/BUILD.bazel index 0fe2c4a8b..01ab0b112 100644 --- a/tachyon/math/finite_fields/BUILD.bazel +++ b/tachyon/math/finite_fields/BUILD.bazel @@ -16,7 +16,6 @@ tachyon_cc_library( ":cyclotomic_multiplicative_subgroup", "//tachyon/base/buffer:copyable", "//tachyon/base/json", - "//tachyon/math/geometry:point3", "@com_google_absl//absl/strings", "@com_google_absl//absl/types:span", ], @@ -231,7 +230,6 @@ tachyon_cc_library( ":cyclotomic_multiplicative_subgroup", "//tachyon/base/buffer:copyable", "//tachyon/base/json", - "//tachyon/math/geometry:point2", "@com_google_absl//absl/strings", "@com_google_absl//absl/types:span", ], diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index 341a22fac..e4fb9483b 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -16,7 +16,6 @@ #include "tachyon/base/buffer/copyable.h" #include "tachyon/base/json/json.h" #include "tachyon/math/finite_fields/cyclotomic_multiplicative_subgroup.h" -#include "tachyon/math/geometry/point3.h" namespace tachyon { namespace math { diff --git a/tachyon/math/finite_fields/quadratic_extension_field.h b/tachyon/math/finite_fields/quadratic_extension_field.h index 869958034..81e1166fa 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field.h +++ b/tachyon/math/finite_fields/quadratic_extension_field.h @@ -16,7 +16,6 @@ #include "tachyon/base/buffer/copyable.h" #include "tachyon/base/json/json.h" #include "tachyon/math/finite_fields/cyclotomic_multiplicative_subgroup.h" -#include "tachyon/math/geometry/point2.h" namespace tachyon { namespace math { From 8bdc9dba9bab3c914f156d78e22f4294bfe1667b Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Wed, 19 Jun 2024 22:48:16 +0900 Subject: [PATCH 02/15] fix(math): fix bug in `CubicExtensionField::Norm()` See https://github.com/arkworks-rs/algebra/blob/a4dc41de8579d24fc982f13863fa35644afb34aa/ff/src/fields/models/cubic_extension.rs#L138. --- tachyon/math/finite_fields/cubic_extension_field.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index e4fb9483b..be84a76d7 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -94,7 +94,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { self_to_p *= (self_to_p2 * (*this)); // NOTE(chokobole): below CHECK() is not a device code. // See https://github.com/kroma-network/tachyon/issues/76 - CHECK(!(self_to_p.c1().IsZero() && self_to_p.c2().IsZero())); + CHECK(self_to_p.c1().IsZero() && self_to_p.c2().IsZero()); return self_to_p.c0(); } From f7cd8939358877d658dc30f286d166ce103b99c6 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Thu, 20 Jun 2024 15:53:09 +0900 Subject: [PATCH 03/15] fix(math): fix broken build in `CubicExtensionField::Norm()` --- tachyon/math/finite_fields/cubic_extension_field.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index be84a76d7..bb2a03eea 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -87,11 +87,11 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // NOTE(chokobole): This assumes that |BaseField::ExtensionDegree()| // never overflows even on 32 bit machine. size_t index_multiplier = size_t{BaseField::ExtensionDegree()}; - CubicExtensionField self_to_p = *this; + Derived self_to_p = static_cast(*this); self_to_p.FrobeniusMapInPlace(index_multiplier); - CubicExtensionField self_to_p2 = *this; + Derived self_to_p2 = static_cast(*this); self_to_p2.FrobeniusMapInPlace(2 * index_multiplier); - self_to_p *= (self_to_p2 * (*this)); + self_to_p *= (self_to_p2 * static_cast(*this)); // NOTE(chokobole): below CHECK() is not a device code. // See https://github.com/kroma-network/tachyon/issues/76 CHECK(self_to_p.c1().IsZero() && self_to_p.c2().IsZero()); From 5353bd03cd55e150713fb2bde940f35f9bfb83f9 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Thu, 20 Jun 2024 15:53:45 +0900 Subject: [PATCH 04/15] test(math): add `Norm()` test for `(Quadratic|Cubic)ExtensionField` --- .../math/finite_fields/cubic_extension_field_unittest.cc | 8 ++++++++ .../finite_fields/quadratic_extension_field_unittest.cc | 7 +++++++ 2 files changed, 15 insertions(+) diff --git a/tachyon/math/finite_fields/cubic_extension_field_unittest.cc b/tachyon/math/finite_fields/cubic_extension_field_unittest.cc index 4b35afe4d..d74817fbb 100644 --- a/tachyon/math/finite_fields/cubic_extension_field_unittest.cc +++ b/tachyon/math/finite_fields/cubic_extension_field_unittest.cc @@ -35,6 +35,14 @@ TEST_F(CubicExtensionFieldTest, Random) { EXPECT_TRUE(success); } +TEST_F(CubicExtensionFieldTest, Norm) { + constexpr static uint32_t kModulus = GF7::Config::kModulus; + GF7_3 r = GF7_3::Random(); + GF7_3 r_to_p = r.Pow(kModulus); + GF7_3 r_to_p2 = r_to_p.Pow(kModulus); + EXPECT_EQ(r.Norm(), (r * r_to_p * r_to_p2).c0()); +} + TEST_F(CubicExtensionFieldTest, EqualityOperators) { GF7_3 f(GF7(3), GF7(4), GF7(5)); GF7_3 f2(GF7(4), GF7(4), GF7(5)); diff --git a/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc b/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc index ee7be6c67..14671c1d6 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc +++ b/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc @@ -36,6 +36,13 @@ TEST_F(QuadraticExtensionFieldTest, Random) { EXPECT_TRUE(success); } +TEST_F(QuadraticExtensionFieldTest, Norm) { + constexpr static uint32_t kModulus = GF7::Config::kModulus; + GF7_2 r = GF7_2::Random(); + GF7_2 r_to_p = r.Pow(kModulus); + EXPECT_EQ(r.Norm(), (r * r_to_p).c0()); +} + TEST_F(QuadraticExtensionFieldTest, ConjugateInPlace) { GF7_2 f = GF7_2::Random(); GF7_2 f2 = f; From baacaa0bcea9648601f9db4f9bf35da6cfb990b7 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Wed, 19 Jun 2024 22:41:29 +0900 Subject: [PATCH 05/15] style(math): remove leading whitespace --- tachyon/math/finite_fields/cubic_extension_field.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index bb2a03eea..efac9f7f0 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -84,7 +84,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // Since Frobenius coefficients on the towered extensions are // indexed w.r.t. to |BasePrimeField|, we need to calculate the correct // index. - // NOTE(chokobole): This assumes that |BaseField::ExtensionDegree()| + // NOTE(chokobole): This assumes that |BaseField::ExtensionDegree()| // never overflows even on 32 bit machine. size_t index_multiplier = size_t{BaseField::ExtensionDegree()}; Derived self_to_p = static_cast(*this); From 07dee73250fbc51a6644bcbeb491cdc582ffeadc Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Wed, 19 Jun 2024 23:38:44 +0900 Subject: [PATCH 06/15] feat(math): add quadratic extension for `Goldilocks` and `KoalaBear` See https://github.com/Plonky3/Plonky3/blob/11ac8745b21295d5699f47089dbd927836f53fff/goldilocks/src/extension.rs#L6-L11 and https://github.com/Plonky3/Plonky3/blob/11ac8745b21295d5699f47089dbd927836f53fff/koala-bear/src/extension.rs#L8-L13. --- tachyon/math/finite_fields/goldilocks/BUILD.bazel | 11 +++++++++++ tachyon/math/finite_fields/koala_bear/BUILD.bazel | 11 +++++++++++ 2 files changed, 22 insertions(+) diff --git a/tachyon/math/finite_fields/goldilocks/BUILD.bazel b/tachyon/math/finite_fields/goldilocks/BUILD.bazel index 2829d55f6..308a82a11 100644 --- a/tachyon/math/finite_fields/goldilocks/BUILD.bazel +++ b/tachyon/math/finite_fields/goldilocks/BUILD.bazel @@ -1,6 +1,7 @@ load("@bazel_skylib//rules:common_settings.bzl", "string_flag") load("//bazel:tachyon.bzl", "if_has_avx512", "if_x86_64") load("//bazel:tachyon_cc.bzl", "tachyon_asm_prime_field_defines", "tachyon_cc_library", "tachyon_cc_unittest") +load("//tachyon/math/finite_fields/generator/ext_prime_field_generator:build_defs.bzl", "generate_fp2s") load( "//tachyon/math/finite_fields/generator/prime_field_generator:build_defs.bzl", "SMALL_SUBGROUP_ADICITY", @@ -40,6 +41,16 @@ generate_large_fft_prime_fields( use_montgomery = if_x86_64(False, True), ) +generate_fp2s( + name = "goldilocks2", + base_field = "Goldilocks", + base_field_hdr = "tachyon/math/finite_fields/goldilocks/goldilocks.h", + class_name = "Goldilocks2", + namespace = "tachyon::math", + non_residue = ["7"], + deps = [":goldilocks"], +) + tachyon_cc_library( name = "goldilocks_prime_field_x86_special", srcs = if_x86_64(["goldilocks_prime_field_x86_special.cc"]), diff --git a/tachyon/math/finite_fields/koala_bear/BUILD.bazel b/tachyon/math/finite_fields/koala_bear/BUILD.bazel index 032b30887..9a8a6a842 100644 --- a/tachyon/math/finite_fields/koala_bear/BUILD.bazel +++ b/tachyon/math/finite_fields/koala_bear/BUILD.bazel @@ -1,5 +1,6 @@ load("//bazel:tachyon.bzl", "if_aarch64", "if_has_avx512", "if_x86_64") load("//bazel:tachyon_cc.bzl", "tachyon_avx512_defines", "tachyon_cc_library") +load("//tachyon/math/finite_fields/generator/ext_prime_field_generator:build_defs.bzl", "generate_fp2s") load("//tachyon/math/finite_fields/generator/prime_field_generator:build_defs.bzl", "generate_prime_fields") package(default_visibility = ["//visibility:public"]) @@ -14,6 +15,16 @@ generate_prime_fields( use_montgomery = True, ) +generate_fp2s( + name = "koala_bear2", + base_field = "KoalaBear", + base_field_hdr = "tachyon/math/finite_fields/koala_bear/koala_bear.h", + class_name = "KoalaBear2", + namespace = "tachyon::math", + non_residue = ["3"], + deps = [":koala_bear"], +) + tachyon_cc_library( name = "packed_koala_bear", hdrs = ["packed_koala_bear.h"], From bd3c3b601967ca1352462c9b1b59f41e44714737 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Thu, 20 Jun 2024 00:54:55 +0900 Subject: [PATCH 07/15] refac(math): remove redundant type alias --- .../generator/ext_prime_field_generator/fq.h.tpl | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl b/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl index 6f7d24c10..b49ace189 100644 --- a/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl +++ b/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl @@ -11,15 +11,12 @@ class %{class}Config { using BaseField = _BaseField; using BasePrimeField = %{base_prime_field}; using FrobeniusCoefficient = %{frobenius_coefficient}; -%{if FrobeniusCoefficient2} - using FrobeniusCoefficient2 = %{frobenius_coefficient}; -%{endif FrobeniusCoefficient2} // TODO(chokobole): Make them constexpr. static BaseField kNonResidue; static FrobeniusCoefficient kFrobeniusCoeffs[%{frobenius_coeffs_size}]; %{if FrobeniusCoefficient2} - static FrobeniusCoefficient2 kFrobeniusCoeffs2[%{frobenius_coeffs_size}]; + static FrobeniusCoefficient kFrobeniusCoeffs2[%{frobenius_coeffs_size}]; %{endif FrobeniusCoefficient2} constexpr static bool kNonResidueIsMinusOne = %{non_residue_is_minus_one}; From 95ba7b058546ce8956e6eab41d61c1bc061d27d8 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Thu, 20 Jun 2024 13:06:28 +0900 Subject: [PATCH 08/15] fix(math): initialize variable in `constexpr` function --- tachyon/math/base/arithmetics.h | 2 +- tachyon/math/base/big_int.h | 24 +++++++++---------- tachyon/math/elliptic_curves/msm/msm_ctx.h | 2 +- .../short_weierstrass/affine_point.h | 2 +- .../short_weierstrass/jacobian_point.h | 2 +- .../short_weierstrass/point_xyzz.h | 2 +- .../short_weierstrass/projective_point.h | 2 +- .../short_weierstrass/sw_curve.h | 4 ++-- .../finite_fields/cubic_extension_field.h | 6 ++--- .../prime_field_x86.h.tpl | 16 ++++++------- .../math/finite_fields/prime_field_fallback.h | 18 +++++++------- tachyon/math/finite_fields/prime_field_gpu.h | 14 +++++------ .../finite_fields/prime_field_gpu_debug.h | 16 ++++++------- .../finite_fields/quadratic_extension_field.h | 6 ++--- .../finite_fields/small_prime_field_mont.h | 12 +++++----- .../multilinear_dense_evaluations.h | 2 +- 16 files changed, 65 insertions(+), 65 deletions(-) diff --git a/tachyon/math/base/arithmetics.h b/tachyon/math/base/arithmetics.h index 2c7d7a4e6..a9a8ebe1c 100644 --- a/tachyon/math/base/arithmetics.h +++ b/tachyon/math/base/arithmetics.h @@ -62,7 +62,7 @@ ALWAYS_INLINE SubResult SubWithBorrow(uint32_t a, uint32_t b, ALWAYS_INLINE constexpr MulResult MulAddWithCarry( uint32_t a, uint32_t b, uint32_t c, uint32_t carry = 0) { uint64_t tmp = uint64_t{a} + uint64_t{b} * uint64_t{c} + uint64_t{carry}; - MulResult result; + MulResult result{}; result.lo = static_cast(tmp); result.hi = static_cast(tmp >> 32); return result; diff --git a/tachyon/math/base/big_int.h b/tachyon/math/base/big_int.h index fadc60fd1..eb3db221e 100644 --- a/tachyon/math/base/big_int.h +++ b/tachyon/math/base/big_int.h @@ -103,7 +103,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { // Returns the maximum representable value for BigInt. constexpr static BigInt Max() { - BigInt ret; + BigInt ret{}; for (uint64_t& limb : ret.limbs) { limb = std::numeric_limits::max(); } @@ -112,7 +112,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { // Generate a random BigInt between [0, |max|). constexpr static BigInt Random(const BigInt& max = Max()) { - BigInt ret; + BigInt ret{}; for (size_t i = 0; i < N; ++i) { ret[i] = base::Uniform(base::Range::All()); } @@ -141,7 +141,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { template constexpr static BigInt FromBitsLE(const std::bitset& bits) { static_assert(BitNums <= kBitNums); - BigInt ret; + BigInt ret{}; size_t bit_idx = 0; size_t limb_idx = 0; std::bitset limb_bits; @@ -167,7 +167,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { template constexpr static BigInt FromBitsBE(const std::bitset& bits) { static_assert(BitNums <= kBitNums); - BigInt ret; + BigInt ret{}; std::bitset limb_bits; size_t bit_idx = 0; size_t limb_idx = 0; @@ -196,7 +196,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { // ordering. template constexpr static BigInt FromBytesLE(const ByteContainer& bytes) { - BigInt ret; + BigInt ret{}; size_t byte_idx = 0; size_t limb_idx = 0; uint64_t limb = 0; @@ -224,7 +224,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { // ordering. template constexpr static BigInt FromBytesBE(const ByteContainer& bytes) { - BigInt ret; + BigInt ret{}; size_t byte_idx = 0; size_t limb_idx = 0; uint64_t limb = 0; @@ -446,7 +446,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt operator&(const BigInt& other) const { - BigInt ret; + BigInt ret{}; if constexpr (N == 1) { ret[0] = limbs[0] & other[0]; } else if constexpr (N == 2) { @@ -471,7 +471,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt operator|(const BigInt& other) const { - BigInt ret; + BigInt ret{}; if constexpr (N == 1) { ret[0] = limbs[0] | other[0]; } else if constexpr (N == 2) { @@ -496,7 +496,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt operator^(const BigInt& other) const { - BigInt ret; + BigInt ret{}; if constexpr (N == 1) { ret[0] = limbs[0] ^ other[0]; } else if constexpr (N == 2) { @@ -548,7 +548,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt Add(const BigInt& other, uint64_t& carry) const { - BigInt ret; + BigInt ret{}; DoAdd(*this, other, carry, ret); return ret; } @@ -569,7 +569,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt Sub(const BigInt& other, uint64_t& borrow) const { - BigInt ret; + BigInt ret{}; DoSub(*this, other, borrow, ret); return ret; } @@ -590,7 +590,7 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } constexpr BigInt MulBy2(uint64_t& carry) const { - BigInt ret; + BigInt ret{}; DoMulBy2(*this, carry, ret); return ret; } diff --git a/tachyon/math/elliptic_curves/msm/msm_ctx.h b/tachyon/math/elliptic_curves/msm/msm_ctx.h index 3def1c796..fa3dc7850 100644 --- a/tachyon/math/elliptic_curves/msm/msm_ctx.h +++ b/tachyon/math/elliptic_curves/msm/msm_ctx.h @@ -18,7 +18,7 @@ struct TACHYON_EXPORT MSMCtx { template constexpr static MSMCtx CreateDefault(size_t size) { - MSMCtx ctx; + MSMCtx ctx{}; ctx.window_bits = ComputeWindowsBits(size); ctx.window_count = ComputeWindowsCount(ctx.window_bits); ctx.size = size; diff --git a/tachyon/math/elliptic_curves/short_weierstrass/affine_point.h b/tachyon/math/elliptic_curves/short_weierstrass/affine_point.h index 38df4bd81..37e28d743 100644 --- a/tachyon/math/elliptic_curves/short_weierstrass/affine_point.h +++ b/tachyon/math/elliptic_curves/short_weierstrass/affine_point.h @@ -61,7 +61,7 @@ class AffinePoint< constexpr static std::optional CreateFromX(const BaseField& x, bool pick_odd) { - AffinePoint point; + AffinePoint point{}; if (!Curve::GetPointFromX(x, pick_odd, &point)) return std::nullopt; return point; } diff --git a/tachyon/math/elliptic_curves/short_weierstrass/jacobian_point.h b/tachyon/math/elliptic_curves/short_weierstrass/jacobian_point.h index ab583cc35..46b1b1a2c 100644 --- a/tachyon/math/elliptic_curves/short_weierstrass/jacobian_point.h +++ b/tachyon/math/elliptic_curves/short_weierstrass/jacobian_point.h @@ -65,7 +65,7 @@ class JacobianPoint< constexpr static std::optional CreateFromX(const BaseField& x, bool pick_odd) { - JacobianPoint point; + JacobianPoint point{}; if (!Curve::GetPointFromX(x, pick_odd, &point)) return std::nullopt; return point; } diff --git a/tachyon/math/elliptic_curves/short_weierstrass/point_xyzz.h b/tachyon/math/elliptic_curves/short_weierstrass/point_xyzz.h index 2a5358dac..8e12909a7 100644 --- a/tachyon/math/elliptic_curves/short_weierstrass/point_xyzz.h +++ b/tachyon/math/elliptic_curves/short_weierstrass/point_xyzz.h @@ -69,7 +69,7 @@ class PointXYZZ<_Curve, constexpr static std::optional CreateFromX(const BaseField& x, bool pick_odd) { - PointXYZZ point; + PointXYZZ point{}; if (!Curve::GetPointFromX(x, pick_odd, &point)) return std::nullopt; return point; } diff --git a/tachyon/math/elliptic_curves/short_weierstrass/projective_point.h b/tachyon/math/elliptic_curves/short_weierstrass/projective_point.h index 09304d746..6543ed8f7 100644 --- a/tachyon/math/elliptic_curves/short_weierstrass/projective_point.h +++ b/tachyon/math/elliptic_curves/short_weierstrass/projective_point.h @@ -63,7 +63,7 @@ class ProjectivePoint< constexpr static std::optional CreateFromX( const BaseField& x, bool pick_odd) { - ProjectivePoint point; + ProjectivePoint point{}; if (!Curve::GetPointFromX(x, pick_odd, &point)) return std::nullopt; return point; } diff --git a/tachyon/math/elliptic_curves/short_weierstrass/sw_curve.h b/tachyon/math/elliptic_curves/short_weierstrass/sw_curve.h index 78fa6101f..cb60d9358 100644 --- a/tachyon/math/elliptic_curves/short_weierstrass/sw_curve.h +++ b/tachyon/math/elliptic_curves/short_weierstrass/sw_curve.h @@ -47,8 +47,8 @@ class SWCurve { template constexpr static bool GetPointFromX(const BaseField& x, bool pick_odd, Point* point) { - BaseField even_y; - BaseField odd_y; + BaseField even_y{}; + BaseField odd_y{}; if (!GetYsFromX(x, &even_y, &odd_y)) return false; if constexpr (std::is_same_v) { *point = AffinePoint(x, pick_odd ? odd_y : even_y); diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index efac9f7f0..74986915f 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -228,7 +228,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // MultiplicativeSemigroup methods constexpr Derived Mul(const Derived& other) const { - Derived ret; + Derived ret{}; DoMul(*static_cast(this), other, ret); return ret; } @@ -255,7 +255,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { } constexpr Derived SquareImpl() const { - Derived ret; + Derived ret{}; DoSquareImpl(*static_cast(this), ret); return ret; } @@ -268,7 +268,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // MultiplicativeGroup methods constexpr std::optional Inverse() const { - Derived ret; + Derived ret{}; if (LIKELY(DoInverse(*static_cast(this), ret))) return ret; LOG_IF_NOT_GPU(ERROR) << "Inverse of zero attempted"; return std::nullopt; diff --git a/tachyon/math/finite_fields/generator/prime_field_generator/prime_field_x86.h.tpl b/tachyon/math/finite_fields/generator/prime_field_generator/prime_field_x86.h.tpl index ed5feb64f..2cc5467df 100644 --- a/tachyon/math/finite_fields/generator/prime_field_generator/prime_field_x86.h.tpl +++ b/tachyon/math/finite_fields/generator/prime_field_generator/prime_field_x86.h.tpl @@ -56,7 +56,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final constexpr static PrimeField Zero() { return PrimeField(); } constexpr static PrimeField One() { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::kOne; return ret; } @@ -89,7 +89,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final } constexpr static PrimeField FromMontgomery(const BigInt& big_int) { - PrimeField ret; + PrimeField ret{}; ret.value_ = big_int; return ret; } @@ -157,7 +157,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final // AdditiveSemigroup methods constexpr PrimeField Add(const PrimeField& other) const { - PrimeField ret; + PrimeField ret{}; %{prefix}_rawAdd(ret.value_.limbs, value_.limbs, other.value_.limbs); return ret; } @@ -169,7 +169,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final // AdditiveGroup methods constexpr PrimeField Sub(const PrimeField& other) const { - PrimeField ret; + PrimeField ret{}; %{prefix}_rawSub(ret.value_.limbs, value_.limbs, other.value_.limbs); return ret; } @@ -180,7 +180,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final } constexpr PrimeField Negate() const { - PrimeField ret; + PrimeField ret{}; %{prefix}_rawNeg(ret.value_.limbs, value_.limbs); return ret; } @@ -193,7 +193,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final // TODO(chokobole): Support bigendian. // MultiplicativeSemigroup methods constexpr PrimeField Mul(const PrimeField& other) const { - PrimeField ret; + PrimeField ret{}; %{prefix}_rawMMul(ret.value_.limbs, value_.limbs, other.value_.limbs); return ret; } @@ -204,7 +204,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final } constexpr PrimeField Square() const { - PrimeField ret; + PrimeField ret{}; %{prefix}_rawMSquare(ret.value_.limbs, value_.limbs); return ret; } @@ -216,7 +216,7 @@ class PrimeField<_Config, std::enable_if_t<_Config::%{asm_flag}>> final // MultiplicativeGroup methods constexpr std::optional Inverse() const { - PrimeField ret; + PrimeField ret{}; if (LIKELY(value_.template MontgomeryInverse( Config::kModulus, Config::kMontgomeryR2, ret.value_))) { return ret; diff --git a/tachyon/math/finite_fields/prime_field_fallback.h b/tachyon/math/finite_fields/prime_field_fallback.h index acc77b3b4..068a5c5ea 100644 --- a/tachyon/math/finite_fields/prime_field_fallback.h +++ b/tachyon/math/finite_fields/prime_field_fallback.h @@ -61,7 +61,7 @@ class PrimeField<_Config, std::enable_if_t& mont) { - PrimeField ret; + PrimeField ret{}; ret.value_ = mont; return ret; } @@ -167,7 +167,7 @@ class PrimeField<_Config, std::enable_if_t::template Clamp(Config::kModulus, @@ -184,7 +184,7 @@ class PrimeField<_Config, std::enable_if_t::template Clamp(Config::kModulus, @@ -202,7 +202,7 @@ class PrimeField<_Config, std::enable_if_t value_) { ret.value_ = value_.Add(Config::kModulus); ret.value_.SubInPlace(other.value_); @@ -221,7 +221,7 @@ class PrimeField<_Config, std::enable_if_t Inverse() const { - PrimeField ret; + PrimeField ret{}; if (LIKELY(value_.template MontgomeryInverse( Config::kModulus, Config::kMontgomeryR2, ret.value_))) { return ret; diff --git a/tachyon/math/finite_fields/prime_field_gpu.h b/tachyon/math/finite_fields/prime_field_gpu.h index acc6ef272..469312e54 100644 --- a/tachyon/math/finite_fields/prime_field_gpu.h +++ b/tachyon/math/finite_fields/prime_field_gpu.h @@ -60,7 +60,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { constexpr static PrimeFieldGpu Zero() { return PrimeFieldGpu(); } constexpr static PrimeFieldGpu One() { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; ret.value_ = GetOne(); return ret; } @@ -95,7 +95,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { } constexpr static PrimeFieldGpu FromMontgomery(const BigInt& mont) { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; ret.value_ = mont; return ret; } @@ -201,7 +201,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { // AdditiveSemigroup methods __device__ constexpr PrimeFieldGpu Add(const PrimeFieldGpu& other) const { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; AddLimbs(value_, other.value_, ret.value_); return Clamp(ret); } @@ -214,7 +214,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { // AdditiveGroup methods __device__ constexpr PrimeFieldGpu Sub(const PrimeFieldGpu& other) const { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; uint64_t carry = SubLimbs(value_, other.value_, ret.value_); if (carry == 0) return ret; AddLimbs(ret.value_, GetModulus(), ret.value_); @@ -229,7 +229,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { } __device__ constexpr PrimeFieldGpu Negate() const { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; SubLimbs(GetModulus(), value_, ret.value_); return ret; } @@ -246,7 +246,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { // Forces us to think more carefully about the last carry bit if we use a // modulus with fewer than 2 leading zeroes of slack. static_assert(!(Config::kModulus[N - 1] >> 62)); - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; MulLimbs(value_, other.value_, ret.value_); return Clamp(ret); } @@ -264,7 +264,7 @@ class PrimeFieldGpu final : public PrimeFieldBase> { // MultiplicativeGroup methods __device__ constexpr std::optional Inverse() const { - PrimeFieldGpu ret; + PrimeFieldGpu ret{}; if (UNLIKELY(!DoInverse(*this, ret))) { // TODO(ashjeong): add CUDA error logging return std::nullopt; diff --git a/tachyon/math/finite_fields/prime_field_gpu_debug.h b/tachyon/math/finite_fields/prime_field_gpu_debug.h index 81eaafb68..497d5ccb8 100644 --- a/tachyon/math/finite_fields/prime_field_gpu_debug.h +++ b/tachyon/math/finite_fields/prime_field_gpu_debug.h @@ -57,13 +57,13 @@ class PrimeFieldGpuDebug final constexpr static PrimeFieldGpuDebug Zero() { return PrimeFieldGpuDebug(); } constexpr static PrimeFieldGpuDebug One() { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; ret.value_ = Config::kOne; return ret; } static PrimeFieldGpuDebug Random() { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; ret.value_ = PrimeField::Random().value(); return ret; } @@ -94,7 +94,7 @@ class PrimeFieldGpuDebug final } constexpr static PrimeFieldGpuDebug FromMontgomery(const BigInt& mont) { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; ret.value_ = mont; return ret; } @@ -170,7 +170,7 @@ class PrimeFieldGpuDebug final // AdditiveSemigroup methods constexpr PrimeFieldGpuDebug Add(const PrimeFieldGpuDebug& other) const { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; AddLimbs(value_, other.value_, ret.value_); return Clamp(ret); } @@ -183,7 +183,7 @@ class PrimeFieldGpuDebug final // AdditiveGroup methods constexpr PrimeFieldGpuDebug Sub(const PrimeFieldGpuDebug& other) const { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; uint64_t carry = SubLimbs(value_, other.value_, ret.value_); if (carry == 0) return ret; AddLimbs(ret.value_, Config::kModulus, ret.value_); @@ -198,7 +198,7 @@ class PrimeFieldGpuDebug final } constexpr PrimeFieldGpuDebug Negate() const { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; SubLimbs(Config::kModulus, value_, ret.value_); return ret; } @@ -215,7 +215,7 @@ class PrimeFieldGpuDebug final // Forces us to think more carefully about the last carry bit if we use a // modulus with fewer than 2 leading zeroes of slack. static_assert(!(Config::kModulus[N - 1] >> 62)); - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; MulLimbs(value_, other.value_, ret.value_); return Clamp(ret); } @@ -234,7 +234,7 @@ class PrimeFieldGpuDebug final // MultiplicativeGroup methods // TODO(chokobole): Share codes with PrimeField and PrimeFieldGpu. constexpr std::optional Inverse() const { - PrimeFieldGpuDebug ret; + PrimeFieldGpuDebug ret{}; if (LIKELY(DoInverse(*this, ret))) return ret; LOG_IF_NOT_GPU(ERROR) << "Inverse of zero attempted"; return std::nullopt; diff --git a/tachyon/math/finite_fields/quadratic_extension_field.h b/tachyon/math/finite_fields/quadratic_extension_field.h index 81e1166fa..796b93ce4 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field.h +++ b/tachyon/math/finite_fields/quadratic_extension_field.h @@ -191,7 +191,7 @@ class QuadraticExtensionField // MultiplicativeSemigroup methods constexpr Derived Mul(const Derived& other) const { - Derived ret; + Derived ret{}; DoMul(*static_cast(this), other, ret); return ret; } @@ -216,7 +216,7 @@ class QuadraticExtensionField } constexpr Derived SquareImpl() const { - Derived ret; + Derived ret{}; DoSquareImpl(*static_cast(this), ret); return ret; } @@ -229,7 +229,7 @@ class QuadraticExtensionField // MultiplicativeGroup methods constexpr std::optional Inverse() const { - Derived ret; + Derived ret{}; if (LIKELY(DoInverse(*static_cast(this), ret))) { return ret; } diff --git a/tachyon/math/finite_fields/small_prime_field_mont.h b/tachyon/math/finite_fields/small_prime_field_mont.h index e9397d967..50af9bb53 100644 --- a/tachyon/math/finite_fields/small_prime_field_mont.h +++ b/tachyon/math/finite_fields/small_prime_field_mont.h @@ -53,7 +53,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && constexpr static PrimeField Zero() { return PrimeField(); } constexpr static PrimeField One() { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::kOne; return ret; } @@ -87,7 +87,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && } constexpr static PrimeField FromMontgomery(const uint32_t value) { - PrimeField ret; + PrimeField ret{}; ret.value_ = value; return ret; } @@ -142,7 +142,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && // AdditiveSemigroup methods constexpr PrimeField Add(PrimeField other) const { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::AddMod(value_, other.value_); return ret; } @@ -154,7 +154,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && // AdditiveGroup methods constexpr PrimeField Sub(PrimeField other) const { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::SubMod(value_, other.value_); return ret; } @@ -165,7 +165,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && } constexpr PrimeField Negate() const { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::SubMod(0, value_); return ret; } @@ -177,7 +177,7 @@ class PrimeField<_Config, std::enable_if_t<(_Config::kModulusBits <= 32) && // MultiplicativeSemigroup methods constexpr PrimeField Mul(PrimeField other) const { - PrimeField ret; + PrimeField ret{}; ret.value_ = Config::FromMontgomery(uint64_t{value_} * other.value_); return ret; } diff --git a/tachyon/math/polynomials/multivariate/multilinear_dense_evaluations.h b/tachyon/math/polynomials/multivariate/multilinear_dense_evaluations.h index 331513eaa..eb712fb0d 100644 --- a/tachyon/math/polynomials/multivariate/multilinear_dense_evaluations.h +++ b/tachyon/math/polynomials/multivariate/multilinear_dense_evaluations.h @@ -149,7 +149,7 @@ class MultilinearDenseEvaluations { // NOTE(chokobole): This creates a polynomial that contains |F::Zero()| up to // |degree| + 1. constexpr static MultilinearDenseEvaluations Zero(size_t degree) { - MultilinearDenseEvaluations ret; + MultilinearDenseEvaluations ret{}; ret.evaluations_ = std::vector(size_t{1} << degree); return ret; } From 4172702785a434c2152504cd4f10fb71e3ee260b Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Mon, 24 Jun 2024 10:41:58 +0900 Subject: [PATCH 09/15] fix(math): remove `constexpr` from non-constexpr methods --- tachyon/math/base/big_int.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tachyon/math/base/big_int.h b/tachyon/math/base/big_int.h index eb3db221e..23cefa38f 100644 --- a/tachyon/math/base/big_int.h +++ b/tachyon/math/base/big_int.h @@ -123,14 +123,14 @@ struct ALIGNAS(internal::LimbsAlignment(N)) BigInt { } // Convert a decimal string to a BigInt. - constexpr static std::optional FromDecString(std::string_view str) { + static std::optional FromDecString(std::string_view str) { BigInt ret; if (!internal::StringToLimbs(str, ret.limbs, N)) return std::nullopt; return ret; } // Convert a hexadecimal string to a BigInt. - constexpr static std::optional FromHexString(std::string_view str) { + static std::optional FromHexString(std::string_view str) { BigInt ret; if (!(internal::HexStringToLimbs(str, ret.limbs, N))) return std::nullopt; return ret; From 99a741db41452eb6eec68cf81e3b7fd5df37a90e Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Mon, 24 Jun 2024 10:19:51 +0900 Subject: [PATCH 10/15] chore(math): update the comments in `Fpx` --- .../finite_fields/cubic_extension_field.h | 4 +-- tachyon/math/finite_fields/fp12.h | 28 ++++++++-------- tachyon/math/finite_fields/fp2.h | 2 +- tachyon/math/finite_fields/fp3.h | 10 +++--- tachyon/math/finite_fields/fp4.h | 12 +++---- tachyon/math/finite_fields/fp6.h | 32 +++++++++---------- .../finite_fields/quadratic_extension_field.h | 4 +-- 7 files changed, 46 insertions(+), 46 deletions(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index 74986915f..4e03459e3 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -304,7 +304,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // = (a.c0 * b.c0 + (a.c1 * b.c2 + a.c2 * b.c1) * q, // a.c0 * b.c1 + a.c1 * b.c0 + a.c2 * b.c2 * q, // a.c0 * b.c2 + a.c1 * b.c1 + a.c2 * b.c0) - // Where q is Config::kNonResidue. + // where q is |Config::kNonResidue|. // See https://eprint.iacr.org/2006/471.pdf // Devegili OhEig Scott Dahab --- Multiplication and Squaring on AbstractPairing-Friendly Fields.pdf; Section 4 (Karatsuba) @@ -338,7 +338,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { // = c0² + 2 * c1 * c2 * x³ + 2 * c0 * c1 * x + c2² * x⁴ + (c1² + 2 * c0 * c2) * x² // = c0² + 2 * c1 * c2 * q + (2 * c0 * c1 + c2² * q) * x + (c1² + 2 * c0 * c2) * x² // = (c0² + 2 * c1 * c2 * q, 2 * c0 * c1 + c2² * q, c1² + 2 * c0 * c2) - // Where q is Config::kNonResidue. + // where q is |Config::kNonResidue|. // See https://eprint.iacr.org/2006/471.pdf // Devegili OhEig Scott Dahab --- Multiplication and Squaring on AbstractPairing-Friendly Fields.pdf; Section 4 (CH-SQR2) diff --git a/tachyon/math/finite_fields/fp12.h b/tachyon/math/finite_fields/fp12.h index 523e487c1..85129c824 100644 --- a/tachyon/math/finite_fields/fp12.h +++ b/tachyon/math/finite_fields/fp12.h @@ -24,7 +24,7 @@ class Fp12 final : public QuadraticExtensionField> { using Fp2 = typename Fp6::BaseField; using CpuField = Fp12; - // TODO(chokobole): Implements Fp12Gpu + // TODO(chokobole): Implement Fp12Gpu using GpuField = Fp12; using QuadraticExtensionField>::QuadraticExtensionField; @@ -36,7 +36,7 @@ class Fp12 final : public QuadraticExtensionField> { static void Init() { using BaseFieldConfig = typename BaseField::Config; - // x⁶ = q = BaseFieldConfig::kNonResidue + // x⁶ = q = |BaseFieldConfig::kNonResidue| Config::Init(); @@ -107,33 +107,33 @@ class Fp12 final : public QuadraticExtensionField> { #undef SET_EXP_GMP - // kFrobeniusCoeffs[0] = q^((P⁰ - 1) / 6) + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 6) Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); #define SET_FROBENIUS_COEFF(d) \ BigInt exp##d; \ gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ Config::kFrobeniusCoeffs[d] = BaseFieldConfig::kNonResidue.Pow(exp##d) - // kFrobeniusCoeffs[1] = q^(exp₁) = q^((P¹ - 1) / 6) + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 6) SET_FROBENIUS_COEFF(1); - // kFrobeniusCoeffs[2] = q^(exp₂) = q^((P² - 1) / 6) + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 6) SET_FROBENIUS_COEFF(2); - // kFrobeniusCoeffs[3] = q^(exp₃) = q^((P³ - 1) / 6) + // |kFrobeniusCoeffs[3]| = q^(exp₃) = q^((P³ - 1) / 6) SET_FROBENIUS_COEFF(3); - // kFrobeniusCoeffs[4] = q^(exp₄) = q^((P⁴ - 1) / 6) + // |kFrobeniusCoeffs[4]| = q^(exp₄) = q^((P⁴ - 1) / 6) SET_FROBENIUS_COEFF(4); - // kFrobeniusCoeffs[5] = q^(exp₅) = q^((P⁵ - 1) / 6) + // |kFrobeniusCoeffs[5]| = q^(exp₅) = q^((P⁵ - 1) / 6) SET_FROBENIUS_COEFF(5); - // kFrobeniusCoeffs[6] = q^(exp₆) = q^((P⁶ - 1) / 6) + // |kFrobeniusCoeffs[6]| = q^(exp₆) = q^((P⁶ - 1) / 6) SET_FROBENIUS_COEFF(6); - // kFrobeniusCoeffs[7] = q^(exp₇) = q^((P⁷ - 1) / 6) + // |kFrobeniusCoeffs[7]| = q^(exp₇) = q^((P⁷ - 1) / 6) SET_FROBENIUS_COEFF(7); - // kFrobeniusCoeffs[8] = q^(exp₈) = q^((P⁸ - 1) / 6) + // |kFrobeniusCoeffs[8]| = q^(exp₈) = q^((P⁸ - 1) / 6) SET_FROBENIUS_COEFF(8); - // kFrobeniusCoeffs[9] = q^(exp₉) = q^((P⁹ - 1) / 6) + // |kFrobeniusCoeffs[9]| = q^(exp₉) = q^((P⁹ - 1) / 6) SET_FROBENIUS_COEFF(9); - // kFrobeniusCoeffs[10] = q^(exp₁₀) = q^((P¹⁰ - 1) / 6) + // |kFrobeniusCoeffs[10]| = q^(exp₁₀) = q^((P¹⁰ - 1) / 6) SET_FROBENIUS_COEFF(10); - // kFrobeniusCoeffs[11] = q^(exp₁₁) = q^((P¹¹ - 1) / 6) + // |kFrobeniusCoeffs[11]| = q^(exp₁₁) = q^((P¹¹ - 1) / 6) SET_FROBENIUS_COEFF(11); #undef SET_FROBENIUS_COEFF diff --git a/tachyon/math/finite_fields/fp2.h b/tachyon/math/finite_fields/fp2.h index 6f72c9c8e..3691104e1 100644 --- a/tachyon/math/finite_fields/fp2.h +++ b/tachyon/math/finite_fields/fp2.h @@ -18,7 +18,7 @@ class Fp2 final : public QuadraticExtensionField> { using FrobeniusCoefficient = typename Config::FrobeniusCoefficient; using CpuField = Fp2; - // TODO(chokobole): Implements Fp2Gpu + // TODO(chokobole): Implement Fp2Gpu using GpuField = Fp2; using QuadraticExtensionField>::QuadraticExtensionField; diff --git a/tachyon/math/finite_fields/fp3.h b/tachyon/math/finite_fields/fp3.h index 8ec3789f6..35e10684c 100644 --- a/tachyon/math/finite_fields/fp3.h +++ b/tachyon/math/finite_fields/fp3.h @@ -19,7 +19,7 @@ class Fp3 final : public CubicExtensionField> { using FrobeniusCoefficient = typename Config::FrobeniusCoefficient; using CpuField = Fp3; - // TODO(chokobole): Implements Fp3Gpu + // TODO(chokobole): Implement Fp3Gpu using GpuField = Fp3; using CubicExtensionField>::CubicExtensionField; @@ -31,7 +31,7 @@ class Fp3 final : public CubicExtensionField> { static void Init() { Config::Init(); - // x³ = q = Config::kNonResidue + // x³ = q = |Config::kNonResidue| // αᴾ = (α₀ + α₁x + α₂x²)ᴾ // = α₀ᴾ + α₁ᴾxᴾ + α₂ᴾx²ᴾ @@ -65,16 +65,16 @@ class Fp3 final : public CubicExtensionField> { #undef SET_EXP_GMP - // kFrobeniusCoeffs[0] = q^((P⁰ - 1) / 3) = 1 + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 3) = 1 Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); #define SET_FROBENIUS_COEFF(d) \ BigInt exp##d; \ gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ Config::kFrobeniusCoeffs[d] = Config::kNonResidue.Pow(exp##d) - // kFrobeniusCoeffs[1] = q^(exp₁) = q^((P¹ - 1) / 3) = ω + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 3) = ω SET_FROBENIUS_COEFF(1); - // kFrobeniusCoeffs[2] = q^(exp₂) = q^((P² - 1) / 3) + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 3) SET_FROBENIUS_COEFF(2); #undef SET_FROBENIUS_COEFF diff --git a/tachyon/math/finite_fields/fp4.h b/tachyon/math/finite_fields/fp4.h index 18cab4f2c..aa892073d 100644 --- a/tachyon/math/finite_fields/fp4.h +++ b/tachyon/math/finite_fields/fp4.h @@ -19,7 +19,7 @@ class Fp4 final : public QuadraticExtensionField> { using FrobeniusCoefficient = typename Config::FrobeniusCoefficient; using CpuField = Fp4; - // TODO(chokobole): Implements Fp4Gpu + // TODO(chokobole): Implement Fp4Gpu using GpuField = Fp4; using QuadraticExtensionField>::QuadraticExtensionField; @@ -31,7 +31,7 @@ class Fp4 final : public QuadraticExtensionField> { static void Init() { using BaseFieldConfig = typename BaseField::Config; - // x⁴ = q = BaseFieldConfig::kNonResidue + // x⁴ = q = |BaseFieldConfig::kNonResidue| Config::Init(); @@ -71,18 +71,18 @@ class Fp4 final : public QuadraticExtensionField> { #undef SET_EXP_GMP - // kFrobeniusCoeffs[0] = q^((P⁰ - 1) / 4) = 1 + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 4) = 1 Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); #define SET_FROBENIUS_COEFF(d) \ BigInt exp##d; \ gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ Config::kFrobeniusCoeffs[d] = BaseFieldConfig::kNonResidue.Pow(exp##d) - // kFrobeniusCoeffs[1] = q^(exp₁) = q^((P¹ - 1) / 4) = ω + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 4) = ω SET_FROBENIUS_COEFF(1); - // kFrobeniusCoeffs[2] = q^(exp₂) = q^((P² - 1) / 4) + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 4) SET_FROBENIUS_COEFF(2); - // kFrobeniusCoeffs[3] = q^(exp₃) = q^((P³ - 1) / 4) + // |kFrobeniusCoeffs[3]| = q^(exp₃) = q^((P³ - 1) / 4) SET_FROBENIUS_COEFF(3); #undef SET_FROBENIUS_COEFF diff --git a/tachyon/math/finite_fields/fp6.h b/tachyon/math/finite_fields/fp6.h index a0d441e14..b3cc8bc1e 100644 --- a/tachyon/math/finite_fields/fp6.h +++ b/tachyon/math/finite_fields/fp6.h @@ -26,7 +26,7 @@ class Fp6> final using Fp3 = BaseField; using CpuField = Fp6; - // TODO(chokobole): Implements Fp6Gpu + // TODO(chokobole): Implement Fp6Gpu using GpuField = Fp6; using QuadraticExtensionField>::QuadraticExtensionField; @@ -37,7 +37,7 @@ class Fp6> final static void Init() { using BaseFieldConfig = typename BaseField::Config; - // x⁶ = q = BaseFieldConfig::kNonResidue + // x⁶ = q = |BaseFieldConfig::kNonResidue| Config::Init(); @@ -84,22 +84,22 @@ class Fp6> final #undef SET_EXP_GMP - // kFrobeniusCoeffs[0] = q^((P⁰ - 1) / 6) = 1 + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 6) = 1 Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); #define SET_FROBENIUS_COEFF(d) \ BigInt exp##d; \ gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ Config::kFrobeniusCoeffs[d] = BaseFieldConfig::kNonResidue.Pow(exp##d) - // kFrobeniusCoeffs[1] = q^(exp₁) = q^((P¹ - 1) / 6) = ω + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 6) = ω SET_FROBENIUS_COEFF(1); - // kFrobeniusCoeffs[2] = q^(exp₂) = q^((P² - 1) / 6) + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 6) SET_FROBENIUS_COEFF(2); - // kFrobeniusCoeffs[3] = q^(exp₃) = q^((P³ - 1) / 6) + // |kFrobeniusCoeffs[3]| = q^(exp₃) = q^((P³ - 1) / 6) SET_FROBENIUS_COEFF(3); - // kFrobeniusCoeffs[4] = q^(exp₄) = q^((P⁴ - 1) / 6) + // |kFrobeniusCoeffs[4]| = q^(exp₄) = q^((P⁴ - 1) / 6) SET_FROBENIUS_COEFF(4); - // kFrobeniusCoeffs[5] = q^(exp₅) = q^((P⁵ - 1) / 6) + // |kFrobeniusCoeffs[5]| = q^(exp₅) = q^((P⁵ - 1) / 6) SET_FROBENIUS_COEFF(5); #undef SET_FROBENIUS_COEFF @@ -221,7 +221,7 @@ class Fp6> final using Fp2 = BaseField; using CpuField = Fp6; - // TODO(chokobole): Implements Fp6Gpu + // TODO(chokobole): Implement Fp6Gpu using GpuField = Fp6; using CubicExtensionField>::CubicExtensionField; @@ -232,7 +232,7 @@ class Fp6> final static void Init() { Config::Init(); - // x³ = q = Config::kNonResidue + // x³ = q = |Config::kNonResidue| // αᴾ = (α₀ + α₁x + α₂x²)ᴾ // = α₀ᴾ + α₁ᴾxᴾ + α₂ᴾx²ᴾ @@ -278,22 +278,22 @@ class Fp6> final #undef SET_EXP_GMP - // kFrobeniusCoeffs[0] = q^((P⁰ - 1) / 3) + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 3) Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); #define SET_FROBENIUS_COEFF(d) \ BigInt exp##d; \ gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ Config::kFrobeniusCoeffs[d] = Config::kNonResidue.Pow(exp##d) - // kFrobeniusCoeffs[1] = q^(exp₁) = q^((P¹ - 1) / 3) + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 3) SET_FROBENIUS_COEFF(1); - // kFrobeniusCoeffs[2] = q^(exp₂) = q^((P² - 1) / 3) + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 3) SET_FROBENIUS_COEFF(2); - // kFrobeniusCoeffs[3] = q^(exp₃) = q^((P³ - 1) / 3) + // |kFrobeniusCoeffs[3]| = q^(exp₃) = q^((P³ - 1) / 3) SET_FROBENIUS_COEFF(3); - // kFrobeniusCoeffs[4] = q^(exp₄) = q^((P⁴ - 1) / 3) + // |kFrobeniusCoeffs[4]| = q^(exp₄) = q^((P⁴ - 1) / 3) SET_FROBENIUS_COEFF(4); - // kFrobeniusCoeffs[5] = q^(exp₅) = q^((P⁵ - 1) / 3) + // |kFrobeniusCoeffs[5]| = q^(exp₅) = q^((P⁵ - 1) / 3) SET_FROBENIUS_COEFF(5); #undef SET_FROBENIUS_COEFF diff --git a/tachyon/math/finite_fields/quadratic_extension_field.h b/tachyon/math/finite_fields/quadratic_extension_field.h index 796b93ce4..046ac965e 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field.h +++ b/tachyon/math/finite_fields/quadratic_extension_field.h @@ -281,7 +281,7 @@ class QuadraticExtensionField // = a.c0 * b.c0 + a.c1 * b.c1 * x² + (a.c0 * b.c1 + a.c1 * b.c0) * x // = a.c0 * b.c0 + a.c1 * b.c1 * q + (a.c0 * b.c1 + a.c1 * b.c0) * x // = (a.c0 * b.c0 + a.c1 * b.c1 * q, a.c0 * b.c1 + a.c1 * b.c0) - // Where q is Config::kNonResidue. + // where q is |Config::kNonResidue|. // clang-format on if constexpr (ExtensionDegree() == 2) { BaseField c0; @@ -325,7 +325,7 @@ class QuadraticExtensionField // = c0² + c1² * x² + 2 * c0 * c1 * x // = c0² + c1² * q + 2 * c0 * c1 * x // = (c0² + c1² * q, 2 * c0 * c1) - // Where q is Config::kNonResidue. + // where q is |Config::kNonResidue|. // When q = -1, we can re-use intermediate additions to improve performance. // v0 = c0 - c1 From 68adf93c33cfd7be2906d6e6f5e111d815942173 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Mon, 24 Jun 2024 13:10:55 +0900 Subject: [PATCH 11/15] chore(math): update the comments in `cubic_extension_field.h` --- tachyon/math/finite_fields/cubic_extension_field.h | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index 4e03459e3..f66a654bd 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -75,24 +75,23 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { return 3 * BaseField::ExtensionDegree(); } - // Calculate the norm of an element with respect to the base field - // |BaseField|. The norm maps an element |a| in the extension field - // Fqᵐ to an element in the |BaseField| Fq. - // |a.Norm() = a * a^q * a^q²| + // Calculate the norm of an element with respect to |BaseField|. + // The norm maps an element |a| in the extension field Fqᵐ to an element + // in the |BaseField| Fq. |a.Norm() = a * a^q * a^q²| constexpr BaseField Norm() const { - // w.r.t to |BaseField|, we need the 0th, 1st & 2nd powers of q + // w.r.t to |BaseField|, we need the 0th, 1st & 2nd powers of q. // Since Frobenius coefficients on the towered extensions are // indexed w.r.t. to |BasePrimeField|, we need to calculate the correct // index. // NOTE(chokobole): This assumes that |BaseField::ExtensionDegree()| - // never overflows even on 32 bit machine. + // never overflows even on a 32-bit machine. size_t index_multiplier = size_t{BaseField::ExtensionDegree()}; Derived self_to_p = static_cast(*this); self_to_p.FrobeniusMapInPlace(index_multiplier); Derived self_to_p2 = static_cast(*this); self_to_p2.FrobeniusMapInPlace(2 * index_multiplier); self_to_p *= (self_to_p2 * static_cast(*this)); - // NOTE(chokobole): below CHECK() is not a device code. + // NOTE(chokobole): The |CHECK()| below is not device code. // See https://github.com/kroma-network/tachyon/issues/76 CHECK(self_to_p.c1().IsZero() && self_to_p.c2().IsZero()); return self_to_p.c0(); From f2f80b44a91340c0009676b43e1a706847f58990 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Tue, 25 Jun 2024 13:18:10 +0900 Subject: [PATCH 12/15] refac(math): simplify comparison in `(Cubic|Quaratic)ExtensionField` --- tachyon/math/finite_fields/cubic_extension_field.h | 14 +++----------- .../math/finite_fields/quadratic_extension_field.h | 8 +++----- 2 files changed, 6 insertions(+), 16 deletions(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index f66a654bd..fad6b6c06 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -128,7 +128,7 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { } constexpr bool operator!=(const Derived& other) const { - return c0_ != other.c0_ || c1_ != other.c1_ || c2_ != other.c2_; + return !operator==(other); } constexpr bool operator<(const Derived& other) const { @@ -148,19 +148,11 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { } constexpr bool operator<=(const Derived& other) const { - if (c2_ == other.c2_) { - if (c1_ == other.c1_) return c0_ <= other.c0_; - return c1_ <= other.c1_; - } - return c2_ <= other.c2_; + return !operator>(other); } constexpr bool operator>=(const Derived& other) const { - if (c2_ == other.c2_) { - if (c1_ == other.c1_) return c0_ >= other.c0_; - return c1_ >= other.c1_; - } - return c2_ >= other.c2_; + return !operator<(other); } // AdditiveSemigroup methods diff --git a/tachyon/math/finite_fields/quadratic_extension_field.h b/tachyon/math/finite_fields/quadratic_extension_field.h index 046ac965e..b5e416a0c 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field.h +++ b/tachyon/math/finite_fields/quadratic_extension_field.h @@ -112,7 +112,7 @@ class QuadraticExtensionField } constexpr bool operator!=(const Derived& other) const { - return c0_ != other.c0_ || c1_ != other.c1_; + return !operator==(other); } constexpr bool operator<(const Derived& other) const { @@ -126,13 +126,11 @@ class QuadraticExtensionField } constexpr bool operator<=(const Derived& other) const { - if (c1_ == other.c1_) return c0_ <= other.c0_; - return c1_ <= other.c1_; + return !operator>(other); } constexpr bool operator>=(const Derived& other) const { - if (c1_ == other.c1_) return c0_ >= other.c0_; - return c1_ >= other.c1_; + return !operator<(other); } // AdditiveSemigroup methods From e4d727035583580b527120cfd24709ac3b37086a Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Mon, 24 Jun 2024 10:49:40 +0900 Subject: [PATCH 13/15] refac: follow constant naming convention --- .../math/finite_fields/cubic_extension_field.h | 10 +++++----- .../finite_fields/quadratic_extension_field.h | 8 ++++---- .../univariate_evaluation_domain_unittest.cc | 18 +++++++++--------- tachyon/zk/plonk/vanishing/vanishing_utils.h | 8 ++++---- .../linear_combination_unittest.cc | 9 +++++---- tachyon/zk/r1cs/groth16/groth16_unittest.cc | 6 +++--- vendors/circom/benchmark/circom_benchmark.cc | 8 ++++---- 7 files changed, 34 insertions(+), 33 deletions(-) diff --git a/tachyon/math/finite_fields/cubic_extension_field.h b/tachyon/math/finite_fields/cubic_extension_field.h index fad6b6c06..9383b12dc 100644 --- a/tachyon/math/finite_fields/cubic_extension_field.h +++ b/tachyon/math/finite_fields/cubic_extension_field.h @@ -49,16 +49,16 @@ class CubicExtensionField : public CyclotomicMultiplicativeSubgroup { static Derived FromBasePrimeFields( absl::Span prime_fields) { CHECK_EQ(prime_fields.size(), ExtensionDegree()); - constexpr size_t base_field_degree = BaseField::ExtensionDegree(); - if constexpr (base_field_degree == 1) { + constexpr size_t kBaseFieldDegree = BaseField::ExtensionDegree(); + if constexpr (kBaseFieldDegree == 1) { return Derived(prime_fields[0], prime_fields[1], prime_fields[2]); } else { BaseField c0 = BaseField::FromBasePrimeFields( - prime_fields.subspan(0, base_field_degree)); + prime_fields.subspan(0, kBaseFieldDegree)); BaseField c1 = BaseField::FromBasePrimeFields( - prime_fields.subspan(base_field_degree, base_field_degree)); + prime_fields.subspan(kBaseFieldDegree, kBaseFieldDegree)); BaseField c2 = BaseField::FromBasePrimeFields( - prime_fields.subspan(2 * base_field_degree, base_field_degree)); + prime_fields.subspan(2 * kBaseFieldDegree, kBaseFieldDegree)); return Derived(std::move(c0), std::move(c1), std::move(c2)); } } diff --git a/tachyon/math/finite_fields/quadratic_extension_field.h b/tachyon/math/finite_fields/quadratic_extension_field.h index b5e416a0c..01452fc54 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field.h +++ b/tachyon/math/finite_fields/quadratic_extension_field.h @@ -47,14 +47,14 @@ class QuadraticExtensionField static Derived FromBasePrimeFields( absl::Span prime_fields) { CHECK_EQ(prime_fields.size(), ExtensionDegree()); - constexpr size_t base_field_degree = BaseField::ExtensionDegree(); - if constexpr (base_field_degree == 1) { + constexpr size_t kBaseFieldDegree = BaseField::ExtensionDegree(); + if constexpr (kBaseFieldDegree == 1) { return Derived(prime_fields[0], prime_fields[1]); } else { BaseField c0 = BaseField::FromBasePrimeFields( - prime_fields.subspan(0, base_field_degree)); + prime_fields.subspan(0, kBaseFieldDegree)); BaseField c1 = BaseField::FromBasePrimeFields( - prime_fields.subspan(base_field_degree)); + prime_fields.subspan(kBaseFieldDegree)); return Derived(std::move(c0), std::move(c1)); } } diff --git a/tachyon/math/polynomials/univariate/univariate_evaluation_domain_unittest.cc b/tachyon/math/polynomials/univariate/univariate_evaluation_domain_unittest.cc index a0564bd56..1e1044dca 100644 --- a/tachyon/math/polynomials/univariate/univariate_evaluation_domain_unittest.cc +++ b/tachyon/math/polynomials/univariate/univariate_evaluation_domain_unittest.cc @@ -257,13 +257,13 @@ TYPED_TEST(UnivariateEvaluationDomainTest, FFTCorrectness) { using DensePoly = typename Domain::DensePoly; using Evals = typename Domain::Evals; - // NOTE(TomTaehoonKim): |degree| is set to 2⁵ -1 for default, but for the + // NOTE(TomTaehoonKim): |kDegree| is set to 2⁵ -1 for default, but for the // |Radix2EvaluationDomain| test with openmp, it is updated to // |Radix2EvaluationDomain::kDefaultMinNumChunksForCompaction| - 1. - size_t log_degree = 5; - size_t degree = (size_t{1} << log_degree) - 1; - DensePoly rand_poly = DensePoly::Random(degree); - for (size_t log_domain_size = log_degree; log_domain_size < log_degree + 2; + const size_t kLogDegree = 5; + const size_t kDegree = (size_t{1} << kLogDegree) - 1; + DensePoly rand_poly = DensePoly::Random(kDegree); + for (size_t log_domain_size = kLogDegree; log_domain_size < kLogDegree + 2; ++log_domain_size) { size_t domain_size = size_t{1} << log_domain_size; this->TestDomains( @@ -287,10 +287,10 @@ TYPED_TEST(UnivariateEvaluationDomainTest, DegreeAwareFFTCorrectness) { using Evals = typename Domain::Evals; if constexpr (std::is_same_v) { - const size_t log_degree = 5; - const size_t degree = (size_t{1} << log_degree) - 1; - DensePoly rand_poly = DensePoly::Random(degree); - size_t domain_size = (degree + 1) * Domain::kDegreeAwareFFTThresholdFactor; + const size_t kLogDegree = 5; + const size_t kDegree = (size_t{1} << kLogDegree) - 1; + DensePoly rand_poly = DensePoly::Random(kDegree); + size_t domain_size = (kDegree + 1) * Domain::kDegreeAwareFFTThresholdFactor; this->TestDomains(domain_size, [domain_size, &rand_poly](const BaseDomain& d) { Evals deg_aware_fft_evals = d.FFT(rand_poly); diff --git a/tachyon/zk/plonk/vanishing/vanishing_utils.h b/tachyon/zk/plonk/vanishing/vanishing_utils.h index 815ab4a46..c8993e57d 100644 --- a/tachyon/zk/plonk/vanishing/vanishing_utils.h +++ b/tachyon/zk/plonk/vanishing/vanishing_utils.h @@ -57,13 +57,13 @@ ExtendedEvals& DivideByVanishingPolyInPlace( // Compute the evaluations of t(X) = Xⁿ - 1 in the coset evaluation domain. // We don't have to compute all of them, because it will repeat. - const size_t t_evaluations_size = size_t{1} - << (extended_domain->log_size_of_group() - - domain->log_size_of_group()); + const size_t kTEvaluationsSize = size_t{1} + << (extended_domain->log_size_of_group() - + domain->log_size_of_group()); // |coset_gen_pow_n| = w'ⁿ where w' is generator of extended domain. const F coset_gen_pow_n = extended_domain->group_gen().Pow(domain->size()); const F zeta_pow_n = zeta.Pow(domain->size()); - std::vector t_evaluations(t_evaluations_size); + std::vector t_evaluations(kTEvaluationsSize); // |t_evaluations| = [ζⁿ - 1, ζⁿ * w'ⁿ - 1, ζⁿ * w'²ⁿ - 1, ...] base::Parallelize(t_evaluations, [&coset_gen_pow_n, &zeta_pow_n]( absl::Span chunk, size_t chunk_offset, diff --git a/tachyon/zk/r1cs/constraint_system/linear_combination_unittest.cc b/tachyon/zk/r1cs/constraint_system/linear_combination_unittest.cc index 9896d0ec6..96cb4a178 100644 --- a/tachyon/zk/r1cs/constraint_system/linear_combination_unittest.cc +++ b/tachyon/zk/r1cs/constraint_system/linear_combination_unittest.cc @@ -138,11 +138,12 @@ TEST_F(LinearCombinationTest, ScalarMul) { } TEST_F(LinearCombinationTest, BinarySearch) { - const size_t size = 10; - std::vector> terms = base::CreateVector( - size, [](size_t i) { return Term(F(1), Variable::Instance(i << 1)); }); + const size_t kSize = 10; + std::vector> terms = base::CreateVector(kSize, [](size_t i) { + return Term(F(1), Variable::Instance(i << 1)); + }); LinearCombination lc(std::move(terms)); - for (size_t i = 0; i < (size << 1); ++i) { + for (size_t i = 0; i < (kSize << 1); ++i) { size_t index; if (i % 2 == 0) { ASSERT_TRUE(lc.BinarySearch(Variable::Instance(i), &index)); diff --git a/tachyon/zk/r1cs/groth16/groth16_unittest.cc b/tachyon/zk/r1cs/groth16/groth16_unittest.cc index 3d70560ef..48c4babc2 100644 --- a/tachyon/zk/r1cs/groth16/groth16_unittest.cc +++ b/tachyon/zk/r1cs/groth16/groth16_unittest.cc @@ -14,7 +14,7 @@ namespace { using F = math::bn254::Fr; using Curve = math::bn254::BN254Curve; -constexpr size_t MaxDegree = 31; +constexpr size_t kMaxDegree = 31; class Groth16Test : public testing::Test { public: @@ -28,10 +28,10 @@ TEST_F(Groth16Test, ProveAndVerify) { ToxicWaste toxic_waste = ToxicWaste::RandomWithoutX(); ProvingKey pk; bool loaded = - pk.Load>(toxic_waste, circuit); + pk.Load>(toxic_waste, circuit); ASSERT_TRUE(loaded); Proof proof = - CreateProofWithReductionZK>( + CreateProofWithReductionZK>( circuit, pk); PreparedVerifyingKey pvk = std::move(pk).TakeVerifyingKey().ToPreparedVerifyingKey(); diff --git a/vendors/circom/benchmark/circom_benchmark.cc b/vendors/circom/benchmark/circom_benchmark.cc index 243cce87d..4355d1840 100644 --- a/vendors/circom/benchmark/circom_benchmark.cc +++ b/vendors/circom/benchmark/circom_benchmark.cc @@ -23,7 +23,7 @@ using namespace math; using F = bn254::Fr; using Curve = math::bn254::BN254Curve; -constexpr size_t MaxDegree = (size_t{1} << 16) - 1; +constexpr size_t kMaxDegree = (size_t{1} << 16) - 1; void CheckPublicInput(const std::vector& in, absl::Span public_inputs) { @@ -59,7 +59,7 @@ int RealMain(int argc, char** argv) { Curve::Init(); std::vector>> runners; - runners.push_back(std::make_unique>( + runners.push_back(std::make_unique>( base::FilePath("benchmark/sha256_512_cpp/sha256_512.dat"))); runners.push_back(std::make_unique>( base::FilePath("benchmark/sha256_512_verification_key.json"))); @@ -75,8 +75,8 @@ int RealMain(int argc, char** argv) { runner->LoadZkey(base::FilePath("benchmark/sha256_512.zkey")); if (i == 0) { - TachyonRunner* tachyon_runner = - reinterpret_cast*>(runner.get()); + TachyonRunner* tachyon_runner = + reinterpret_cast*>(runner.get()); WitnessLoader& witness_loader = tachyon_runner->witness_loader(); From 6c6534d35ccce9e02c143b50acf2d247cd0f66cc Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Tue, 25 Jun 2024 15:21:46 +0900 Subject: [PATCH 14/15] test(math): use `EXPECT_(LT|LE|GT|GE)` for testing comparison --- .../binary_fields/binary_fields_unittest.cc | 8 +++---- .../cubic_extension_field_unittest.cc | 24 +++++++++---------- .../prime_field_generator_unittest.cc | 8 +++---- .../finite_fields/prime_field_unittest.cc | 8 +++---- .../quadratic_extension_field_unittest.cc | 16 ++++++------- 5 files changed, 32 insertions(+), 32 deletions(-) diff --git a/tachyon/math/finite_fields/binary_fields/binary_fields_unittest.cc b/tachyon/math/finite_fields/binary_fields/binary_fields_unittest.cc index 62fad20fc..a670d05b2 100644 --- a/tachyon/math/finite_fields/binary_fields/binary_fields_unittest.cc +++ b/tachyon/math/finite_fields/binary_fields/binary_fields_unittest.cc @@ -79,10 +79,10 @@ TYPED_TEST(BinaryFieldsTest, ComparisonOperator) { if constexpr (BinaryField::Config::kModulusBits > 3) { BinaryField f(3); BinaryField f2(4); - EXPECT_TRUE(f < f2); - EXPECT_TRUE(f <= f2); - EXPECT_FALSE(f > f2); - EXPECT_FALSE(f >= f2); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); } else { GTEST_SKIP() << "Modulus is too small"; } diff --git a/tachyon/math/finite_fields/cubic_extension_field_unittest.cc b/tachyon/math/finite_fields/cubic_extension_field_unittest.cc index d74817fbb..d2286a260 100644 --- a/tachyon/math/finite_fields/cubic_extension_field_unittest.cc +++ b/tachyon/math/finite_fields/cubic_extension_field_unittest.cc @@ -64,24 +64,24 @@ TEST_F(CubicExtensionFieldTest, EqualityOperators) { TEST_F(CubicExtensionFieldTest, ComparisonOperator) { GF7_3 f(GF7(3), GF7(4), GF7(5)); GF7_3 f2(GF7(4), GF7(4), GF7(5)); - EXPECT_TRUE(f < f2); - EXPECT_TRUE(f <= f2); - EXPECT_FALSE(f > f2); - EXPECT_FALSE(f >= f2); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); GF7_3 f3(GF7(4), GF7(3), GF7(5)); GF7_3 f4(GF7(3), GF7(4), GF7(5)); - EXPECT_TRUE(f3 < f4); - EXPECT_TRUE(f3 <= f4); - EXPECT_FALSE(f3 > f4); - EXPECT_FALSE(f3 >= f4); + EXPECT_LT(f3, f4); + EXPECT_LE(f3, f4); + EXPECT_GT(f4, f3); + EXPECT_GE(f4, f3); GF7_3 f5(GF7(4), GF7(5), GF7(3)); GF7_3 f6(GF7(3), GF7(2), GF7(5)); - EXPECT_TRUE(f5 < f6); - EXPECT_TRUE(f5 <= f6); - EXPECT_FALSE(f5 > f6); - EXPECT_FALSE(f5 >= f6); + EXPECT_LT(f5, f6); + EXPECT_LE(f5, f6); + EXPECT_GT(f6, f5); + EXPECT_GE(f6, f5); } TEST_F(CubicExtensionFieldTest, AdditiveOperators) { diff --git a/tachyon/math/finite_fields/prime_field_generator_unittest.cc b/tachyon/math/finite_fields/prime_field_generator_unittest.cc index 7b98d3c68..818647f43 100644 --- a/tachyon/math/finite_fields/prime_field_generator_unittest.cc +++ b/tachyon/math/finite_fields/prime_field_generator_unittest.cc @@ -82,10 +82,10 @@ TYPED_TEST(PrimeFieldGeneratorTest, ComparisonOperator) { using PrimeField = TypeParam; PrimeField f(3); PrimeField f2(4); - EXPECT_TRUE(f < f2); - EXPECT_TRUE(f <= f2); - EXPECT_FALSE(f > f2); - EXPECT_FALSE(f >= f2); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); } TYPED_TEST(PrimeFieldGeneratorTest, AdditiveGroupOperators) { diff --git a/tachyon/math/finite_fields/prime_field_unittest.cc b/tachyon/math/finite_fields/prime_field_unittest.cc index 29d9e7259..f14525b02 100644 --- a/tachyon/math/finite_fields/prime_field_unittest.cc +++ b/tachyon/math/finite_fields/prime_field_unittest.cc @@ -76,10 +76,10 @@ TYPED_TEST(PrimeFieldTest, ComparisonOperator) { F f(3); F f2(4); - EXPECT_TRUE(f < f2); - EXPECT_TRUE(f <= f2); - EXPECT_FALSE(f > f2); - EXPECT_FALSE(f >= f2); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); } TYPED_TEST(PrimeFieldTest, AdditiveOperators) { diff --git a/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc b/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc index 14671c1d6..75a389e41 100644 --- a/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc +++ b/tachyon/math/finite_fields/quadratic_extension_field_unittest.cc @@ -68,17 +68,17 @@ TEST_F(QuadraticExtensionFieldTest, EqualityOperators) { TEST_F(QuadraticExtensionFieldTest, ComparisonOperator) { GF7_2 f(GF7(3), GF7(4)); GF7_2 f2(GF7(4), GF7(4)); - EXPECT_TRUE(f < f2); - EXPECT_TRUE(f <= f2); - EXPECT_FALSE(f > f2); - EXPECT_FALSE(f >= f2); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); GF7_2 f3(GF7(4), GF7(3)); GF7_2 f4(GF7(4), GF7(4)); - EXPECT_TRUE(f3 < f4); - EXPECT_TRUE(f3 <= f4); - EXPECT_FALSE(f3 > f4); - EXPECT_FALSE(f3 >= f4); + EXPECT_LT(f3, f4); + EXPECT_LE(f3, f4); + EXPECT_GT(f4, f3); + EXPECT_GE(f4, f3); } TEST_F(QuadraticExtensionFieldTest, AdditiveOperators) { From b1786019729430dd8e6ebca4c9a5219bf3d25406 Mon Sep 17 00:00:00 2001 From: Ryan Kim Date: Thu, 20 Jun 2024 21:07:06 +0900 Subject: [PATCH 15/15] feat(math): add quartic extension for `BabyBear` and `KoalaBear` See https://github.com/Plonky3/Plonky3/blob/11ac8745b21295d5699f47089dbd927836f53fff/baby-bear/src/extension.rs#L6-L11 and https://github.com/Plonky3/Plonky3/blob/11ac8745b21295d5699f47089dbd927836f53fff/koala-bear/src/extension.rs#L40-L45. --- tachyon/math/circle/stark/BUILD.bazel | 1 + tachyon/math/finite_fields/BUILD.bazel | 15 + .../math/finite_fields/baby_bear/BUILD.bazel | 12 + .../finite_fields/finite_field_forwards.h | 2 +- tachyon/math/finite_fields/fp4.h | 112 ++- .../ext_prime_field_generator/build_defs.bzl | 1 - .../ext_prime_field_generator.cc | 7 +- .../ext_prime_field_generator/fq.h.tpl | 7 + .../math/finite_fields/koala_bear/BUILD.bazel | 13 +- .../finite_fields/quartic_extension_field.h | 646 ++++++++++++++++++ .../quartic_extension_field_unittest.cc | 204 ++++++ 11 files changed, 1014 insertions(+), 6 deletions(-) create mode 100644 tachyon/math/finite_fields/quartic_extension_field.h create mode 100644 tachyon/math/finite_fields/quartic_extension_field_unittest.cc diff --git a/tachyon/math/circle/stark/BUILD.bazel b/tachyon/math/circle/stark/BUILD.bazel index 903d4b266..f1387b30e 100644 --- a/tachyon/math/circle/stark/BUILD.bazel +++ b/tachyon/math/circle/stark/BUILD.bazel @@ -48,6 +48,7 @@ generate_fp2s( generate_fp4s( name = "fq4", base_field = "Fq2", + base_field_degree = 2, base_field_hdr = "tachyon/math/circle/stark/fq2.h", class_name = "Fq4", namespace = "tachyon::math::stark", diff --git a/tachyon/math/finite_fields/BUILD.bazel b/tachyon/math/finite_fields/BUILD.bazel index 01ab0b112..0d60514ff 100644 --- a/tachyon/math/finite_fields/BUILD.bazel +++ b/tachyon/math/finite_fields/BUILD.bazel @@ -72,6 +72,7 @@ tachyon_cc_library( hdrs = ["fp4.h"], deps = [ ":quadratic_extension_field", + ":quartic_extension_field", "//tachyon/math/base/gmp:gmp_util", ], ) @@ -235,6 +236,18 @@ tachyon_cc_library( ], ) +tachyon_cc_library( + name = "quartic_extension_field", + hdrs = ["quartic_extension_field.h"], + deps = [ + ":cyclotomic_multiplicative_subgroup", + "//tachyon/base/buffer:copyable", + "//tachyon/base/json", + "@com_google_absl//absl/strings", + "@com_google_absl//absl/types:span", + ], +) + tachyon_cc_library( name = "small_prime_field", hdrs = ["small_prime_field.h"], @@ -275,6 +288,7 @@ tachyon_cc_unittest( "prime_field_generator_unittest.cc", "prime_field_unittest.cc", "quadratic_extension_field_unittest.cc", + "quartic_extension_field_unittest.cc", ] + select({ "@platforms//cpu:x86_64": ["packed_prime_field_unittest.cc"], "@platforms//cpu:aarch64": ["packed_prime_field_unittest.cc"], @@ -301,6 +315,7 @@ tachyon_cc_unittest( "//tachyon/math/elliptic_curves/secp/secp256k1:fq", "//tachyon/math/elliptic_curves/secp/secp256k1:fr", "//tachyon/math/finite_fields/baby_bear", + "//tachyon/math/finite_fields/baby_bear:baby_bear4", "//tachyon/math/finite_fields/binary_fields", "//tachyon/math/finite_fields/goldilocks:goldilocks_prime_field", "//tachyon/math/finite_fields/koala_bear", diff --git a/tachyon/math/finite_fields/baby_bear/BUILD.bazel b/tachyon/math/finite_fields/baby_bear/BUILD.bazel index 0d11b0f9c..93ae5338d 100644 --- a/tachyon/math/finite_fields/baby_bear/BUILD.bazel +++ b/tachyon/math/finite_fields/baby_bear/BUILD.bazel @@ -1,5 +1,6 @@ load("//bazel:tachyon.bzl", "if_aarch64", "if_has_avx512", "if_x86_64") load("//bazel:tachyon_cc.bzl", "tachyon_avx512_defines", "tachyon_cc_library") +load("//tachyon/math/finite_fields/generator/ext_prime_field_generator:build_defs.bzl", "generate_fp4s") load("//tachyon/math/finite_fields/generator/prime_field_generator:build_defs.bzl", "generate_prime_fields") package(default_visibility = ["//visibility:public"]) @@ -14,6 +15,17 @@ generate_prime_fields( use_montgomery = True, ) +generate_fp4s( + name = "baby_bear4", + base_field = "BabyBear", + base_field_degree = 1, + base_field_hdr = "tachyon/math/finite_fields/baby_bear/baby_bear.h", + class_name = "BabyBear4", + namespace = "tachyon::math", + non_residue = ["11"], + deps = [":baby_bear"], +) + tachyon_cc_library( name = "packed_baby_bear", hdrs = ["packed_baby_bear.h"], diff --git a/tachyon/math/finite_fields/finite_field_forwards.h b/tachyon/math/finite_fields/finite_field_forwards.h index fc7292f73..950c54fb6 100644 --- a/tachyon/math/finite_fields/finite_field_forwards.h +++ b/tachyon/math/finite_fields/finite_field_forwards.h @@ -21,7 +21,7 @@ class Fp2; template class Fp3; -template +template class Fp4; template diff --git a/tachyon/math/finite_fields/fp4.h b/tachyon/math/finite_fields/fp4.h index aa892073d..a5c29634c 100644 --- a/tachyon/math/finite_fields/fp4.h +++ b/tachyon/math/finite_fields/fp4.h @@ -8,11 +8,13 @@ #include "tachyon/math/base/gmp/gmp_util.h" #include "tachyon/math/finite_fields/quadratic_extension_field.h" +#include "tachyon/math/finite_fields/quartic_extension_field.h" namespace tachyon::math { template -class Fp4 final : public QuadraticExtensionField> { +class Fp4> final + : public QuadraticExtensionField> { public: using BaseField = typename Config::BaseField; using BasePrimeField = typename Config::BasePrimeField; @@ -89,6 +91,114 @@ class Fp4 final : public QuadraticExtensionField> { } }; +template +class Fp4> final + : public QuarticExtensionField> { + public: + using BaseField = typename Config::BaseField; + using BasePrimeField = typename Config::BasePrimeField; + using FrobeniusCoefficient = typename Config::FrobeniusCoefficient; + + using CpuField = Fp4; + // TODO(chokobole): Implement Fp4Gpu + using GpuField = Fp4; + + using QuarticExtensionField>::QuarticExtensionField; + + static_assert(Config::kDegreeOverBaseField == 4); + static_assert(BaseField::ExtensionDegree() == 1); + + constexpr static uint64_t kDegreeOverBasePrimeField = 4; + + static void Init() { + Config::Init(); + // x⁴ = q = |Config::kNonResidue| + + // αᴾ = (α₀ + α₁x + α₂x² + α₃x³)ᴾ + // = α₀ᴾ + α₁ᴾxᴾ + α₂ᴾx²ᴾ + α₃ᴾx³ᴾ + // = α₀ + α₁xᴾ + α₂x²ᴾ + α₃x³ᴾ <- Fermat's little theorem + // = α₀ + α₁xᴾ⁻¹x + α₂x²ᴾ⁻²x² + α₃x³ᴾ⁻³x³ + // = α₀ + α₁(x⁴)^((P - 1) / 4) * x + α₂(x⁴)^(2 * (P - 1) / 4) * x² + + // α₃(x⁴)^(3 * (P - 1) / 4) * x³ + // = α₀ + α₁ωx + α₂ω²x² + α₃ω³x³, where ω is a quartic root of unity. + + constexpr uint64_t N = BasePrimeField::kLimbNums; + // m₁ = P + mpz_class m1; + if constexpr (BasePrimeField::Config::kModulusBits <= 32) { + m1 = mpz_class(BasePrimeField::Config::kModulus); + } else { + gmp::WriteLimbs(BasePrimeField::Config::kModulus.limbs, N, &m1); + } + +#define SET_M(d, d_prev) mpz_class m##d = m##d_prev * m1 + + // m₂ = m₁ * P = P² + SET_M(2, 1); + // m₃ = m₂ * P = P³ + SET_M(3, 2); + +#undef SET_M + +#define SET_EXP_GMP(d) mpz_class exp##d##_gmp = (m##d - 1) / mpz_class(4) + + // exp₁ = (m₁ - 1) / 4 = (P¹ - 1) / 4 + SET_EXP_GMP(1); + // exp₂ = (m₂ - 1) / 4 = (P² - 1) / 4 + SET_EXP_GMP(2); + // exp₃ = (m₃ - 1) / 4 = (P³ - 1) / 4 + SET_EXP_GMP(3); + +#undef SET_EXP_GMP + + // |kFrobeniusCoeffs[0]| = q^((P⁰ - 1) / 4) = 1 + Config::kFrobeniusCoeffs[0] = FrobeniusCoefficient::One(); +#define SET_FROBENIUS_COEFF(d) \ + BigInt exp##d; \ + gmp::CopyLimbs(exp##d##_gmp, exp##d.limbs); \ + Config::kFrobeniusCoeffs[d] = Config::kNonResidue.Pow(exp##d) + + // |kFrobeniusCoeffs[1]| = q^(exp₁) = q^((P¹ - 1) / 4) = ω + SET_FROBENIUS_COEFF(1); + // |kFrobeniusCoeffs[2]| = q^(exp₂) = q^((P² - 1) / 4) + SET_FROBENIUS_COEFF(2); + // |kFrobeniusCoeffs[3]| = q^(exp₃) = q^((P³ - 1) / 4) + SET_FROBENIUS_COEFF(3); + +#undef SET_FROBENIUS_COEFF + + // |kFrobeniusCoeffs2[0]| = q^(2 * (P⁰ - 1) / 4) = 1 + Config::kFrobeniusCoeffs2[0] = FrobeniusCoefficient::One(); +#define SET_FROBENIUS_COEFF2(d) \ + gmp::CopyLimbs(mpz_class(2) * exp##d##_gmp, exp##d.limbs); \ + Config::kFrobeniusCoeffs2[d] = Config::kNonResidue.Pow(exp##d) + + // |kFrobeniusCoeffs2[1]| = q^(2 * exp₁) = q^(2 * (P¹ - 1) / 4) = ω² + SET_FROBENIUS_COEFF2(1); + // |kFrobeniusCoeffs2[2]| = q^(2 * exp₂) = q^(2 * (P² - 1) / 4) + SET_FROBENIUS_COEFF2(2); + // |kFrobeniusCoeffs2[3]| = q^(2 * exp₃) = q^(2 * (P³ - 1) / 4) + SET_FROBENIUS_COEFF2(3); + +#undef SET_FROBENIUS_COEFF2 + + // |kFrobeniusCoeffs3[0]| = q^(3 * (P⁰ - 1) / 4) = 1 + Config::kFrobeniusCoeffs3[0] = FrobeniusCoefficient::One(); +#define SET_FROBENIUS_COEFF3(d) \ + gmp::CopyLimbs(mpz_class(3) * exp##d##_gmp, exp##d.limbs); \ + Config::kFrobeniusCoeffs3[d] = Config::kNonResidue.Pow(exp##d) + + // |kFrobeniusCoeffs3[1]| = q^(3 * exp₁) = q^(3 * (P¹ - 1) / 4) = ω³ + SET_FROBENIUS_COEFF3(1); + // |kFrobeniusCoeffs3[2]| = q^(3 * exp₂) = q^(3 * (P² - 1) / 4) + SET_FROBENIUS_COEFF3(2); + // |kFrobeniusCoeffs3[3]| = q^(3 * exp₃) = q^(3 * (P³ - 1) / 4) + SET_FROBENIUS_COEFF3(3); + +#undef SET_FROBENIUS_COEFF3 + } +}; + } // namespace tachyon::math #endif // TACHYON_MATH_FINITE_FIELDS_FP4_H_ diff --git a/tachyon/math/finite_fields/generator/ext_prime_field_generator/build_defs.bzl b/tachyon/math/finite_fields/generator/ext_prime_field_generator/build_defs.bzl index bbe3d2f51..e7a70b186 100644 --- a/tachyon/math/finite_fields/generator/ext_prime_field_generator/build_defs.bzl +++ b/tachyon/math/finite_fields/generator/ext_prime_field_generator/build_defs.bzl @@ -122,7 +122,6 @@ def generate_fp4s( _generate_ext_prime_fields( name = name, degree = 4, - base_field_degree = 2, ext_prime_field_deps = ["//tachyon/math/finite_fields:fp4"], **kwargs ) diff --git a/tachyon/math/finite_fields/generator/ext_prime_field_generator/ext_prime_field_generator.cc b/tachyon/math/finite_fields/generator/ext_prime_field_generator/ext_prime_field_generator.cc index c909e9673..0185c2fbd 100644 --- a/tachyon/math/finite_fields/generator/ext_prime_field_generator/ext_prime_field_generator.cc +++ b/tachyon/math/finite_fields/generator/ext_prime_field_generator/ext_prime_field_generator.cc @@ -112,7 +112,8 @@ int GenerationConfig::GenerateConfigHdr() const { } replacements["%{frobenius_coefficient}"] = - (degree == 4 || (degree == 6 && base_field_degree == 3) || degree == 12) + ((degree == 4 && base_field_degree == 2) || + (degree == 6 && base_field_degree == 3) || degree == 12) ? "typename BaseField::BaseField" : "BaseField"; @@ -122,7 +123,9 @@ int GenerationConfig::GenerateConfigHdr() const { std::vector tpl_lines = absl::StrSplit(tpl_content, '\n'); RemoveOptionalLines(tpl_lines, "FrobeniusCoefficient2", - degree_over_base_field == 3); + degree_over_base_field >= 3); + RemoveOptionalLines(tpl_lines, "FrobeniusCoefficient3", + degree_over_base_field >= 4); tpl_content = absl::StrJoin(tpl_lines, "\n"); std::string content = absl::StrReplaceAll(tpl_content, replacements); diff --git a/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl b/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl index b49ace189..1dceda156 100644 --- a/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl +++ b/tachyon/math/finite_fields/generator/ext_prime_field_generator/fq.h.tpl @@ -18,6 +18,9 @@ class %{class}Config { %{if FrobeniusCoefficient2} static FrobeniusCoefficient kFrobeniusCoeffs2[%{frobenius_coeffs_size}]; %{endif FrobeniusCoefficient2} +%{if FrobeniusCoefficient3} + static FrobeniusCoefficient kFrobeniusCoeffs3[%{frobenius_coeffs_size}]; +%{endif FrobeniusCoefficient3} constexpr static bool kNonResidueIsMinusOne = %{non_residue_is_minus_one}; constexpr static uint64_t kDegreeOverBaseField = %{degree_over_base_field}; @@ -42,6 +45,10 @@ typename %{class}Config::FrobeniusCoefficient %{class}Config typename %{class}Config::FrobeniusCoefficient %{class}Config::kFrobeniusCoeffs2[%{frobenius_coeffs_size}]; %{endif FrobeniusCoefficient2} +%{if FrobeniusCoefficient3} +template +typename %{class}Config::FrobeniusCoefficient %{class}Config::kFrobeniusCoeffs3[%{frobenius_coeffs_size}]; +%{endif FrobeniusCoefficient3} using %{class} = Fp%{degree}<%{class}Config<%{base_field}>>; } // namespace %{namespace} diff --git a/tachyon/math/finite_fields/koala_bear/BUILD.bazel b/tachyon/math/finite_fields/koala_bear/BUILD.bazel index 9a8a6a842..38238dbef 100644 --- a/tachyon/math/finite_fields/koala_bear/BUILD.bazel +++ b/tachyon/math/finite_fields/koala_bear/BUILD.bazel @@ -1,6 +1,6 @@ load("//bazel:tachyon.bzl", "if_aarch64", "if_has_avx512", "if_x86_64") load("//bazel:tachyon_cc.bzl", "tachyon_avx512_defines", "tachyon_cc_library") -load("//tachyon/math/finite_fields/generator/ext_prime_field_generator:build_defs.bzl", "generate_fp2s") +load("//tachyon/math/finite_fields/generator/ext_prime_field_generator:build_defs.bzl", "generate_fp2s", "generate_fp4s") load("//tachyon/math/finite_fields/generator/prime_field_generator:build_defs.bzl", "generate_prime_fields") package(default_visibility = ["//visibility:public"]) @@ -25,6 +25,17 @@ generate_fp2s( deps = [":koala_bear"], ) +generate_fp4s( + name = "koala_bear4", + base_field = "KoalaBear", + base_field_degree = 1, + base_field_hdr = "tachyon/math/finite_fields/koala_bear/koala_bear.h", + class_name = "KoalaBear4", + namespace = "tachyon::math", + non_residue = ["3"], + deps = [":koala_bear"], +) + tachyon_cc_library( name = "packed_koala_bear", hdrs = ["packed_koala_bear.h"], diff --git a/tachyon/math/finite_fields/quartic_extension_field.h b/tachyon/math/finite_fields/quartic_extension_field.h new file mode 100644 index 000000000..9edc0b72a --- /dev/null +++ b/tachyon/math/finite_fields/quartic_extension_field.h @@ -0,0 +1,646 @@ +#ifndef TACHYON_MATH_FINITE_FIELDS_QUARTIC_EXTENSION_FIELD_H_ +#define TACHYON_MATH_FINITE_FIELDS_QUARTIC_EXTENSION_FIELD_H_ + +#include +#include +#include + +#include "absl/strings/substitute.h" +#include "absl/types/span.h" + +#include "tachyon/base/buffer/copyable.h" +#include "tachyon/base/json/json.h" +#include "tachyon/math/finite_fields/cyclotomic_multiplicative_subgroup.h" + +namespace tachyon { +namespace math { + +template +class QuarticExtensionField : public CyclotomicMultiplicativeSubgroup { + public: + using Config = typename FiniteField::Config; + using BaseField = typename Config::BaseField; + using BasePrimeField = typename Config::BasePrimeField; + + constexpr QuarticExtensionField() = default; + constexpr QuarticExtensionField(const BaseField& c0, const BaseField& c1, + const BaseField& c2, const BaseField& c3) + : c0_(c0), c1_(c1), c2_(c2), c3_(c3) {} + constexpr QuarticExtensionField(BaseField&& c0, BaseField&& c1, + BaseField&& c2, BaseField&& c3) + : c0_(std::move(c0)), + c1_(std::move(c1)), + c2_(std::move(c2)), + c3_(std::move(c3)) {} + + constexpr static Derived Zero() { + return {BaseField::Zero(), BaseField::Zero(), BaseField::Zero(), + BaseField::Zero()}; + } + + constexpr static Derived One() { + return {BaseField::One(), BaseField::Zero(), BaseField::Zero(), + BaseField::Zero()}; + } + + static Derived Random() { + return {BaseField::Random(), BaseField::Random(), BaseField::Random(), + BaseField::Random()}; + } + + static Derived FromBasePrimeFields( + absl::Span prime_fields) { + CHECK_EQ(prime_fields.size(), ExtensionDegree()); + constexpr size_t kBaseFieldDegree = BaseField::ExtensionDegree(); + if constexpr (kBaseFieldDegree == 1) { + return Derived(prime_fields[0], prime_fields[1], prime_fields[2], + prime_fields[3]); + } else { + BaseField c0 = BaseField::FromBasePrimeFields( + prime_fields.subspan(0, kBaseFieldDegree)); + prime_fields.remove_prefix(kBaseFieldDegree); + BaseField c1 = BaseField::FromBasePrimeFields( + prime_fields.subspan(0, kBaseFieldDegree)); + prime_fields.remove_prefix(kBaseFieldDegree); + BaseField c2 = BaseField::FromBasePrimeFields( + prime_fields.subspan(0, kBaseFieldDegree)); + prime_fields.remove_prefix(kBaseFieldDegree); + BaseField c3 = BaseField::FromBasePrimeFields( + prime_fields.subspan(kBaseFieldDegree)); + return Derived(std::move(c0), std::move(c1), std::move(c2), + std::move(c3)); + } + } + + constexpr bool IsZero() const { + return c0_.IsZero() && c1_.IsZero() && c2_.IsZero() && c3_.IsZero(); + } + + constexpr bool IsOne() const { + return c0_.IsOne() && c1_.IsZero() && c2_.IsZero() && c3_.IsZero(); + } + + constexpr static uint64_t ExtensionDegree() { + return 4 * BaseField::ExtensionDegree(); + } + + // Calculate the norm of an element with respect to |BaseField|. + // The norm maps an element |a| in the extension field Fqᵐ to an element + // in the |BaseField| Fq. |a.Norm() = a * a^q * a^q² * a^q³| + constexpr BaseField Norm() const { + // w.r.t to |BaseField|, we need the 0th, 1st, 2nd & 3rd powers of q. + // Since Frobenius coefficients on the towered extensions are + // indexed w.r.t. to |BasePrimeField|, we need to calculate the correct + // index. + // NOTE(chokobole): This assumes that |BaseField::ExtensionDegree()| + // never overflows even on a 32-bit machine. + size_t index_multiplier = size_t{BaseField::ExtensionDegree()}; + Derived self_to_p = static_cast(*this); + self_to_p.FrobeniusMapInPlace(index_multiplier); + Derived self_to_p2 = static_cast(*this); + self_to_p2.FrobeniusMapInPlace(2 * index_multiplier); + Derived self_to_p3 = static_cast(*this); + self_to_p3.FrobeniusMapInPlace(3 * index_multiplier); + self_to_p *= (self_to_p2 * self_to_p3 * static_cast(*this)); + // NOTE(chokobole): The |CHECK()| below is not device code. + // See https://github.com/kroma-network/tachyon/issues/76 + CHECK(self_to_p.c1().IsZero() && self_to_p.c2().IsZero() && + self_to_p.c3().IsZero()); + return self_to_p.c0(); + } + + constexpr Derived& FrobeniusMapInPlace(uint64_t exponent) { + c0_.FrobeniusMapInPlace(exponent); + c1_.FrobeniusMapInPlace(exponent); + c2_.FrobeniusMapInPlace(exponent); + c3_.FrobeniusMapInPlace(exponent); + c1_ *= + Config::kFrobeniusCoeffs[exponent % Config::kDegreeOverBasePrimeField]; + c2_ *= + Config::kFrobeniusCoeffs2[exponent % Config::kDegreeOverBasePrimeField]; + c3_ *= + Config::kFrobeniusCoeffs3[exponent % Config::kDegreeOverBasePrimeField]; + return *static_cast(this); + } + + std::string ToString() const { + return absl::Substitute("($0, $1, $2, $3)", c0_.ToString(), c1_.ToString(), + c2_.ToString(), c3_.ToString()); + } + + std::string ToHexString(bool pad_zero = false) const { + return absl::Substitute("($0, $1, $2, $3)", c0_.ToHexString(pad_zero), + c1_.ToHexString(pad_zero), + c2_.ToHexString(pad_zero), + c3_.ToHexString(pad_zero)); + } + + constexpr const BaseField& c0() const { return c0_; } + constexpr const BaseField& c1() const { return c1_; } + constexpr const BaseField& c2() const { return c2_; } + constexpr const BaseField& c3() const { return c3_; } + + constexpr bool operator==(const Derived& other) const { + return c0_ == other.c0_ && c1_ == other.c1_ && c2_ == other.c2_ && + c3_ == other.c3_; + } + + constexpr bool operator!=(const Derived& other) const { + return !operator==(other); + } + + constexpr bool operator<(const Derived& other) const { + if (c3_ == other.c3_) { + if (c2_ == other.c2_) { + if (c1_ == other.c1_) return c0_ < other.c0_; + return c1_ < other.c1_; + } + return c2_ < other.c2_; + } + return c3_ < other.c3_; + } + + constexpr bool operator>(const Derived& other) const { + if (c3_ == other.c3_) { + if (c2_ == other.c2_) { + if (c1_ == other.c1_) return c0_ > other.c0_; + return c1_ > other.c1_; + } + return c2_ > other.c2_; + } + return c3_ > other.c3_; + } + + constexpr bool operator<=(const Derived& other) const { + return !operator>(other); + } + + constexpr bool operator>=(const Derived& other) const { + return !operator<(other); + } + + // AdditiveSemigroup methods + constexpr Derived Add(const Derived& other) const { + return { + c0_ + other.c0_, + c1_ + other.c1_, + c2_ + other.c2_, + c3_ + other.c3_, + }; + } + + constexpr Derived& AddInPlace(const Derived& other) { + c0_ += other.c0_; + c1_ += other.c1_; + c2_ += other.c2_; + c3_ += other.c3_; + return *static_cast(this); + } + + constexpr Derived DoubleImpl() const { + return { + c0_.Double(), + c1_.Double(), + c2_.Double(), + c3_.Double(), + }; + } + + constexpr Derived& DoubleImplInPlace() { + c0_.DoubleInPlace(); + c1_.DoubleInPlace(); + c2_.DoubleInPlace(); + c3_.DoubleInPlace(); + return *static_cast(this); + } + + // AdditiveGroup methods + constexpr Derived Sub(const Derived& other) const { + return { + c0_ - other.c0_, + c1_ - other.c1_, + c2_ - other.c2_, + c3_ - other.c3_, + }; + } + + constexpr Derived& SubInPlace(const Derived& other) { + c0_ -= other.c0_; + c1_ -= other.c1_; + c2_ -= other.c2_; + c3_ -= other.c3_; + return *static_cast(this); + } + + constexpr Derived Negate() const { + return { + -c0_, + -c1_, + -c2_, + -c3_, + }; + } + + constexpr Derived& NegateInPlace() { + c0_.NegateInPlace(); + c1_.NegateInPlace(); + c2_.NegateInPlace(); + c3_.NegateInPlace(); + return *static_cast(this); + } + + // MultiplicativeSemigroup methods + constexpr Derived Mul(const Derived& other) const { + Derived ret{}; + DoMul(*static_cast(this), other, ret); + return ret; + } + + constexpr Derived& MulInPlace(const Derived& other) { + DoMul(*static_cast(this), other, + *static_cast(this)); + return *static_cast(this); + } + + constexpr Derived Mul(const BaseField& element) const { + return { + c0_ * element, + c1_ * element, + c2_ * element, + c3_ * element, + }; + } + + constexpr Derived& MulInPlace(const BaseField& element) { + c0_ *= element; + c1_ *= element; + c2_ *= element; + c3_ *= element; + return *static_cast(this); + } + + constexpr Derived SquareImpl() const { + Derived ret{}; + DoSquareImpl(*static_cast(this), ret); + return ret; + } + + constexpr Derived& SquareImplInPlace() { + DoSquareImpl(*static_cast(this), + *static_cast(this)); + return *static_cast(this); + } + + // MultiplicativeGroup methods + constexpr std::optional Inverse() const { + Derived ret{}; + if (LIKELY(DoInverse(*static_cast(this), ret))) { + return ret; + } + LOG_IF_NOT_GPU(ERROR) << "Inverse of zero attempted"; + return std::nullopt; + } + + [[nodiscard]] constexpr std::optional InverseInPlace() { + if (LIKELY(DoInverse(*static_cast(this), + *static_cast(this)))) { + return static_cast(this); + } + LOG_IF_NOT_GPU(ERROR) << "Inverse of zero attempted"; + return std::nullopt; + } + + protected: + constexpr static void DoMul(const Derived& a, const Derived& b, Derived& c) { + // clang-format off + // (a.c0, a.c1, a.c2, a.c3) * (b.c0, b.c1, b.c2, b.c3) + // = (a.c0 + a.c1 * x + a.c2 * x² + a.c3 * x³) * (b.c0 + b.c1 * x + b.c2 * x² + b.c3 * x³) + // = a.c0 * b.c0 + (a.c0 * b.c1 + a.c1 * b.c0) * x + (a.c0 * b.c2 + a.c1 * b.c1 + a.c2 * b.c0) * x² + + // (a.c0 * b.c3 + a.c1 * b.c2 + a.c2 * b.c1 + a.c3 * b.c0) * x³ + (a.c1 * b.c3 + a.c2 * b.c2 + a.c3 * b.c1) * x⁴ + + // (a.c2 * b.c3 + a.c3 * b.c2) * x⁵ + a.c3 * b.c3 * x⁶ + // = a.c0 * b.c0 + (a.c1 * b.c3 + a.c2 * b.c2 + a.c3 * b.c1) * x⁴ + + // (a.c0 * b.c1 + a.c1 * b.c0) * x + (a.c2 * b.c3 + a.c3 * b.c2) * x⁵ + + // (a.c0 * b.c2 + a.c1 * b.c1 + a.c2 * b.c0) * x² + a.c3 * b.c3 * x⁶ + + // (a.c0 * b.c3 + a.c1 * b.c2 + a.c2 * b.c1 + a.c3 * b.c0) * x³ + // = a.c0 * b.c0 + (a.c1 * b.c3 + a.c2 * b.c2 + a.c3 * b.c1) * q + + // (a.c0 * b.c1 + a.c1 * b.c0) * x + (a.c2 * b.c3 + a.c3 * b.c2) * q * x + + // (a.c0 * b.c2 + a.c1 * b.c1 + a.c2 * b.c0) * x² + a.c3 * b.c3 * q * x² + + // (a.c0 * b.c3 + a.c1 * b.c2 + a.c2 * b.c1 + a.c3 * b.c0) * x³ + // = (a.c0 * b.c0 + (a.c1 * b.c3 + a.c2 * b.c2 + a.c3 * b.c1) * q, + // a.c0 * b.c1 + a.c1 * b.c0 + (a.c2 * b.c3 + a.c3 * b.c2) * q, + // a.c0 * b.c2 + a.c1 * b.c1 + a.c2 * b.c0 + a.c3 * b.c3 * q, + // a.c0 * b.c3 + a.c1 * b.c2 + a.c2 * b.c1 + a.c3 * b.c0) + // where q is |Config::kNonResidue|. + + // See https://eprint.iacr.org/2006/471.pdf + // Devegili OhEig Scott Dahab --- Multiplication and Squaring on AbstractPairing-Friendly Fields.pdf; Section 5.2 + // clang-format on + + constexpr BaseField kInv2 = *BaseField(2).Inverse(); + constexpr BaseField kInv3 = *BaseField(3).Inverse(); + constexpr BaseField kInv4 = *BaseField(4).Inverse(); + constexpr BaseField kInv6 = *BaseField(6).Inverse(); + constexpr BaseField kInv12 = *BaseField(12).Inverse(); + constexpr BaseField kInv20 = *BaseField(20).Inverse(); + constexpr BaseField kInv24 = *BaseField(24).Inverse(); + constexpr BaseField kInv30 = *BaseField(30).Inverse(); + constexpr BaseField kInv120 = *BaseField(120).Inverse(); + constexpr BaseField kNeg5 = -BaseField(5); + constexpr BaseField kNegInv2 = -kInv2; + constexpr BaseField kNegInv3 = -kInv3; + constexpr BaseField kNegInv4 = -kInv4; + constexpr BaseField kNegInv6 = -kInv6; + constexpr BaseField kNegInv12 = -kInv12; + constexpr BaseField kNegInv24 = -kInv24; + constexpr BaseField kNegInv120 = -kInv120; + + // h1 = 2 * a.c1 + BaseField h1 = a.c1_.Double(); + // h2 = 4 * a.c2 + BaseField h2 = a.c2_.Double(); + h2.DoubleInPlace(); + // h3 = 8 * a.c3 + BaseField h3 = a.c3_.Double(); + h3.DoubleInPlace().DoubleInPlace(); + // h4 = 2 * b.c1 + BaseField h4 = b.c1_.Double(); + // h5 = 4 * b.c2 + BaseField h5 = b.c2_.Double(); + h5.DoubleInPlace(); + // h6 = 8 * b.c3 + BaseField h6 = b.c3_.Double(); + h6.DoubleInPlace().DoubleInPlace(); + + // v0 = a.c0 * b.c0 + BaseField v0 = a.c0_ * b.c0_; + // v1 = (a.c0 + a.c1 + a.c2 + a.c3) * (b.c0 + b.c1 + b.c2 + b.c3) + BaseField v1 = + (a.c0_ + a.c1_ + a.c2_ + a.c3_) * (b.c0_ + b.c1_ + b.c2_ + b.c3_); + // v2 = (a.c0 - a.c1 + a.c2 - a.c3) * (b.c0 - b.c1 + b.c2 - b.c3) + BaseField v2 = + (a.c0_ - a.c1_ + a.c2_ - a.c3_) * (b.c0_ - b.c1_ + b.c2_ - b.c3_); + // v3 = (a.c0 + 2 * a.c1 + 4 * a.c2 + 8 * a.c3) * + // (b.c0 + 2 * b.c1 + 4 * b.c2 + 8 * b.c3) + BaseField v3 = (a.c0_ + h1 + h2 + h3) * (b.c0_ + h4 + h5 + h6); + // v4 = (a.c0 - 2 * a.c1 + 4 * a.c2 - 8 * a.c3) * + // (b.c0 - 2 * b.c1 + 4 * b.c2 - 8 * b.c3) + BaseField v4 = (a.c0_ - h1 + h2 - h3) * (b.c0_ - h4 + h5 - h6); + // h1 = 3 * a.c1 + h1 += a.c1_; + // h2 = 9 * a.c2 + h2.DoubleInPlace().AddInPlace(a.c2_); + // h3 = 27 * a.c3 + h3 += a.c3_; + h3 += h3.Double(); + // h4 = 3 * b.c1 + h4 += b.c1_; + // h5 = 9 * b.c2 + h5.DoubleInPlace().AddInPlace(b.c2_); + // h6 = 27 * b.c3 + h6 += b.c3_; + h6 += h6.Double(); + // v5 = (a.c0 + 3 * a.c1 + 9 * a.c2 + 27 * a.c3) * + // (b.c0 + 3 * b.c1 + 9 * b.c2 + 27 * b.c3) + BaseField v5 = (a.c0_ + h1 + h2 + h3) * (b.c0_ + h4 + h5 + h6); + // v6 = a.c3 * b.c3 + BaseField v6 = a.c3_ * b.c3_; + + // v0_5 = 5 * v0 + BaseField v0_5 = v0.Double(); + v0_5.DoubleInPlace().AddInPlace(v0); + // v6_3 = 3 * v6 + BaseField v6_3 = v6.Double(); + v6_3 += v6; + + // clang-format off + // c.c0 = v0 + + // q * ((1 / 4) * v0 - (1 / 6) * (v1 + v2) + (1 / 24) * (v3 + v4) - 5 * v6) + c.c0_ = v0 + + Config::MulByNonResidue(kInv4 * v0 + kNegInv6 * (v1 + v2) + kInv24 * (v3 + v4) + kNeg5 * v6); + // c.c1 = -(1 / 3) * v0 + v1 - (1 / 2) * v2 - (1 / 4) * v3 + (1 / 20) * v4 + (1 / 30) * v5 - 12 * v6 + + // q * (-(1 / 12) * (v0 - v1) + (1 / 24) * (v2 - v3) - (1 / 120) * (v4 - v5) - 3 * v6) + c.c1_ = kNegInv3 * v0 + v1 + kNegInv2 * v2 + kNegInv4 * v3 + kInv20 * v4 + kInv30 * v5 - v6_3.Double().Double() + + Config::MulByNonResidue(kNegInv12 * (v0 - v1) + kInv24 * (v2 - v3) + kNegInv120 * (v4 - v5) - v6_3); + // c.c2 = -(5 / 4) * v0 + (2 / 3) * (v1 + v2) - (1 / 24) * (v3 + v4) + 4 * v6 + + // q * v6 + c.c2_ = kNegInv4 * v0_5 + kInv3 * (v1 + v2).Double() + kNegInv24 * (v3 + v4) + v6.Double().Double() + + Config::MulByNonResidue(v6); + // c.c3 = (1 / 12) * (5 * v0 - 7 * v1) - (1 / 24) * (v2 - 7 * v3 + v4 + v5) + 15 * v6 + c.c3_ = kInv12 * (v0_5 - v1.Double().Double().Double() + v1) + kNegInv24 * (v2 - v3.Double().Double().Double() + v3 + v4 + v5) + v6_3.Double().Double() + v6_3; + // clang-format on + } + + constexpr static void DoSquareImpl(const Derived& a, Derived& b) { + // clang-format off + // (c0, c1, c2, c3)² + // = (c0 + c1 * x + c2 * x² + c3 * x³)² + // = c0² + 2 * c0 * c1 * x + (c1² + 2 * c0 * c2) * x² + 2 * (c0 * c3 + c1 * c2) * x³ + (c2² + 2 * c1 * c3) * x⁴ + 2 * c2 * c3 * x⁵ + c3 * x⁶ + // = c0² + (c2² + 2 * c1 * c3) * x⁴ + 2 * c0 * c1 * x + 2 * c2 * c3 * x⁵ + (c1² + 2 * c0 * c2) * x² + c3 * x⁶ + 2 * (c0 * c3 + c1 * c2) * x³ + // = c0² + (c2² + 2 * c1 * c3) * q + 2 * (c0 * c1 + c2 * c3 * q) * x + (c1² + 2 * c0 * c2 + c3 * q) * x² + 2 * (c0 * c3 + c1 * c2) * x³ + // = (c0² + (c2² + 2 * c1 * c3) * q, 2 * (c0 * c1 + c2 * c3 * q), c1² + 2 * c0 * c2 + c3 * q, 2 * (c0 * c3 + c1 * c2)) + // where q is |Config::kNonResidue|. + + // See https://eprint.iacr.org/2006/471.pdf + // Devegili OhEig Scott Dahab --- Multiplication and Squaring on AbstractPairing-Friendly Fields.pdf; Section 5 + // clang-format on + + constexpr BaseField kInv2 = *BaseField(2).Inverse(); + constexpr BaseField kInv3 = *BaseField(3).Inverse(); + constexpr BaseField kInv4 = *BaseField(4).Inverse(); + constexpr BaseField kInv6 = *BaseField(6).Inverse(); + constexpr BaseField kInv12 = *BaseField(12).Inverse(); + constexpr BaseField kInv20 = *BaseField(20).Inverse(); + constexpr BaseField kInv24 = *BaseField(24).Inverse(); + constexpr BaseField kInv30 = *BaseField(30).Inverse(); + constexpr BaseField kInv120 = *BaseField(120).Inverse(); + constexpr BaseField kNeg5 = -BaseField(5); + constexpr BaseField kNegInv2 = -kInv2; + constexpr BaseField kNegInv3 = -kInv3; + constexpr BaseField kNegInv4 = -kInv4; + constexpr BaseField kNegInv6 = -kInv6; + constexpr BaseField kNegInv12 = -kInv12; + constexpr BaseField kNegInv24 = -kInv24; + constexpr BaseField kNegInv120 = -kInv120; + + // h1 = 2 * c1 + BaseField h1 = a.c1_.Double(); + // h2 = 4 * c2 + BaseField h2 = a.c2_.Double(); + h2.DoubleInPlace(); + // h3 = 8 * c3 + BaseField h3 = a.c3_.Double(); + h3.DoubleInPlace().DoubleInPlace(); + + // v0 = c0² + BaseField v0 = a.c0_.Square(); + // v1 = (c0 + c1 + c2 + c3)² + BaseField v1 = (a.c0_ + a.c1_ + a.c2_ + a.c3_).Square(); + // v2 = (c0 - c1 + c2 - c3)² + BaseField v2 = (a.c0_ - a.c1_ + a.c2_ - a.c3_).Square(); + // v3 = (c0 + 2 * c1 + 4 * c2 + 8 * c3)² + BaseField v3 = (a.c0_ + h1 + h2 + h3).Square(); + // v4 = (c0 - 2 * c1 + 4 * c2 - 8 * c3)² + BaseField v4 = (a.c0_ - h1 + h2 - h3).Square(); + // h1 = 3 * c1 + h1 += a.c1_; + // h2 = 9 * c2 + h2.DoubleInPlace().AddInPlace(a.c2_); + // h3 = 27 * c3 + h3 += a.c3_; + h3 += h3.Double(); + // v5 = (c0 + 3 * c1 + 9 * c2 + 27 * c3)² + BaseField v5 = (a.c0_ + h1 + h2 + h3).Square(); + // v6 = c3² + BaseField v6 = a.c3_.Square(); + + // v0_5 = 5 * v0 + BaseField v0_5 = v0.Double(); + v0_5.DoubleInPlace().AddInPlace(v0); + // v6_3 = 3 * v6 + BaseField v6_3 = v6.Double(); + v6_3 += v6; + + // clang-format off + // b.c0 = v0 + + // q * ((1 / 4) * v0 - (1 / 6) * (v1 + v2) + (1 / 24) * (v3 + v4) - 5 * v6) + b.c0_ = v0 + + Config::MulByNonResidue(kInv4 * v0 + kNegInv6 * (v1 + v2) + kInv24 * (v3 + v4) + kNeg5 * v6); + // b.c1 = -(1 / 3) * v0 + v1 - (1 / 2) * v2 - (1 / 4) * v3 + (1 / 20) * v4 + (1 / 30) * v5 - 12 * v6 + + // q * (-(1 / 12) * (v0 - v1) + (1 / 24) * (v2 - v3) - (1 / 120) * (v4 - v5) - 3 * v6) + b.c1_ = kNegInv3 * v0 + v1 + kNegInv2 * v2 + kNegInv4 * v3 + kInv20 * v4 + kInv30 * v5 - v6_3.Double().Double() + + Config::MulByNonResidue(kNegInv12 * (v0 - v1) + kInv24 * (v2 - v3) + kNegInv120 * (v4 - v5) - v6_3); + // b.c2 = -(5 / 4) * v0 + (2 / 3) * (v1 + v2) - (1 / 24) * (v3 + v4) + 4 * v6 + + // q * v6 + b.c2_ = kNegInv4 * v0_5 + kInv3 * (v1 + v2).Double() + kNegInv24 * (v3 + v4) + v6.Double().Double() + + Config::MulByNonResidue(v6); + // b.c3 = (1 / 12) * (5 * v0 - 7 * v1) - (1 / 24) * (v2 - 7 * v3 + v4 + v5) + 15 * v6 + b.c3_ = kInv12 * (v0_5 - v1.Double().Double().Double() + v1) + kNegInv24 * (v2 - v3.Double().Double().Double() + v3 + v4 + v5) + v6_3.Double().Double() + v6_3; + // clang-format on + } + + [[nodiscard]] constexpr static bool DoInverse(const Derived& a, Derived& b) { + if (UNLIKELY(a.IsZero())) { + LOG_IF_NOT_GPU(ERROR) << "Inverse of zero attempted"; + return false; + } + + // See Algorithm 11.3.4 in Handbook of Elliptic and Hyperelliptic Curve + // Cryptography. + // Compute aʳ⁻¹, where r = (p⁴ - 1) / (p - 1) = p³ + p² + p + 1 + size_t index_multiplier = size_t{BaseField::ExtensionDegree()}; + // f = a^{p³ + p² + p} + Derived a_to_r_minus_1 = a; + a_to_r_minus_1.FrobeniusMapInPlace(index_multiplier); + Derived a_to_p2 = a; + a_to_p2.FrobeniusMapInPlace(2 * index_multiplier); + Derived a_to_p3 = a; + a_to_p3.FrobeniusMapInPlace(3 * index_multiplier); + a_to_r_minus_1 *= (a_to_p2 * a_to_p3); + + // Since aʳ, which is |Norm()|, is in the base field, computing the constant + // is enough. + BaseField a_to_r = BaseField::Zero(); + a_to_r += a.c1_ * a_to_r_minus_1.c3_; + a_to_r += a.c2_ * a_to_r_minus_1.c2_; + a_to_r += a.c3_ * a_to_r_minus_1.c1_; + a_to_r = Config::MulByNonResidue(a_to_r); + a_to_r += a.c0_ * a_to_r_minus_1.c0_; + + // a⁻¹ = aʳ⁻¹ * a⁻ʳ + b = a_to_r_minus_1 * *a_to_r.Inverse(); + return true; + } + + // c = c0_ + c1_ * X + c2_ * X² + c3_ * X³ + BaseField c0_; + BaseField c1_; + BaseField c2_; + BaseField c3_; +}; + +template < + typename BaseField, typename Derived, + std::enable_if_t>* = + nullptr> +Derived operator*(const BaseField& element, + const QuarticExtensionField& f) { + return static_cast(f) * element; +} + +} // namespace math + +namespace base { + +template +class Copyable, Derived>>> { + public: + static bool WriteTo( + const math::QuarticExtensionField& quadratic_extension_field, + Buffer* buffer) { + return buffer->WriteMany( + quadratic_extension_field.c0(), quadratic_extension_field.c1(), + quadratic_extension_field.c2(), quadratic_extension_field.c3()); + } + + static bool ReadFrom( + const ReadOnlyBuffer& buffer, + math::QuarticExtensionField* quadratic_extension_field) { + typename Derived::BaseField c0; + typename Derived::BaseField c1; + typename Derived::BaseField c2; + typename Derived::BaseField c3; + if (!buffer.ReadMany(&c0, &c1, &c2, &c3)) return false; + + *quadratic_extension_field = math::QuarticExtensionField( + std::move(c0), std::move(c1), std::move(c2), std::move(c3)); + return true; + } + + static size_t EstimateSize( + const math::QuarticExtensionField& quadratic_extension_field) { + return base::EstimateSize( + quadratic_extension_field.c0(), quadratic_extension_field.c1(), + quadratic_extension_field.c2(), quadratic_extension_field.c3()); + } +}; + +template +class RapidJsonValueConverter< + Derived, std::enable_if_t, Derived>>> { + public: + using BaseField = typename math::QuarticExtensionField::BaseField; + + template + static rapidjson::Value From( + const math::QuarticExtensionField& value, Allocator& allocator) { + rapidjson::Value object(rapidjson::kObjectType); + AddJsonElement(object, "c0", value.c0(), allocator); + AddJsonElement(object, "c1", value.c1(), allocator); + AddJsonElement(object, "c2", value.c2(), allocator); + AddJsonElement(object, "c3", value.c3(), allocator); + return object; + } + + static bool To(const rapidjson::Value& json_value, std::string_view key, + math::QuarticExtensionField* value, + std::string* error) { + BaseField c0; + BaseField c1; + BaseField c2; + BaseField c3; + if (!ParseJsonElement(json_value, "c0", &c0, error)) return false; + if (!ParseJsonElement(json_value, "c1", &c1, error)) return false; + if (!ParseJsonElement(json_value, "c2", &c2, error)) return false; + if (!ParseJsonElement(json_value, "c3", &c3, error)) return false; + *value = math::QuarticExtensionField(std::move(c0), std::move(c1), + std::move(c2), std::move(c3)); + return true; + } +}; + +} // namespace base +} // namespace tachyon + +#endif // TACHYON_MATH_FINITE_FIELDS_QUARTIC_EXTENSION_FIELD_H_ diff --git a/tachyon/math/finite_fields/quartic_extension_field_unittest.cc b/tachyon/math/finite_fields/quartic_extension_field_unittest.cc new file mode 100644 index 000000000..de3cc4a65 --- /dev/null +++ b/tachyon/math/finite_fields/quartic_extension_field_unittest.cc @@ -0,0 +1,204 @@ +#include + +#include "gtest/gtest.h" + +#include "tachyon/math/finite_fields/baby_bear/baby_bear4.h" +#include "tachyon/math/finite_fields/test/finite_field_test.h" + +namespace tachyon::math { + +namespace { + +using F4 = BabyBear4; +using F = BabyBear; + +class QuaticExtensionFieldTest : public FiniteFieldTest {}; + +} // namespace + +TEST_F(QuaticExtensionFieldTest, Zero) { + EXPECT_TRUE(F4::Zero().IsZero()); + EXPECT_FALSE(F4::One().IsZero()); +} + +TEST_F(QuaticExtensionFieldTest, One) { + EXPECT_TRUE(F4::One().IsOne()); + EXPECT_FALSE(F4::Zero().IsOne()); +} + +TEST_F(QuaticExtensionFieldTest, Random) { + bool success = false; + F4 r = F4::Random(); + for (size_t i = 0; i < 100; ++i) { + if (r != F4::Random()) { + success = true; + break; + } + } + EXPECT_TRUE(success); +} + +TEST_F(QuaticExtensionFieldTest, Norm) { + constexpr static uint32_t kModulus = BabyBear::Config::kModulus; + F4 r = F4::Random(); + F4 r_to_p = r.Pow(kModulus); + F4 r_to_p2 = r_to_p.Pow(kModulus); + F4 r_to_p3 = r_to_p2.Pow(kModulus); + EXPECT_EQ(r.Norm(), (r * r_to_p * r_to_p2 * r_to_p3).c0()); +} + +TEST_F(QuaticExtensionFieldTest, EqualityOperators) { + F4 f(F(3), F(4), F(5), F(6)); + F4 f2(F(4), F(4), F(5), F(6)); + EXPECT_NE(f, f2); + F4 f3(F(4), F(3), F(5), F(6)); + EXPECT_NE(f2, f3); + + F4 f4(F(3), F(4), F(5), F(7)); + EXPECT_NE(f, f4); + + F4 f5(F(3), F(4), F(5), F(6)); + EXPECT_EQ(f, f5); +} + +TEST_F(QuaticExtensionFieldTest, ComparisonOperator) { + F4 f(F(3), F(4), F(5), F(6)); + F4 f2(F(4), F(4), F(5), F(6)); + EXPECT_LT(f, f2); + EXPECT_LE(f, f2); + EXPECT_GT(f2, f); + EXPECT_GE(f2, f); + + F4 f3(F(4), F(3), F(5), F(6)); + F4 f4(F(3), F(4), F(5), F(6)); + EXPECT_LT(f3, f4); + EXPECT_LE(f3, f4); + EXPECT_GT(f4, f3); + EXPECT_GE(f4, f3); + + F4 f5(F(4), F(5), F(6), F(3)); + F4 f6(F(3), F(2), F(6), F(5)); + EXPECT_LT(f5, f6); + EXPECT_LE(f5, f6); + EXPECT_GT(f6, f5); + EXPECT_GE(f6, f5); +} + +TEST_F(QuaticExtensionFieldTest, AdditiveOperators) { + struct { + F4 a; + F4 b; + F4 sum; + F4 amb; + F4 bma; + } tests[] = { + { + {F(1), F(2), F(3), F(4)}, + {F(3), F(5), F(6), F(8)}, + {F(4), F(7), F(9), F(12)}, + {-F(2), -F(3), -F(3), -F(4)}, + {F(2), F(3), F(3), F(4)}, + }, + }; + + for (const auto& test : tests) { + EXPECT_EQ(test.a + test.b, test.sum); + EXPECT_EQ(test.b + test.a, test.sum); + EXPECT_EQ(test.a - test.b, test.amb); + EXPECT_EQ(test.b - test.a, test.bma); + + F4 tmp = test.a; + tmp += test.b; + EXPECT_EQ(tmp, test.sum); + tmp -= test.b; + EXPECT_EQ(tmp, test.a); + } +} + +TEST_F(QuaticExtensionFieldTest, AdditiveGroupOperators) { + F4 f(F(3), F(4), F(5), F(6)); + F4 f_neg(-F(3), -F(4), -F(5), -F(6)); + EXPECT_EQ(-f, f_neg); + f.NegateInPlace(); + EXPECT_EQ(f, f_neg); + + f = F4(F(3), F(4), F(5), F(6)); + F4 f_dbl(F(6), F(8), F(10), F(12)); + EXPECT_EQ(f.Double(), f_dbl); + f.DoubleInPlace(); + EXPECT_EQ(f, f_dbl); +} + +TEST_F(QuaticExtensionFieldTest, MultiplicativeOperators) { + struct { + F4 a; + F4 b; + F4 mul; + F4 adb; + F4 bda; + } tests[] = { + { + {F(1), F(2), F(3), F(4)}, + {F(3), F(5), F(6), F(8)}, + {F(597), F(539), F(377), F(47)}, + {F(1144494179), F(1502926259), F(1509084158), F(151175067)}, + {F(653096429), F(494869942), F(67683040), F(1807436149)}, + }, + }; + + for (const auto& test : tests) { + EXPECT_EQ(test.a * test.b, test.mul); + EXPECT_EQ(test.b * test.a, test.mul); + EXPECT_EQ(test.a / test.b, test.adb); + EXPECT_EQ(test.b / test.a, test.bda); + + F4 tmp = test.a; + tmp *= test.b; + EXPECT_EQ(tmp, test.mul); + ASSERT_TRUE(tmp /= test.b); + EXPECT_EQ(tmp, test.a); + } +} + +TEST_F(QuaticExtensionFieldTest, MultiplicativeOperators2) { + F4 f(F(3), F(4), F(5), F(6)); + F4 f_mul(F(6), F(8), F(10), F(12)); + EXPECT_EQ(f * F(2), f_mul); + f *= F(2); + EXPECT_EQ(f, f_mul); +} + +TEST_F(QuaticExtensionFieldTest, MultiplicativeGroupOperators) { + F4 f = F4::Random(); + std::optional f_inv = f.Inverse(); + if (UNLIKELY(f.IsZero())) { + ASSERT_FALSE(f_inv); + ASSERT_FALSE(f.InverseInPlace()); + } else { + EXPECT_EQ(f * *f_inv, F4::One()); + F4 f_tmp = f; + EXPECT_EQ(**f.InverseInPlace() * f_tmp, F4::One()); + } + + f = F4(F(3), F(4), F(5), F(6)); + F4 f_sqr = F4(F(812), F(684), F(442), F(76)); + EXPECT_EQ(f.Square(), f_sqr); + f.SquareInPlace(); + EXPECT_EQ(f, f_sqr); +} + +TEST_F(QuaticExtensionFieldTest, JsonValueConverter) { + F4 expected_point(F(1), F(2), F(3), F(4)); + std::string expected_json = R"({"c0":1,"c1":2,"c2":3,"c3":4})"; + + F4 p; + std::string error; + ASSERT_TRUE(base::ParseJson(expected_json, &p, &error)); + ASSERT_TRUE(error.empty()); + EXPECT_EQ(p, expected_point); + + std::string json = base::WriteToJson(p); + EXPECT_EQ(json, expected_json); +} + +} // namespace tachyon::math