Skip to content

Commit

Permalink
updated msm for Icicle 1.0
Browse files Browse the repository at this point in the history
  • Loading branch information
@svpolonsky
svpolonsky committed Jan 10, 2024
1 parent bdbfb34 commit 9c1d893
Show file tree
Hide file tree
Showing 2 changed files with 117 additions and 103 deletions.
84 changes: 49 additions & 35 deletions benchmarks/msm/benchmark.cu
View file
Original file line number Diff line number Diff line change
@@ -1,39 +1,33 @@
#define CURVE_BN254 1
#define CURVE_BLS12_381 2
#define CURVE_BLS12_377 3
// #define CURVE_BN254 1
// #define CURVE_BLS12_381 2
// #define CURVE_BLS12_377 3

#define CURVE CURVE_BN254
// #define CURVE_ID CURVE_BN254

#include <stdio.h>
#include <iostream>
#include <string>
#include <cuda_runtime.h>
// #include <cuda_runtime.h>
#include <nvml.h>
#include <benchmark/benchmark.h>
#include "icicle/primitives/field.cuh"
#include "icicle/utils/storage.cuh"
#include "icicle/primitives/projective.cuh"
// #include "icicle/primitives/field.cuh"
// #include "icicle/utils/storage.cuh"
// #include "icicle/primitives/projective.cuh"

#include "icicle/appUtils/msm/msm.cu"

#if CURVE == CURVE_BN254

#include "icicle/curves/bn254/curve_config.cuh"
using namespace BN254;
const std::string curve = "BN254";
// select the curve
#define CURVE_ID 1

#elif CURVE == CURVE_BLS12_381
// include MSM template
#include "icicle/appUtils/msm/msm.cu"
using namespace curve_config;

#include "icicle/curves/bls12_381/curve_config.cuh"
using namespace BLS12_381;
#if CURVE_ID == BN254
const std::string curve = "BN254";
#elif CURVE_ID == BLS12_381
const std::string curve = "BLS12-381";

#elif CURVE == CURVE_BLS12_377

#include "icicle/curves/bls12_377/curve_config.cuh"
using namespace BLS12_377;
#elif CURVE_ID == BLS12_377
const std::string curve = "BLS12-377";

#endif

const unsigned max_msm_size = 1<<22;
Expand All @@ -54,11 +48,31 @@ cudaStream_t stream;

static void BM_msm(benchmark::State& state) {
const uint32_t msm_size=state.range(0);
bool on_device = true;
bool big_triangle = false;
// Create a device context
device_context::DeviceContext ctx = {
stream, // stream
0, // device_id
0, // mempool
};
// Create a MSM configuration
msm::MSMConfig config = {
ctx, // DeviceContext
0, // points_size
1, // precompute_factor
0, // c
0, // bitsize
10, // large_bucket_factor
1, // batch_size
true, // are_scalars_on_device
false, // are_scalars_montgomery_form
true, // are_points_on_device
false, // are_points_montgomery_form
true, // are_results_on_device
false, // is_big_triangle
false, // is_async
};
for (auto _ : state) {
large_msm<scalar_t, projective_t, affine_t>(scalars_d, points_d, msm_size, result_d, on_device, big_triangle, bucket_factor, stream);
// cudaDeviceSynchronize();
msm::MSM<scalar_t, affine_t, projective_t>(scalars_d, points_d, msm_size, config, result_d);
}
unsigned int power;
nvmlDeviceGetPowerUsage(device, &power);
Expand Down Expand Up @@ -92,21 +106,21 @@ int main(int argc, char** argv) {
std::cout << gpu_full_name << std::endl;
std::string gpu_name = gpu_full_name;
int gpu_clock_mhz = deviceProperties.clockRate/1000.;

nvmlInit();
nvmlDeviceGetHandleByIndex(0, &device); // for GPU 0

std::cout << "Setting host data" << std::endl;

scalars = (scalar_t*) malloc(sizeof(scalar_t) * max_msm_size);
points = (affine_t*)malloc(sizeof(affine_t) * max_msm_size);
for (unsigned i = 0; i < max_msm_size; i++) {
points[i] = (i % max_msm_size < 10) ? projective_t::to_affine(projective_t::rand_host()) : points[i - 10];
scalars[i] = scalar_t::rand_host();
}
scalar_t::RandHostMany(scalars, max_msm_size);
projective_t::RandHostManyAffine(points, max_msm_size);

std::cout << "Moving data to device" << std::endl;
// for (unsigned i = 0; i < max_msm_size; i++) {
// points[i] = (i % max_msm_size < 10) ? projective_t::to_affine(projective_t::rand_host()) : points[i - 10];
// scalars[i] = scalar_t::rand_host();
// }

std::cout << "Moving data to device" << std::endl;
cudaMalloc(&scalars_d, sizeof(scalar_t) * max_msm_size);
cudaMalloc(&points_d, sizeof(affine_t) * max_msm_size);
cudaMalloc(&result_d, sizeof(projective_t));
Expand All @@ -115,7 +129,6 @@ int main(int argc, char** argv) {


std::cout << "Running benchmark" << std::endl;

cudaStreamCreate(&stream);

// Run all benchmarks
Expand All @@ -130,6 +143,7 @@ int main(int argc, char** argv) {
::benchmark::AddCustomContext("coefficient_C", std::to_string(bucket_factor));
::benchmark::RunSpecifiedBenchmarks();

std::cout << "Cleaning up" << std::endl;
cudaFree(scalars_d);
cudaFree(points_d);
cudaFree(result_d);
Expand Down
136 changes: 68 additions & 68 deletions benchmarks/msm/benchmark.json
View file
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
{
"context": {
"date": "2023-12-12T18:47:36+00:00",
"host_name": "0f4f6719d34b",
"date": "2024-01-10T18:31:05+00:00",
"host_name": "9bb077e68cdf",
"executable": "/icicle-benchmark/build/benchmark",
"num_cpus": 12,
"mhz_per_cpu": 4208,
Expand Down Expand Up @@ -32,7 +32,7 @@
"num_sharing": 6
}
],
"load_avg": [0.892578,1.04199,1.06738],
"load_avg": [1.67236,1.54883,1.26611],
"library_build_type": "release",
"coefficient_C": "12",
"comment": "on-device API",
Expand All @@ -52,12 +52,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3934,
"real_time": 1.0719420736669561e-02,
"cpu_time": 1.0706253246314182e-02,
"iterations": 3850,
"real_time": 1.0958126205981213e-02,
"cpu_time": 1.0935583719740259e-02,
"time_unit": "s",
"PowerUsage": 1.7900000000000000e+02,
"Temperature": 5.8000000000000000e+01
"PowerUsage": 1.8300000000000000e+02,
"Temperature": 6.0000000000000000e+01
},
{
"name": "BM_msm/2048/min_time:30.000",
Expand All @@ -68,12 +68,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3905,
"real_time": 1.0774707772873971e-02,
"cpu_time": 1.0763022053777207e-02,
"iterations": 3833,
"real_time": 1.0995332787121713e-02,
"cpu_time": 1.0964406394469089e-02,
"time_unit": "s",
"PowerUsage": 1.8200000000000000e+02,
"Temperature": 6.1000000000000000e+01
"PowerUsage": 1.8700000000000000e+02,
"Temperature": 6.3000000000000000e+01
},
{
"name": "BM_msm/4096/min_time:30.000",
Expand All @@ -84,12 +84,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3901,
"real_time": 1.0820364909526127e-02,
"cpu_time": 1.0802259874647524e-02,
"iterations": 3837,
"real_time": 1.1030317013532288e-02,
"cpu_time": 1.1005096382851187e-02,
"time_unit": "s",
"PowerUsage": 1.8400000000000000e+02,
"Temperature": 6.2000000000000000e+01
"PowerUsage": 1.8900000000000000e+02,
"Temperature": 6.4000000000000000e+01
},
{
"name": "BM_msm/8192/min_time:30.000",
Expand All @@ -100,12 +100,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3592,
"real_time": 1.1705802419521837e-02,
"cpu_time": 1.1692981692928735e-02,
"iterations": 3520,
"real_time": 1.1957617815344499e-02,
"cpu_time": 1.1933407338352272e-02,
"time_unit": "s",
"PowerUsage": 1.8600000000000000e+02,
"Temperature": 6.3000000000000000e+01
"PowerUsage": 1.9000000000000000e+02,
"Temperature": 6.5000000000000000e+01
},
{
"name": "BM_msm/16384/min_time:30.000",
Expand All @@ -116,12 +116,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3586,
"real_time": 1.1728714219186968e-02,
"cpu_time": 1.1715881023703301e-02,
"iterations": 3508,
"real_time": 1.1985734673315169e-02,
"cpu_time": 1.1963134809293039e-02,
"time_unit": "s",
"PowerUsage": 1.8900000000000000e+02,
"Temperature": 6.4000000000000000e+01
"PowerUsage": 1.9400000000000000e+02,
"Temperature": 6.5000000000000000e+01
},
{
"name": "BM_msm/32768/min_time:30.000",
Expand All @@ -132,12 +132,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3571,
"real_time": 1.1769651148988034e-02,
"cpu_time": 1.1759971484738174e-02,
"iterations": 3497,
"real_time": 1.2017841417779722e-02,
"cpu_time": 1.1998319173863311e-02,
"time_unit": "s",
"PowerUsage": 1.9500000000000000e+02,
"Temperature": 6.4000000000000000e+01
"PowerUsage": 2.0000000000000000e+02,
"Temperature": 6.6000000000000000e+01
},
{
"name": "BM_msm/65536/min_time:30.000",
Expand All @@ -148,12 +148,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3235,
"real_time": 1.2998765338193538e-02,
"cpu_time": 1.2985645372488402e-02,
"iterations": 3199,
"real_time": 1.3154686054723730e-02,
"cpu_time": 1.3129190219443565e-02,
"time_unit": "s",
"PowerUsage": 2.0000000000000000e+02,
"Temperature": 6.5000000000000000e+01
"PowerUsage": 2.0600000000000000e+02,
"Temperature": 6.7000000000000000e+01
},
{
"name": "BM_msm/131072/min_time:30.000",
Expand All @@ -164,12 +164,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 3079,
"real_time": 1.3650311346873023e-02,
"cpu_time": 1.3640242942838589e-02,
"iterations": 3068,
"real_time": 1.3715011163613476e-02,
"cpu_time": 1.3691055463168176e-02,
"time_unit": "s",
"PowerUsage": 2.1200000000000000e+02,
"Temperature": 6.5000000000000000e+01
"PowerUsage": 2.1700000000000000e+02,
"Temperature": 6.7000000000000000e+01
},
{
"name": "BM_msm/262144/min_time:30.000",
Expand All @@ -180,12 +180,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 2777,
"real_time": 1.5142664100817509e-02,
"cpu_time": 1.5132199924738948e-02,
"iterations": 2720,
"real_time": 1.5460519887881888e-02,
"cpu_time": 1.5432111724264697e-02,
"time_unit": "s",
"PowerUsage": 2.3100000000000000e+02,
"Temperature": 6.7000000000000000e+01
"PowerUsage": 2.3500000000000000e+02,
"Temperature": 6.9000000000000000e+01
},
{
"name": "BM_msm/524288/min_time:30.000",
Expand All @@ -196,12 +196,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 2270,
"real_time": 1.8516216871363525e-02,
"cpu_time": 1.8493770291189423e-02,
"iterations": 2239,
"real_time": 1.8804006550705219e-02,
"cpu_time": 1.8765309853059419e-02,
"time_unit": "s",
"PowerUsage": 2.5500000000000000e+02,
"Temperature": 6.9000000000000000e+01
"PowerUsage": 2.6000000000000000e+02,
"Temperature": 7.1000000000000000e+01
},
{
"name": "BM_msm/1048576/min_time:30.000",
Expand All @@ -212,12 +212,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 1662,
"real_time": 2.5298198188919010e-02,
"cpu_time": 2.5272283267749707e-02,
"iterations": 1633,
"real_time": 2.5804960957143186e-02,
"cpu_time": 2.5770517832823046e-02,
"time_unit": "s",
"PowerUsage": 2.8900000000000000e+02,
"Temperature": 7.1000000000000000e+01
"PowerUsage": 2.9200000000000000e+02,
"Temperature": 7.2000000000000000e+01
},
{
"name": "BM_msm/2097152/min_time:30.000",
Expand All @@ -228,12 +228,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 1073,
"real_time": 3.9168599447368266e-02,
"cpu_time": 3.9118837269338277e-02,
"iterations": 1052,
"real_time": 3.9985827347923233e-02,
"cpu_time": 3.9931135256654024e-02,
"time_unit": "s",
"PowerUsage": 2.8900000000000000e+02,
"Temperature": 6.9000000000000000e+01
"PowerUsage": 2.9500000000000000e+02,
"Temperature": 6.7000000000000000e+01
},
{
"name": "BM_msm/4194304/min_time:30.000",
Expand All @@ -244,12 +244,12 @@
"repetitions": 1,
"repetition_index": 0,
"threads": 1,
"iterations": 628,
"real_time": 6.7063869372719104e-02,
"cpu_time": 6.7002969555732456e-02,
"iterations": 612,
"real_time": 6.8616608176585839e-02,
"cpu_time": 6.8533704723856226e-02,
"time_unit": "s",
"PowerUsage": 2.9500000000000000e+02,
"Temperature": 6.6000000000000000e+01
"PowerUsage": 2.9400000000000000e+02,
"Temperature": 6.5000000000000000e+01
}
]
}

0 comments on commit 9c1d893

Please sign in to comment.