diff --git a/icicle/include/curves/macro.h b/icicle/include/curves/macro.h index 6ce3cb664..b4833d447 100644 --- a/icicle/include/curves/macro.h +++ b/icicle/include/curves/macro.h @@ -22,7 +22,7 @@ typedef Affine affine_t; #define G2_CURVE_DEFINITIONS \ - typedef ExtensionField g2_point_field_t; \ + typedef ComplexExtensionField g2_point_field_t; \ static constexpr g2_point_field_t g2_generator_x = \ g2_point_field_t{point_field_t{g2_gen_x_re}, point_field_t{g2_gen_x_im}}; \ static constexpr g2_point_field_t g2_generator_y = \ diff --git a/icicle/include/fields/quartic_extension.cuh b/icicle/include/fields/quartic_extension.cuh index dab83487b..59a95694a 100644 --- a/icicle/include/fields/quartic_extension.cuh +++ b/icicle/include/fields/quartic_extension.cuh @@ -5,7 +5,7 @@ #include "gpu-utils/sharedmem.cuh" template -class QuartExtensionField +class QuarticExtensionField { private: typedef typename T::Wide FWide; @@ -36,88 +36,88 @@ public: FF im2; FF im3; - static constexpr HOST_DEVICE_INLINE QuartExtensionField zero() + static constexpr HOST_DEVICE_INLINE QuarticExtensionField zero() { - return QuartExtensionField{FF::zero(), FF::zero(), FF::zero(), FF::zero()}; + return QuarticExtensionField{FF::zero(), FF::zero(), FF::zero(), FF::zero()}; } - static constexpr HOST_DEVICE_INLINE QuartExtensionField one() + static constexpr HOST_DEVICE_INLINE QuarticExtensionField one() { - return QuartExtensionField{FF::one(), FF::zero(), FF::zero(), FF::zero()}; + return QuarticExtensionField{FF::one(), FF::zero(), FF::zero(), FF::zero()}; } - static constexpr HOST_DEVICE_INLINE QuartExtensionField to_montgomery(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField to_montgomery(const QuarticExtensionField& xs) { - return QuartExtensionField{ + return QuarticExtensionField{ FF::to_montgomery(xs.real), FF::to_montgomery(xs.im1), FF::to_montgomery(xs.im2), FF::to_montgomery(xs.im3)}; } - static constexpr HOST_DEVICE_INLINE QuartExtensionField from_montgomery(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField from_montgomery(const QuarticExtensionField& xs) { - return QuartExtensionField{ + return QuarticExtensionField{ FF::from_montgomery(xs.real), FF::from_montgomery(xs.im1), FF::from_montgomery(xs.im2), FF::from_montgomery(xs.im3)}; } - static HOST_INLINE QuartExtensionField rand_host() + static HOST_INLINE QuarticExtensionField rand_host() { - return QuartExtensionField{FF::rand_host(), FF::rand_host(), FF::rand_host(), FF::rand_host()}; + return QuarticExtensionField{FF::rand_host(), FF::rand_host(), FF::rand_host(), FF::rand_host()}; } - static void rand_host_many(QuartExtensionField* out, int size) + static void rand_host_many(QuarticExtensionField* out, int size) { for (int i = 0; i < size; i++) out[i] = rand_host(); } template - static constexpr HOST_DEVICE_INLINE QuartExtensionField sub_modulus(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField sub_modulus(const QuarticExtensionField& xs) { - return QuartExtensionField{ + return QuarticExtensionField{ FF::sub_modulus(&xs.real), FF::sub_modulus(&xs.im1), FF::sub_modulus(&xs.im2), FF::sub_modulus(&xs.im3)}; } - friend std::ostream& operator<<(std::ostream& os, const QuartExtensionField& xs) + friend std::ostream& operator<<(std::ostream& os, const QuarticExtensionField& xs) { os << "{ Real: " << xs.real << " }; { Im1: " << xs.im1 << " }; { Im2: " << xs.im2 << " }; { Im3: " << xs.im3 << " };"; return os; } - friend HOST_DEVICE_INLINE QuartExtensionField operator+(QuartExtensionField xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator+(QuarticExtensionField xs, const QuarticExtensionField& ys) { - return QuartExtensionField{xs.real + ys.real, xs.im1 + ys.im1, xs.im2 + ys.im2, xs.im3 + ys.im3}; + return QuarticExtensionField{xs.real + ys.real, xs.im1 + ys.im1, xs.im2 + ys.im2, xs.im3 + ys.im3}; } - friend HOST_DEVICE_INLINE QuartExtensionField operator-(QuartExtensionField xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator-(QuarticExtensionField xs, const QuarticExtensionField& ys) { - return QuartExtensionField{xs.real - ys.real, xs.im1 - ys.im1, xs.im2 - ys.im2, xs.im3 - ys.im3}; + return QuarticExtensionField{xs.real - ys.real, xs.im1 - ys.im1, xs.im2 - ys.im2, xs.im3 - ys.im3}; } - friend HOST_DEVICE_INLINE QuartExtensionField operator+(FF xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator+(FF xs, const QuarticExtensionField& ys) { - return QuartExtensionField{xs + ys.real, ys.im1, ys.im2, ys.im3}; + return QuarticExtensionField{xs + ys.real, ys.im1, ys.im2, ys.im3}; } - friend HOST_DEVICE_INLINE QuartExtensionField operator-(FF xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator-(FF xs, const QuarticExtensionField& ys) { - return QuartExtensionField{xs - ys.real, FF::neg(ys.im1), FF::neg(ys.im2), FF::neg(ys.im3)}; + return QuarticExtensionField{xs - ys.real, FF::neg(ys.im1), FF::neg(ys.im2), FF::neg(ys.im3)}; } - friend HOST_DEVICE_INLINE QuartExtensionField operator+(QuartExtensionField xs, const FF& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator+(QuarticExtensionField xs, const FF& ys) { - return QuartExtensionField{xs.real + ys, xs.im1, xs.im2, xs.im3}; + return QuarticExtensionField{xs.real + ys, xs.im1, xs.im2, xs.im3}; } - friend HOST_DEVICE_INLINE QuartExtensionField operator-(QuartExtensionField xs, const FF& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator-(QuarticExtensionField xs, const FF& ys) { - return QuartExtensionField{xs.real - ys, xs.im1, xs.im2, xs.im3}; + return QuarticExtensionField{xs.real - ys, xs.im1, xs.im2, xs.im3}; } template static constexpr HOST_DEVICE_INLINE ExtensionWide - mul_wide(const QuartExtensionField& xs, const QuartExtensionField& ys) + mul_wide(const QuarticExtensionField& xs, const QuarticExtensionField& ys) { if (CONFIG::nonresidue_is_negative) return ExtensionWide{ @@ -144,46 +144,46 @@ public: } template - static constexpr HOST_DEVICE_INLINE ExtensionWide mul_wide(const QuartExtensionField& xs, const FF& ys) + static constexpr HOST_DEVICE_INLINE ExtensionWide mul_wide(const QuarticExtensionField& xs, const FF& ys) { return ExtensionWide{ FF::mul_wide(xs.real, ys), FF::mul_wide(xs.im1, ys), FF::mul_wide(xs.im2, ys), FF::mul_wide(xs.im3, ys)}; } template - static constexpr HOST_DEVICE_INLINE ExtensionWide mul_wide(const FF& xs, const QuartExtensionField& ys) + static constexpr HOST_DEVICE_INLINE ExtensionWide mul_wide(const FF& xs, const QuarticExtensionField& ys) { return ExtensionWide{ FF::mul_wide(xs, ys.real), FF::mul_wide(xs, ys.im1), FF::mul_wide(xs, ys.im2), FF::mul_wide(xs, ys.im3)}; } template - static constexpr HOST_DEVICE_INLINE QuartExtensionField reduce(const ExtensionWide& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField reduce(const ExtensionWide& xs) { - return QuartExtensionField{ + return QuarticExtensionField{ FF::template reduce(xs.real), FF::template reduce(xs.im1), FF::template reduce(xs.im2), FF::template reduce(xs.im3)}; } template - friend HOST_DEVICE_INLINE QuartExtensionField operator*(const T1& xs, const T2& ys) + friend HOST_DEVICE_INLINE QuarticExtensionField operator*(const T1& xs, const T2& ys) { ExtensionWide xy = mul_wide(xs, ys); return reduce(xy); } - friend HOST_DEVICE_INLINE bool operator==(const QuartExtensionField& xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE bool operator==(const QuarticExtensionField& xs, const QuarticExtensionField& ys) { return (xs.real == ys.real) && (xs.im1 == ys.im1) && (xs.im2 == ys.im2) && (xs.im3 == ys.im3); } - friend HOST_DEVICE_INLINE bool operator!=(const QuartExtensionField& xs, const QuartExtensionField& ys) + friend HOST_DEVICE_INLINE bool operator!=(const QuarticExtensionField& xs, const QuarticExtensionField& ys) { return !(xs == ys); } template - static constexpr HOST_DEVICE_INLINE QuartExtensionField mul_unsigned(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField mul_unsigned(const QuarticExtensionField& xs) { return { FF::template mul_unsigned(xs.real), FF::template mul_unsigned(xs.im1), @@ -191,27 +191,27 @@ public: } template - static constexpr HOST_DEVICE_INLINE ExtensionWide sqr_wide(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE ExtensionWide sqr_wide(const QuarticExtensionField& xs) { // TODO: change to a more efficient squaring return mul_wide(xs, xs); } template - static constexpr HOST_DEVICE_INLINE QuartExtensionField sqr(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField sqr(const QuarticExtensionField& xs) { // TODO: change to a more efficient squaring return xs * xs; } template - static constexpr HOST_DEVICE_INLINE QuartExtensionField neg(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField neg(const QuarticExtensionField& xs) { return {FF::neg(xs.real), FF::neg(xs.im1), FF::neg(xs.im2), FF::neg(xs.im3)}; } // inverse of zero is set to be zero which is what we want most of the time - static constexpr HOST_DEVICE_INLINE QuartExtensionField inverse(const QuartExtensionField& xs) + static constexpr HOST_DEVICE_INLINE QuarticExtensionField inverse(const QuarticExtensionField& xs) { FF x, x0, x2; if (CONFIG::nonresidue_is_negative) { @@ -251,10 +251,10 @@ public: }; template -struct SharedMemory> { - __device__ QuartExtensionField* getPointer() +struct SharedMemory> { + __device__ QuarticExtensionField* getPointer() { - extern __shared__ QuartExtensionField s_ext4_scalar_[]; + extern __shared__ QuarticExtensionField s_ext4_scalar_[]; return s_ext4_scalar_; } }; \ No newline at end of file diff --git a/icicle/include/fields/stark_fields/babybear.cuh b/icicle/include/fields/stark_fields/babybear.cuh index 6932d54ed..99f170369 100644 --- a/icicle/include/fields/stark_fields/babybear.cuh +++ b/icicle/include/fields/stark_fields/babybear.cuh @@ -27,5 +27,5 @@ namespace babybear { /** * Extension field of `scalar_t` enabled if `-DEXT_FIELD` env variable is. */ - typedef QuartExtensionField extension_t; + typedef QuarticExtensionField extension_t; } // namespace babybear diff --git a/icicle/include/fields/stark_fields/m31.cuh b/icicle/include/fields/stark_fields/m31.cuh index 724b7bcc2..9ce343dd1 100644 --- a/icicle/include/fields/stark_fields/m31.cuh +++ b/icicle/include/fields/stark_fields/m31.cuh @@ -276,7 +276,7 @@ namespace m31 { /** * Extension field of `scalar_t` enabled if `-DEXT_FIELD` env variable is. */ - typedef QuartExtensionField q_extension_t; + typedef QuarticExtensionField q_extension_t; } // namespace m31 template diff --git a/icicle/src/ntt/kernel_ntt.cu b/icicle/src/ntt/kernel_ntt.cu index 01b6b5b20..d08e78b2d 100644 --- a/icicle/src/ntt/kernel_ntt.cu +++ b/icicle/src/ntt/kernel_ntt.cu @@ -1187,10 +1187,13 @@ namespace mxntt { return CHK_LAST(); } +#ifdef DCCT + uint32_t twiddles_offset = 0; +#endif + // general case: uint32_t nof_blocks = (1UL << (log_size - 9)) * (columns_batch ? ((batch_size + 31) / 32) * 32 : batch_size); if (dit) { - uint32_t twiddles_offset = 0; for (int i = 0; i < 5; i++) { uint32_t stage_size = fast_tw ? STAGE_SIZES_HOST_FT[log_size][i] : STAGE_SIZES_HOST[log_size][i]; uint32_t stride_log = 0; @@ -1232,7 +1235,6 @@ namespace mxntt { } } else { // dif bool first_run = false, prev_stage = false; - uint32_t twiddles_offset = 0; for (int i = 4; i >= 0; i--) { uint32_t stage_size = fast_tw ? STAGE_SIZES_HOST_FT[log_size][i] : STAGE_SIZES_HOST[log_size][i]; uint32_t stride_log = 0; @@ -1240,8 +1242,8 @@ namespace mxntt { stride_log += fast_tw ? STAGE_SIZES_HOST_FT[log_size][j] : STAGE_SIZES_HOST[log_size][j]; first_run = stage_size && !prev_stage; - uint32_t nof_ntt_blocks = (1 << log_size - stage_size) * (columns_batch ? 1 : batch_size); #ifdef DCCT + uint32_t nof_ntt_blocks = (1 << log_size - stage_size) * (columns_batch ? 1 : batch_size); if (stage_size == 6) ntt64_dcct<<>>( first_run ? in : out, out, basic_twiddles, log_size, tw_log_size, columns_batch ? batch_size : 0,