From 07d7a2445750dbea97b47a4c0384b544ee163181 Mon Sep 17 00:00:00 2001 From: Yuval Shekel Date: Thu, 13 Jun 2024 14:37:01 +0300 Subject: [PATCH] montgomery conversion for cuda --- .../cpu/src/curve/cpu_mont_conversion.cpp | 4 +- icicle_v3/backend/cuda/CMakeLists.txt | 5 +- icicle_v3/backend/cuda/include/cuda_mont.cuh | 54 +++++++++++ icicle_v3/backend/cuda/src/curve/cuda_mont.cu | 33 +++++++ icicle_v3/backend/cuda/src/field/cuda_mont.cu | 51 +--------- .../icicle/curves/montgomery_conversion.h | 40 +------- icicle_v3/include/icicle/curves/projective.h | 2 +- icicle_v3/include/icicle/vec_ops.h | 3 +- icicle_v3/src/curves/ffi_extern.cpp | 4 +- .../src/curves/montgomery_conversion.cpp | 32 +++---- icicle_v3/src/vec_ops.cpp | 6 +- icicle_v3/tests/test_curve_api.cpp | 96 ++++++++++++------- icicle_v3/tests/test_field_api.cpp | 2 +- 13 files changed, 184 insertions(+), 148 deletions(-) create mode 100644 icicle_v3/backend/cuda/include/cuda_mont.cuh create mode 100644 icicle_v3/backend/cuda/src/curve/cuda_mont.cu diff --git a/icicle_v3/backend/cpu/src/curve/cpu_mont_conversion.cpp b/icicle_v3/backend/cpu/src/curve/cpu_mont_conversion.cpp index af6f0ea52..5c94ebabb 100644 --- a/icicle_v3/backend/cpu/src/curve/cpu_mont_conversion.cpp +++ b/icicle_v3/backend/cpu/src/curve/cpu_mont_conversion.cpp @@ -10,8 +10,8 @@ using namespace curve_config; using namespace icicle; template -eIcicleError cpu_convert_mont( - const Device& device, const T* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, T* output) +eIcicleError +cpu_convert_mont(const Device& device, const T* input, size_t n, bool is_into, const VecOpsConfig& config, T* output) { for (size_t i = 0; i < n; ++i) { output[i] = is_into ? T::to_montgomery(input[i]) : T::from_montgomery(input[i]); diff --git a/icicle_v3/backend/cuda/CMakeLists.txt b/icicle_v3/backend/cuda/CMakeLists.txt index d53f1ec93..05ba5ff00 100644 --- a/icicle_v3/backend/cuda/CMakeLists.txt +++ b/icicle_v3/backend/cuda/CMakeLists.txt @@ -37,7 +37,10 @@ target_link_libraries(icicle_cuda_field PRIVATE ${CUDA_LIBRARIES}) # Link to CUD # curve API library if (CURVE) - add_library(icicle_cuda_curve SHARED src/curve/cuda_msm.cu) + add_library(icicle_cuda_curve SHARED + src/curve/cuda_msm.cu + src/curve/cuda_mont.cu + ) target_include_directories(icicle_cuda_curve PRIVATE include) target_link_libraries(icicle_cuda_curve PUBLIC icicle_device icicle_curve) set_target_properties(icicle_cuda_curve PROPERTIES OUTPUT_NAME "icicle_cuda_curve_${FIELD}") diff --git a/icicle_v3/backend/cuda/include/cuda_mont.cuh b/icicle_v3/backend/cuda/include/cuda_mont.cuh new file mode 100644 index 000000000..729637bb6 --- /dev/null +++ b/icicle_v3/backend/cuda/include/cuda_mont.cuh @@ -0,0 +1,54 @@ +#include +#include + +#include "icicle/errors.h" +#include "icicle/vec_ops.h" +#include "gpu-utils/error_handler.h" + +namespace montgomery { +#define MAX_THREADS_PER_BLOCK 256 + + template + __global__ void MontgomeryKernel(const E* input, int n, E* output) + { + int tid = blockIdx.x * blockDim.x + threadIdx.x; + if (tid < n) { output[tid] = is_into ? E::to_montgomery(input[tid]) : E::from_montgomery(input[tid]); } + } + + template + cudaError_t ConvertMontgomery(const E* input, size_t n, const VecOpsConfig& config, E* output) + { + cudaStream_t cuda_stream = reinterpret_cast(config.stream); + + E *d_alloc_out = nullptr, *d_alloc_in = nullptr, *d_out; + const E* d_in; + if (!config.is_a_on_device) { + CHK_IF_RETURN(cudaMallocAsync(&d_alloc_in, n * sizeof(E), cuda_stream)); + CHK_IF_RETURN(cudaMemcpyAsync(d_alloc_in, input, n * sizeof(E), cudaMemcpyHostToDevice, cuda_stream)); + d_in = d_alloc_in; + } else { + d_in = input; + } + + if (!config.is_result_on_device) { + CHK_IF_RETURN(cudaMallocAsync(&d_alloc_out, n * sizeof(E), cuda_stream)); + d_out = d_alloc_out; + } else { + d_out = output; + } + + int num_threads = MAX_THREADS_PER_BLOCK; + int num_blocks = (n + num_threads - 1) / num_threads; + MontgomeryKernel<<>>(d_in, n, d_out); + + if (d_alloc_in) { CHK_IF_RETURN(cudaFreeAsync(d_alloc_in, cuda_stream)); } + if (d_alloc_out) { + CHK_IF_RETURN(cudaMemcpyAsync(output, d_out, n * sizeof(E), cudaMemcpyDeviceToHost, cuda_stream)); + CHK_IF_RETURN(cudaFreeAsync(d_out, cuda_stream)); + } + if (!config.is_async) return CHK_STICKY(cudaStreamSynchronize(cuda_stream)); + + return CHK_LAST(); + } + +} // namespace montgomery \ No newline at end of file diff --git a/icicle_v3/backend/cuda/src/curve/cuda_mont.cu b/icicle_v3/backend/cuda/src/curve/cuda_mont.cu new file mode 100644 index 000000000..d0a41e984 --- /dev/null +++ b/icicle_v3/backend/cuda/src/curve/cuda_mont.cu @@ -0,0 +1,33 @@ +#include +#include + +#include "icicle/errors.h" +#include "icicle/curves/montgomery_conversion.h" +#include "gpu-utils/error_handler.h" +#include "error_translation.h" +#include "cuda_mont.cuh" + +#include "icicle/curves/curve_config.h" +using namespace curve_config; +using namespace icicle; + +namespace icicle { + + template + eIcicleError + cuda_convert_mont(const Device& device, const T* input, size_t n, bool is_into, const VecOpsConfig& config, T* output) + { + cudaError_t err = is_into ? montgomery::ConvertMontgomery(input, n, config, output) + : montgomery::ConvertMontgomery(input, n, config, output); + return translateCudaError(err); + } + + REGISTER_AFFINE_CONVERT_MONTGOMERY_BACKEND("CUDA", cuda_convert_mont); + REGISTER_PROJECTIVE_CONVERT_MONTGOMERY_BACKEND("CUDA", cuda_convert_mont); + +#ifdef G2 + REGISTER_AFFINE_G2_CONVERT_MONTGOMERY_BACKEND("CUDA", cuda_convert_mont); + REGISTER_PROJECTIVE_G2_CONVERT_MONTGOMERY_BACKEND("CUDA", cuda_convert_mont); +#endif // G2 + +} // namespace icicle \ No newline at end of file diff --git a/icicle_v3/backend/cuda/src/field/cuda_mont.cu b/icicle_v3/backend/cuda/src/field/cuda_mont.cu index d7cf77919..cfd20f0ee 100644 --- a/icicle_v3/backend/cuda/src/field/cuda_mont.cu +++ b/icicle_v3/backend/cuda/src/field/cuda_mont.cu @@ -5,56 +5,10 @@ #include "icicle/vec_ops.h" #include "gpu-utils/error_handler.h" #include "error_translation.h" +#include "cuda_mont.cuh" namespace icicle { -#define MAX_THREADS_PER_BLOCK 256 - - template - __global__ void MontgomeryKernel(const E* input, int n, bool is_into, E* output) - { - int tid = blockIdx.x * blockDim.x + threadIdx.x; - if (tid < n) { output[tid] = is_into ? E::to_montgomery(input[tid]) : E::from_montgomery(input[tid]); } - } - - template - cudaError_t ConvertMontgomery(const E* input, int n, bool is_into, const VecOpsConfig& config, E* output) - { - cudaStream_t cuda_stream = reinterpret_cast(config.stream); - - E *d_alloc_out = nullptr, *d_alloc_in = nullptr, *d_out; - const E* d_in; - if (!config.is_a_on_device) { - CHK_IF_RETURN(cudaMallocAsync(&d_alloc_in, n * sizeof(E), cuda_stream)); - CHK_IF_RETURN(cudaMemcpyAsync(d_alloc_in, input, n * sizeof(E), cudaMemcpyHostToDevice, cuda_stream)); - d_in = d_alloc_in; - } else { - d_in = input; - } - - if (!config.is_result_on_device) { - CHK_IF_RETURN(cudaMallocAsync(&d_alloc_out, n * sizeof(E), cuda_stream)); - d_out = d_alloc_out; - } else { - d_out = output; - } - - int num_threads = MAX_THREADS_PER_BLOCK; - int num_blocks = (n + num_threads - 1) / num_threads; - MontgomeryKernel<<>>(d_in, n, is_into, d_out); - - if (d_alloc_in) { CHK_IF_RETURN(cudaFreeAsync(d_alloc_in, cuda_stream)); } - if (d_alloc_out) { - CHK_IF_RETURN(cudaMemcpyAsync(output, d_out, n * sizeof(E), cudaMemcpyDeviceToHost, cuda_stream)); - CHK_IF_RETURN(cudaFreeAsync(d_out, cuda_stream)); - } - if (!config.is_async) return CHK_STICKY(cudaStreamSynchronize(cuda_stream)); - - return CHK_LAST(); - } - - /************************************ REGISTRATION ************************************/ - #include "icicle/fields/field_config.h" using namespace field_config; @@ -62,7 +16,8 @@ namespace icicle { eIcicleError convert_montgomery_cuda( const Device& device, const F* input, uint64_t n, bool is_into, const VecOpsConfig& config, F* output) { - auto err = ConvertMontgomery(input, n, is_into, config, output); + auto err = is_into ? montgomery::ConvertMontgomery(input, n, config, output) + : montgomery::ConvertMontgomery(input, n, config, output); return translateCudaError(err); } diff --git a/icicle_v3/include/icicle/curves/montgomery_conversion.h b/icicle_v3/include/icicle/curves/montgomery_conversion.h index a4caa00ea..9772b8e25 100644 --- a/icicle_v3/include/icicle/curves/montgomery_conversion.h +++ b/icicle_v3/include/icicle/curves/montgomery_conversion.h @@ -9,6 +9,7 @@ #include "icicle/curves/affine.h" #include "icicle/curves/projective.h" +#include "icicle/vec_ops.h" #include "icicle/fields/field.h" #include "icicle/curves/curve_config.h" @@ -16,41 +17,10 @@ using namespace curve_config; namespace icicle { - /*************************** Frontend APIs ***************************/ - - struct ConvertMontgomeryConfig { - icicleStreamHandle stream; /**< stream for async execution. */ - bool is_input_on_device; - bool is_output_on_device; - bool is_async; - - ConfigExtension ext; /** backend specific extensions*/ - }; - - static ConvertMontgomeryConfig default_convert_montgomery_config() - { - ConvertMontgomeryConfig config = { - nullptr, // stream - false, // is_input_on_device - false, // is_output_on_device - false, // is_async - }; - return config; - } - - template - eIcicleError - points_convert_montgomery(const T* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, T* output); - /*************************** Backend registration ***************************/ using AffineConvertMontImpl = std::function; + const Device& device, const affine_t* input, size_t n, bool is_into, const VecOpsConfig& config, affine_t* output)>; void register_affine_convert_montgomery(const std::string& deviceType, AffineConvertMontImpl); @@ -67,7 +37,7 @@ namespace icicle { const projective_t* input, size_t n, bool is_into, - const ConvertMontgomeryConfig& config, + const VecOpsConfig& config, projective_t* output)>; void register_projective_convert_montgomery(const std::string& deviceType, ProjectiveConvertMontImpl); @@ -86,7 +56,7 @@ namespace icicle { const g2_affine_t* input, size_t n, bool is_into, - const ConvertMontgomeryConfig& config, + const VecOpsConfig& config, g2_affine_t* output)>; void register_affine_g2_convert_montgomery(const std::string& deviceType, AffineG2ConvertMontImpl); @@ -104,7 +74,7 @@ namespace icicle { const g2_projective_t* input, size_t n, bool is_into, - const ConvertMontgomeryConfig& config, + const VecOpsConfig& config, g2_projective_t* output)>; void register_projective_g2_convert_montgomery(const std::string& deviceType, ProjectiveG2ConvertMontImpl); diff --git a/icicle_v3/include/icicle/curves/projective.h b/icicle_v3/include/icicle/curves/projective.h index fc3a874bf..0055c008d 100644 --- a/icicle_v3/include/icicle/curves/projective.h +++ b/icicle_v3/include/icicle/curves/projective.h @@ -231,7 +231,7 @@ class Projective out[i] = (i % size < 100) ? rand_host() : out[i - 100]; } - static void rand_host_many_affine(Affine* out, int size) + static void rand_host_many(Affine* out, int size) { for (int i = 0; i < size; i++) out[i] = (i % size < 100) ? to_affine(rand_host()) : out[i - 100]; diff --git a/icicle_v3/include/icicle/vec_ops.h b/icicle_v3/include/icicle/vec_ops.h index 3ec3035a5..14fbd728a 100644 --- a/icicle_v3/include/icicle/vec_ops.h +++ b/icicle_v3/include/icicle/vec_ops.h @@ -56,8 +56,7 @@ namespace icicle { eIcicleError vector_mul(const S* vec_a, const S* vec_b, uint64_t n, const VecOpsConfig& config, S* output); template - eIcicleError - scalar_convert_montgomery(const S* input, uint64_t size, bool is_into, const VecOpsConfig& config, S* output); + eIcicleError convert_montgomery(const S* input, uint64_t size, bool is_into, const VecOpsConfig& config, S* output); /*************************** Backend registration ***************************/ diff --git a/icicle_v3/src/curves/ffi_extern.cpp b/icicle_v3/src/curves/ffi_extern.cpp index b774775ee..4e3ff8499 100644 --- a/icicle_v3/src/curves/ffi_extern.cpp +++ b/icicle_v3/src/curves/ffi_extern.cpp @@ -27,7 +27,7 @@ extern "C" void CONCAT_EXPAND(CURVE, generate_projective_points)(projective_t* p extern "C" void CONCAT_EXPAND(CURVE, generate_affine_points)(affine_t* points, int size) { - projective_t::rand_host_many_affine(points, size); + projective_t::rand_host_many(points, size); } /********************************** G2 **********************************/ @@ -53,6 +53,6 @@ extern "C" void CONCAT_EXPAND(CURVE, g2_generate_projective_points)(g2_projectiv extern "C" void CONCAT_EXPAND(CURVE, g2_generate_affine_points)(g2_affine_t* points, int size) { - g2_projective_t::rand_host_many_affine(points, size); + g2_projective_t::rand_host_many(points, size); } #endif // G2 diff --git a/icicle_v3/src/curves/montgomery_conversion.cpp b/icicle_v3/src/curves/montgomery_conversion.cpp index 2887931b0..0162c9b1f 100644 --- a/icicle_v3/src/curves/montgomery_conversion.cpp +++ b/icicle_v3/src/curves/montgomery_conversion.cpp @@ -10,14 +10,14 @@ namespace icicle { ICICLE_DISPATCHER_INST(AffineConvertMont, affine_convert_montgomery, AffineConvertMontImpl); extern "C" eIcicleError CONCAT_EXPAND(CURVE, affine_convert_montgomery)( - const affine_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, affine_t* output) + const affine_t* input, size_t n, bool is_into, const VecOpsConfig& config, affine_t* output) { return AffineConvertMont::execute(input, n, is_into, config, output); } template <> - eIcicleError points_convert_montgomery( - const affine_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, affine_t* output) + eIcicleError + convert_montgomery(const affine_t* input, size_t n, bool is_into, const VecOpsConfig& config, affine_t* output) { return CONCAT_EXPAND(CURVE, affine_convert_montgomery)(input, n, is_into, config, output); } @@ -26,15 +26,15 @@ namespace icicle { ICICLE_DISPATCHER_INST(AffineG2ConvertMont, affine_g2_convert_montgomery, AffineG2ConvertMontImpl); extern "C" eIcicleError CONCAT_EXPAND(CURVE, affine_g2_convert_montgomery)( - const g2_affine_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, g2_affine_t* output) + const g2_affine_t* input, size_t n, bool is_into, const VecOpsConfig& config, g2_affine_t* output) { return AffineG2ConvertMont::execute(input, n, is_into, config, output); } #ifndef G1_AFFINE_SAME_TYPE_AS_G2_AFFINE template <> - eIcicleError points_convert_montgomery( - const g2_affine_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, g2_affine_t* output) + eIcicleError + convert_montgomery(const g2_affine_t* input, size_t n, bool is_into, const VecOpsConfig& config, g2_affine_t* output) { return CONCAT_EXPAND(CURVE, affine_g2_convert_montgomery)(input, n, is_into, config, output); } @@ -44,14 +44,14 @@ namespace icicle { ICICLE_DISPATCHER_INST(ProjectiveConvertMont, projective_convert_montgomery, ProjectiveConvertMontImpl); extern "C" eIcicleError CONCAT_EXPAND(CURVE, projective_convert_montgomery)( - const projective_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, projective_t* output) + const projective_t* input, size_t n, bool is_into, const VecOpsConfig& config, projective_t* output) { return ProjectiveConvertMont::execute(input, n, is_into, config, output); } template <> - eIcicleError points_convert_montgomery( - const projective_t* input, size_t n, bool is_into, const ConvertMontgomeryConfig& config, projective_t* output) + eIcicleError convert_montgomery( + const projective_t* input, size_t n, bool is_into, const VecOpsConfig& config, projective_t* output) { return CONCAT_EXPAND(CURVE, projective_convert_montgomery)(input, n, is_into, config, output); } @@ -60,22 +60,14 @@ namespace icicle { ICICLE_DISPATCHER_INST(ProjectiveG2ConvertMont, projective_g2_convert_montgomery, ProjectiveG2ConvertMontImpl); extern "C" eIcicleError CONCAT_EXPAND(CURVE, projective_g2_convert_montgomery)( - const g2_projective_t* input, - size_t n, - bool is_into, - const ConvertMontgomeryConfig& config, - g2_projective_t* output) + const g2_projective_t* input, size_t n, bool is_into, const VecOpsConfig& config, g2_projective_t* output) { return ProjectiveG2ConvertMont::execute(input, n, is_into, config, output); } template <> - eIcicleError points_convert_montgomery( - const g2_projective_t* input, - size_t n, - bool is_into, - const ConvertMontgomeryConfig& config, - g2_projective_t* output) + eIcicleError convert_montgomery( + const g2_projective_t* input, size_t n, bool is_into, const VecOpsConfig& config, g2_projective_t* output) { return CONCAT_EXPAND(CURVE, projective_g2_convert_montgomery)(input, n, is_into, config, output); } diff --git a/icicle_v3/src/vec_ops.cpp b/icicle_v3/src/vec_ops.cpp index 1d9f3adaa..01e4b2437 100644 --- a/icicle_v3/src/vec_ops.cpp +++ b/icicle_v3/src/vec_ops.cpp @@ -113,8 +113,8 @@ namespace icicle { } template <> - eIcicleError scalar_convert_montgomery( - const scalar_t* input, uint64_t size, bool is_into, const VecOpsConfig& config, scalar_t* output) + eIcicleError + convert_montgomery(const scalar_t* input, uint64_t size, bool is_into, const VecOpsConfig& config, scalar_t* output) { return CONCAT_EXPAND(FIELD, scalar_convert_montgomery)(input, size, is_into, config, output); } @@ -130,7 +130,7 @@ namespace icicle { } template <> - eIcicleError scalar_convert_montgomery( + eIcicleError convert_montgomery( const extension_t* input, uint64_t size, bool is_into, const VecOpsConfig& config, extension_t* output) { return CONCAT_EXPAND(FIELD, convert_montgomery_ext_field)(input, size, is_into, config, output); diff --git a/icicle_v3/tests/test_curve_api.cpp b/icicle_v3/tests/test_curve_api.cpp index ae197e02a..eb5164849 100644 --- a/icicle_v3/tests/test_curve_api.cpp +++ b/icicle_v3/tests/test_curve_api.cpp @@ -21,6 +21,8 @@ using FpMicroseconds = std::chrono::duration + void random_scalars(T* arr, uint64_t count) + { + for (uint64_t i = 0; i < count; i++) + arr[i] = i < 1000 ? T::rand_host() : arr[i - 1000]; + } + template void MSM_test() { @@ -50,7 +67,7 @@ class CurveApiTest : public ::testing::Test auto bases = std::make_unique(N); scalar_t::rand_host_many(scalars.get(), N); - P::rand_host_many_affine(bases.get(), N); + P::rand_host_many(bases.get(), N); P result{}; @@ -73,43 +90,56 @@ class CurveApiTest : public ::testing::Test // TODO test something } - template + template void mont_conversion_test() { - // Note: this test doesn't really test correct mont conversion (since there is no arithmetic in mont) but checks - // that - // it does some conversion and back to original - A affine_point, affine_point_converted; - P projective_point, projective_point_converted; - - projective_point = P::rand_host(); - affine_point = projective_point.to_affine(); - - icicle_set_device({"CPU", 0}); - - // (1) converting to mont and check not equal to original - auto config = default_convert_montgomery_config(); - points_convert_montgomery(&affine_point, 1, true /*into mont*/, config, &affine_point_converted); - points_convert_montgomery(&projective_point, 1, true /*into mont*/, config, &projective_point_converted); - - ASSERT_NE(affine_point, affine_point_converted); // check that it was converted to mont - ASSERT_NE(projective_point.x, projective_point_converted.x); // check that it was converted to mont - - // (2) converting back from mont and check equal - points_convert_montgomery(&projective_point_converted, 1, false /*from mont*/, config, &projective_point_converted); - points_convert_montgomery(&affine_point_converted, 1, false /*from mont*/, config, &affine_point_converted); - - ASSERT_EQ(affine_point, affine_point_converted); - ASSERT_EQ(projective_point, projective_point_converted); + const int N = 1 << 6; + auto elements = std::make_unique(N); + auto main_target_output = std::make_unique(N); + auto ref_target_output = std::make_unique(N); + P::rand_host_many(elements.get(), N); + + auto run = + [&](const std::string& dev_type, T* input, T* output, bool into_mont, bool measure, const char* msg, int iters) { + Device dev = {dev_type.c_str(), 0}; + icicle_set_device(dev); + auto config = default_vec_ops_config(); + + std::ostringstream oss; + oss << dev_type << " " << msg; + + START_TIMER(MONTGOMERY) + for (int i = 0; i < iters; ++i) { + ICICLE_CHECK(convert_montgomery(input, N, into_mont, config, output)); + } + END_TIMER(MONTGOMERY, oss.str().c_str(), measure); + }; + + run( + s_main_target, elements.get(), main_target_output.get(), true /*into*/, VERBOSE /*=measure*/, "to-montgomery", 1); + run( + s_reference_target, elements.get(), ref_target_output.get(), true /*into*/, VERBOSE /*=measure*/, "to-montgomery", + 1); + ASSERT_EQ(0, memcmp(main_target_output.get(), ref_target_output.get(), N * sizeof(T))); + + run( + s_main_target, main_target_output.get(), main_target_output.get(), false /*from*/, VERBOSE /*=measure*/, + "to-montgomery", 1); + run( + s_reference_target, ref_target_output.get(), ref_target_output.get(), false /*from*/, VERBOSE /*=measure*/, + "to-montgomery", 1); + ASSERT_EQ(0, memcmp(main_target_output.get(), ref_target_output.get(), N * sizeof(T))); } }; -TEST_F(CurveApiTest, MSM) { MSM_test(); } -TEST_F(CurveApiTest, MontConversion) { mont_conversion_test(); } +TEST_F(CurveApiTest, msm) { MSM_test(); } +TEST_F(CurveApiTest, MontConversionAffine) { mont_conversion_test(); } +TEST_F(CurveApiTest, MontConversionProjective) { mont_conversion_test(); } #ifdef G2 -TEST_F(CurveApiTest, MSM_G2) { MSM_test(); } -TEST_F(CurveApiTest, MontConversionG2) { mont_conversion_test(); } +TEST_F(CurveApiTest, msmG2) { MSM_test(); } +TEST_F(CurveApiTest, MontConversionG2Affine) { mont_conversion_test(); } +TEST_F(CurveApiTest, MontConversionG2Projective) { mont_conversion_test(); } #endif // G2 #ifdef ECNTT diff --git a/icicle_v3/tests/test_field_api.cpp b/icicle_v3/tests/test_field_api.cpp index a1f95969b..46eb96d11 100644 --- a/icicle_v3/tests/test_field_api.cpp +++ b/icicle_v3/tests/test_field_api.cpp @@ -179,7 +179,7 @@ TYPED_TEST(FieldApiTest, montgomeryConversion) START_TIMER(MONTGOMERY) for (int i = 0; i < iters; ++i) { - ICICLE_CHECK(scalar_convert_montgomery(inout, N, true /*into montgomery*/, config, inout)); + ICICLE_CHECK(convert_montgomery(inout, N, true /*into montgomery*/, config, inout)); } END_TIMER(MONTGOMERY, oss.str().c_str(), measure); };