Compare commits

..

158 Commits

Author SHA1 Message Date
stas
30d4558560 answers Roman's comments 2024-02-14 12:24:03 -05:00
stas
582107fc7c added c++ example Poseidon-hash 2024-02-08 17:36:56 -05:00
ImmanuelSegol
04b1b3dda5 refactor: add a basic example 2024-02-07 20:07:21 -04:00
Jeremy Felder
b6dded89cd Release v1.2.0 (#364)
Release v1.2.0:

- [FEAT] Add Poseidon hash as primitive
- [FEAT] Add Merkle tree using poseidon hash
- [BUG] Fix NTT overflow when using large cosets
2024-02-07 17:01:09 +02:00
Jeremy Felder
5a138367f8 (chore): bump rust crate versions (#362)
bump rust crate versions
2024-02-07 14:49:54 +02:00
ChickenLover
a3fc01d88d Implement Poseidon and TreeBuilder (#352)
* BW scalar field is now the same as BLS base field

* add poseidon

* add merkle tree builder

* poseidon rust bindings

* implement rust bindings

* add doc comments

* remove global poseidon constants

* add custom constants API and script for generating new constants

* add the rest of the curves for poseidon

* add all the curves for real

* misname bls12-377

* typo

* partial rounds

* minor fixes

* small tweak for big performance boost

* add CHK_INIT_IF_RETURN

---------

Co-authored-by: DmytroTym <dmytrotym1@gmail.com>
2024-02-07 14:49:54 +02:00
DmytroTym
d84cab79fd Changed long to int64 2024-02-07 14:49:54 +02:00
DmytroTym
46d7b88f6e Fixed overflow in large coset NTTs 2024-02-07 14:49:54 +02:00
Jeremy Felder
77a7613aa2 Release v1.1.0 (#357)
Release v1.1.0:
- Updated examples to use the new API
- [c++] Curve specific functions using macros
- [c++] Consolidate MSM and Batch MSM to single function
- [CI] Add codespell in CI
- [FIX] Windows rust build
- [FIX] G2 on rust bindings
- [FIX] Bw6 using bls12377
2024-01-31 20:30:44 +02:00
Jeremy Felder
5a96f9937d Bump for release 2024-01-31 16:34:14 +02:00
Stas
aaa3808c81 Update examples for new api (#355)
## Describe the changes

Make sure the examples comply with new API
2024-01-31 08:59:21 +02:00
stas
759b7b26d6 Update examples to use latest API 2024-01-31 08:59:21 +02:00
Leon Hibnik
1874ade68a update readme links (#346)
Update README.md
2024-01-24 12:26:38 +02:00
DmytroTym
96fe5bf283 G2 fix and BW6 scalar field on the Rust side (#341)
* BW scalar field is now the same as BLS base field in Rust

* G2 fixed and added into Rust
2024-01-24 11:51:22 +02:00
Jeremy Felder
f0a0bb5974 Fix/windows build failing (#345)
- Make the curve config's omegas_count conditionally accessed when creating fields
- Remove the extern C function that returns a UDT containing non-POD types and replace it with a default_config function on the Rust bindings side
2024-01-23 10:51:16 +02:00
Jeremy Felder
69af0bef91 [FEAT]: Add codespell to CI and pre-commit hooks (#344)
Add codespell to pre-commit hook/CI and fix typos
2024-01-22 14:27:52 +02:00
Stas
45f6db666b c++ msm uses new API (#338)
## Describe the changes

This PR allows c++ MSM example to compile with new API

## Linked Issues

Resolves #
2024-01-18 11:31:30 -05:00
Stas
4c235bf5f5 Merge branch 'dev' into examples-cpp-msm-new-API 2024-01-18 11:29:54 -05:00
stas
a0f35ea8cd comply with reviewer's comments 2024-01-18 11:13:12 -05:00
yshekel
56fcd2e156 refactor: consolidate msm and batch-msm implementations to one function (#342)
refactor: consolidate msm and batch-msm implementations to one function
    - now batch-msm support parallel BM accumulation in addition to large triangle accumulation
2024-01-18 14:20:34 +02:00
Otsar
c1a32a9879 Update README.md (#339)
Added badge
2024-01-11 18:34:32 +02:00
yshekel
67586e01f8 refactor: generate curve-specific function names with macro (#337)
For exposed functions that are curve-specific, the expected symbol is
based on the curve name.
For example MSMCuda becomes bn254MSMCuda for CURVE=bn254.

currently it is implemented via objcopy by redefining symbols after
compilation.

This PR modifies the function names at preprocessing time instead.
For shared objects, objcopy doesn't work (since it cannot modify the
dynsym section).
2024-01-11 15:04:40 +02:00
Yuval Shekel
9823f87fe2 fix: enable ci for dev branch 2024-01-11 13:49:41 +02:00
Yuval Shekel
5fda751e93 refactor: generate curve-specific function names with macro instead of using objcopy to modify the symbols 2024-01-11 10:54:31 +02:00
stas
0c58fe2b83 c++ msm uses new API 2024-01-09 11:06:10 -05:00
Leon Hibnik
b184befc99 Update readme v1.0.0 (#335)
* Update readme v1.0.0

* update

* update example readme
2024-01-08 23:48:24 +02:00
DmytroTym
91471fbbc6 New API for Version 1.0.0 (#326) 2024-01-08 20:42:16 +02:00
Jeremy Felder
4cc30474d4 Added conditional PIC for compilation 2024-01-08 18:22:50 +02:00
Jeremy Felder
ad932c6b4a Merge branch 'main' into dev 2024-01-08 17:44:03 +02:00
Jeremy Felder
392fff4e8f fix: remove examples directory from clang formatting 2024-01-08 17:40:40 +02:00
DmytroTym
dc3518c0de Smart pointers and documentation (#333)
* Safer smart pointer that covers host and device

* Fixed MSM test

* Scalars and points in MSM are non-mutable in all cases

* change mont API (#332)

* Some Rust doc comments

---------

Co-authored-by: ChickenLover <Romangg81@gmail.com>
2024-01-08 17:35:27 +02:00
Jeremy Felder
79b86e0edf Added examples for rust and c++ 2024-01-08 17:31:46 +02:00
Jeremy Felder
709009a96a Formatting, fixes tests, and general cleanup 2024-01-08 17:31:46 +02:00
Jeremy Felder
b92bc707d1 Fix main build badge to be on push and not pr 2024-01-08 17:28:11 +02:00
Jeremy Felder
a87f2251da Moved ingo_<curve> output to same directory as original library output allowing rust build script to find it when crate is used as a dependency 2024-01-08 17:28:03 +02:00
ImmanuelSegol
31108ee31d Update READMEs and docs (#303)
---------

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
2024-01-08 17:22:36 +02:00
DmytroTym
5f7b36d427 Rust NTT updates (#305)
* NTT and MSM bugs fixed and functionality extended

* More Rust tests, refactored Rust wrappers

* Reworked MSM and NTT public functions on the Rust side

---------

Co-authored-by: ImmanuelSegol <3ditds@gmail.com>
Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
Co-authored-by: BigSky77 <simonjudd2@gmail.com>
Co-authored-by: yanziseeker <153156292+AdventureSeeker987@users.noreply.github.com>
Co-authored-by: BigSky <77446076+bigsky77@users.noreply.github.com>
2024-01-08 17:20:39 +02:00
ChickenLover
41affcdadf Add Montgomery conversions API (#321) 2024-01-08 17:19:46 +02:00
VitaliiH
a22c046410 Adds proper error handling for GPU errors and c++ icicle errors 2024-01-08 17:19:45 +02:00
ChickenLover
0eb8560c63 Feat/roman/msm ntt generics (#299)
Add curves to Rust, use generics for NTT, MSM and tests
2024-01-08 17:19:02 +02:00
Jeremy Felder
605c25f9d2 Fix main build badge to be on push and not pr (#325)
## Describe the changes

This PR fixes the build badge to show the status of the `main` branch
build status instead of the status of a PR branch against `main`
2024-01-01 18:45:14 +02:00
Jeremy Felder
5ccf7ff378 Fix main build badge to be on push and not pr 2023-12-31 12:34:17 +02:00
Jeremy Felder
4beda3a900 Fix curve crates not building when used as dependency (#320)
## Describe the changes

This PR moves `ingo_<curve>` archive output to the same directory as the
original library output allowing rust build scripts to find it when a
curve crate is used as a dependency
2023-12-25 14:19:07 +02:00
Jeremy Felder
da122e7ec9 Moved ingo_<curve> output to same directory as original library output allowing rust build script to find it when crate is used as a dependency 2023-12-25 10:11:50 +02:00
Jeremy Felder
f9e46d158c Fix typos (#319)
## Describe the changes
Fix some typos in comment
2023-12-24 09:25:07 +02:00
BigSky
65d4f2ba2e Add Nix Shell Environment Configuration for CUDA (#318)
This PR introduces a Nix Shell configuration for easing the development
of the ICICLE Core using Nix or NixOS.

Changes:

1. Added instructions on how to use Nix Shell to create a development
environment with all required dependencies and environmental variables.
2. Created `cuda-shell.nix` to configure Nix Shell with necessary CUDA
related packages and environmental variables. Specifically, setting the
`PATH`, `LD_LIBRARY_PATH`, `CUDA_PATH`, `CPATH`, `LIBRARY_PATH`, and
`CMAKE_CUDA_COMPILER` associated with the CUDA toolkit.

This configuration allows us to build and run ICICLE Core more
efficiently in Nix environment.
2023-12-22 10:30:36 -07:00
yanziseeker
263b5d95b5 Fix typos 2023-12-22 06:30:28 +00:00
BigSky77
64c8c89424 Update Nix Shell, add Nix instructions 2023-12-19 12:23:04 -07:00
ImmanuelSegol
dcaa0b4513 update readme (#303)
* refactor

* refactor

* refactor: add rust docs

* Update README.md

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>

* Update README.md

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>

* refactor

* Update icicle/README.md

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>

* Update wrappers/rust/README.md

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>

* Update README.md

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>

* refactor

* remove

---------

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
2023-12-19 09:17:30 -04:00
Jeremy Felder
d5dd16dd9b Fix the need for a dummy bindings.rs file included in the repo (#296) 2023-12-13 12:45:00 +02:00
Jeremy Felder
644aa93d3f New API (#293)
This PR revamps the API and implements the rust bindings for the new API.
2023-12-12 22:30:53 +02:00
Jeremy Felder
73e39c60f6 Fix bw6 test in CI 2023-12-12 21:33:23 +02:00
Jeremy Felder
6d992a9909 Fix CI format workflow 2023-12-12 21:12:20 +02:00
Jeremy Felder
0df6262961 Update CI and prepush hooks 2023-12-12 20:49:57 +02:00
DmytroTym
401ec5afac Merge in 'main' and fmt 2023-12-12 20:08:38 +02:00
DmytroTym
629d8cb309 Update Rust APIs (#292)
* NTT reworked on CUDA side, Rust API updated

* Updates to the Rust-side CUDA runtime wrapper

* Merged in main
2023-12-12 19:31:36 +02:00
DmytroTym
f8610dd5b6 Improved modular multiplier (#289)
Out implementation of Barrett modular multiplication improved by utilising Karatsuba multiplication and more careful optimisations of lsb and msb multipliers in reduction stage
2023-12-05 13:11:44 +02:00
DmytroTym
dfa5b10adb Update Rust apis (#262)
* fix memory error in single_stage_multi_reduction_kernel (#235)

* refactor

* refactor

* revert

* refactor: clang format

* Update icicle/appUtils/msm/msm.cu

* Added separate device context struct, returned lde

* wip - msm and eq

* added lde to cmake

* Montgomery param added in lde.cu mul function

* fixed on_device for ntt and lde

* CamelCase

* fixed msm_test, int unification, google guilde

* wip - ntt crash debugging

* async MSM with a rust wrapper

* wip ntt tests with corretness

* hotfix for correctness > 2^9

* wip on device inout mixing with correctness

* cleanup

* preserving twiddles after first call

* fixed twiddles preserving

* formatting

* removed some printing

* disable ecntt temporarily

* format

* rust fmt

* exclude target from format

* passing ntt after merge

* hotfix for linking issue

* format

* format

* draft of pr comments + correctness restored

* wip refactor + format

* domain wip

* rust format

* Merged feature branch in and Rust MSM correctness

* rust build for correct curve

* Slowdown fixed by passing release flag to cmake

* WIP field and curve

* still wip field and curve

* field and curve in rust 1.0

* Refactored rust into several crates

* Arkworks is now an option, bn254 crate created

* Rust msm and ntt wip

* A version of rust msm done, cuda runtime wrapped

* refactor rust by creating a curve folder

* vec_ops instead of lde for now

* format

---------

Co-authored-by: ImmanuelSegol <3ditds@gmail.com>
Co-authored-by: Vitalii <vitalii@ingonyama.com>
2023-12-03 13:32:50 +02:00
BigSky
fad317ac77 Docs: Clarify enabling tests in cmake build process (#288)
Docs: Emphasize enabling tests in cmake build process
2023-11-28 14:43:29 +02:00
Jeremy Felder
856629d6c8 Rust/large bucket factor msm (#271)
* Update rust bindings to support large_bucket_factor parameter
* Special treatment of ones in MSM removed

---------
Co-authored-by: DmytroTym <dmytrotym1@gmail.com>
2023-11-26 14:10:30 +02:00
Jeremy Felder
2790f180d6 Update feature_request.md (#278) 2023-11-26 14:09:58 +02:00
Jeremy Felder
5813f2955a Update bug_issue.md
Update bug issue template to apply correct label
2023-11-26 12:22:38 +02:00
VitaliiH
7baea7cc5a separable compilation for Rust (#244)
separable compilation for Rust #244
2023-11-16 08:44:56 +01:00
ImmanuelSegol
29cad66ba6 include colab (#261) 2023-11-13 08:16:41 +02:00
DmytroTym
e4e9130340 Two curve NTT correctness issue hotfix (#254) 2023-11-05 08:24:12 +02:00
omahs
5c868abcf9 Fix typos (#257) 2023-11-02 16:57:55 +02:00
vuittont60
a469fb577b fix typos in logs (#255) 2023-11-02 16:57:34 +02:00
Vitalii
133a1b28bc format 2023-10-24 23:49:19 +02:00
Vitalii
c666786fa1 exclude target from format 2023-10-24 23:47:32 +02:00
Vitalii
b108234ce4 formatting 2023-10-24 22:11:06 +02:00
liuxiao
fd62fe5ae8 Support bw6-761 (#188)
Resolves #191 and #113

---------

Co-authored-by: DmytroTym <dmytrotym1@gmail.com>
Co-authored-by: ImmanuelSegol <3ditds@gmail.com>
2023-10-21 18:49:06 +03:00
Jeremy Felder
09d8c5da6a Fix readme badge links, fix CI cpp formatter (#249) 2023-10-17 17:06:11 +03:00
Jeremy Felder
88c9c8584f [CI]: Add concurrency groups at workflow level (#238)
Add concurrency groups at workflow level for CI. Remove dev CI since we no longer use dev branch. Resolves #180
2023-10-17 16:02:31 +03:00
Jeremy Felder
1cf7b2e4ba Exclude target directory from format checks (#247) 2023-10-16 15:56:58 +03:00
DmytroTym
a646c81aaa Linting as per google guide (#241)
* fixed msm_test, int unification, google guilde

* Fixed warnings
2023-10-12 20:22:24 +03:00
DmytroTym
028bed11fa Hotfix to go regression when 2 curves are imported (#245)
Hotfix to slowdown in go when more than one curve is imported.
2023-10-12 15:53:20 +03:00
DmytroTym
d4dd6d3923 Cuda refactoring (#240)
* fix memory error in single_stage_multi_reduction_kernel (#235)

* Added separate device context struct, returned lde

* Montgomery param added in lde.cu mul function

---------

Co-authored-by: ImmanuelSegol <3ditds@gmail.com>
2023-10-04 14:51:59 +03:00
ImmanuelSegol
9114ecb269 fix memory error in single_stage_multi_reduction_kernel (#235)
* refactor

* refactor

* revert

* refactor: clang format

* Update icicle/appUtils/msm/msm.cu
2023-10-03 15:22:28 +03:00
DmytroTym
fc73f274fa remove boost from cmake 2023-10-02 12:13:37 +03:00
DmytroTym
c1dfbd28ad Merge branch 'main' into cuda_refactoring 2023-10-02 12:06:23 +03:00
DmytroTym
2d3ab73eca error codes and extern ntt functions added 2023-09-29 23:10:36 +03:00
DmytroTym
9f884f3022 functional NTT 1.0 2023-09-28 22:28:42 +03:00
Jeremy Felder
97f0079e5c Fix: div by 0 when number of Elements is 1 (#230) 2023-09-28 16:11:17 +03:00
ImmanuelSegol
9f6707581e Make dependency instructions more clear (#227) 2023-09-27 12:25:26 +03:00
DmytroTym
a9c846c44c Draft ntt.cuh 2023-09-27 10:28:02 +03:00
Jeremy Felder
413e1d8b60 [CI]: Adds C++/CUDA CI and conditional workflow runs (#223)
* Add cmake tests for cpp primitives

* Add cpp/cuda formatting

* Add conditional steps based on files changed for faster required checks

* Update runs-on for check files changed
2023-09-26 14:41:06 +03:00
ImmanuelSegol
8612975e36 update docs (#222)
* refactor

* lint

* Typo and spacing

---------

Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
2023-09-26 11:46:11 +03:00
Leon Hibnik
0a639783e2 add mont_r and mont_r_inv and fix args (#187)
Resolves #186
2023-09-26 09:45:10 +03:00
DmytroTym
8068f16d82 Cosmetic changes to the MSM API and build process 2023-09-23 16:50:15 +03:00
ImmanuelSegol
e368f00bb8 Fix CI - dont use deprecated package name (#216)
Update icicle crate name in examples and benches
2023-09-21 16:03:44 +07:00
DmytroTym
82743a098f 1.0 MSM refactoring on CUDA side, cmake build 2023-09-20 19:39:09 +03:00
ImmanuelSegol
08862134e8 fix cuda test - remove boost (#197)
* refactor: remove boost
2023-09-20 10:00:34 +03:00
DmytroTym
d8d2df1cd0 Merge branch 'main' into cuda_refactoring 2023-09-13 15:59:05 +03:00
DmytroTym
323430cddc MSM API prototype 2023-09-13 15:58:48 +03:00
Jeremy Felder
04e5ff5d1a Update root of unity and regenerate omegas (#181)
updates the root of unity and regenerates the omegas for BLS12377
adds an option for the new_curve script to only update the params.cuh file instead of regenerating everything
2023-09-07 08:23:43 +03:00
Jeremy Felder
81e40a1001 Merge pull request #172 from ingonyama-zk/dev 2023-09-05 10:52:10 +03:00
Leon Hibnik
a8c539740a Merge pull request #176 from ingonyama-zk/omershlo-patch-1
Update README.md
2023-09-03 11:09:22 +03:00
Leon Hibnik
1cf711215b Merge pull request #173 from ingonyama-zk/Otsar-Raikou-patch-1
Update README.md
2023-09-03 11:03:08 +03:00
Shlomtz
3d370a2be3 Update README.md [skip ci] 2023-09-01 15:01:39 +03:00
Shlomtz
49212c540c Update README.md
update hall of fame
2023-09-01 13:47:21 +03:00
Otsar
be68cddd1a Update README.md 2023-08-31 14:09:44 +03:00
Jeremy Felder
5667f32bfe Merge branch 'main' into dev 2023-08-31 09:19:44 +03:00
ImmanuelSegol
9ea6bc0bad Update Cargo.toml (#161) 2023-08-31 09:05:21 +03:00
Jeremy Felder
fb52650bbc CI: Additional checks (#155)
adds CI checks for building and testing Golang bindings
adds CI checks for formatting Rust and Golang files
Fixes Golang tests for BN254
Splits Actions checks for PR against main into multiple files

Resolves #108
Resolves #107
Resolves #138
2023-08-31 09:04:53 +03:00
Jeremy Felder
ca8961501e Adding changes from main to dev for clean merge back into main (#170) 2023-08-29 15:53:24 +03:00
DmytroTym
78e20f9add Minimal correct MSM (#162) 2023-08-28 10:14:54 +03:00
Jeremy Felder
dc6893732b Merge pull request #166 from weijiekoh/fix/readme-links
Fix broken links in the readme, resolves #164
2023-08-28 09:22:57 +03:00
Koh Wei Jie
175109a070 Fix link to CRV_CONFIG
Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
2023-08-27 15:48:41 -07:00
Jeremy Felder
b216287de1 Merge pull request #163 from weijiekoh/main
Support cmake versions below 3.24.0
2023-08-27 14:25:28 +03:00
ImmanuelSegol
7a2fa20da7 Remove decimation from API (#165)
Resolves #154
2023-08-27 14:08:56 +03:00
Koh Wei Jie
8e7799b632 fixed broken CRV_TEMPLATE and CRV_CONFIG links in readme 2023-08-24 12:55:15 -07:00
Koh Wei Jie
6dd7722f5d changed minimum cmake version in readme 2023-08-24 10:03:17 -07:00
ImmanuelSegol
27627ed2c1 refactor (#158) 2023-08-24 12:02:43 +03:00
Koh Wei Jie
3284cd8dce updated readme with prerequisites section; added conditional in icicle/CMakeLists.txt to support cmake versions below 3.24 2023-08-23 14:23:46 -07:00
Jeremy Felder
8c7b2bb24a Add citation to repo (#156) 2023-08-21 09:25:24 +03:00
Jeremy Felder
b6c87c3fd8 Fix formatting for all files (#153) 2023-08-20 11:35:28 +03:00
Leon Hibnik
e04bd928e6 Merge pull request #145 from ingonyama-zk/fix/goicicle-setup-script
setup.sh update
2023-08-17 12:08:24 +03:00
Leon Hibnik
cb6ed6af59 Merge branch 'dev' into fix/goicicle-setup-script 2023-08-17 12:07:07 +03:00
Jeremy Felder
9ea3350589 Add language formatters (#132) 2023-08-17 09:41:58 +03:00
Leon Hibnik
f38a9a322c Merge branch 'dev' into fix/goicicle-setup-script 2023-08-16 16:43:32 +03:00
ImmanuelSegol
ad1e482252 missing functions (#152) 2023-08-16 16:38:20 +03:00
Jeremy Felder
273bd536db Merge branch 'dev' into fix/goicicle-setup-script 2023-08-16 16:14:36 +03:00
Jeremy Felder
1463edc413 CI: Run linux on self-hosted, Make windows download smaller and remove caching (#150) (#151) 2023-08-16 16:14:22 +03:00
Jeremy Felder
db93204dc7 CI: Run linux on self-hosted, Make windows download smaller and remove caching (#150) 2023-08-16 15:16:32 +03:00
ImmanuelSegol
e1b692b8ed Merge branch 'dev' into fix/goicicle-setup-script 2023-08-16 12:59:44 +03:00
Leon Hibnik
e6416f4110 Merge pull request #146 from ingonyama-zk/fix/zeroedgecase
bucket_method_msm - address 0 edge case
2023-08-16 12:00:46 +03:00
ImmanuelSegol
96facd58d5 refactor: dont throw error when all scalars are 0 2023-08-15 19:53:45 +03:00
Leon Hibnik
11fe11b071 Merge branch 'dev' into fix/goicicle-setup-script 2023-08-15 14:16:54 +03:00
DmytroTym
19d0730aad Correct MSM for weird scalar distributions (#143) 2023-08-15 13:14:46 +03:00
LeonHibnik
36133ba26c setup.sh update 2023-08-15 12:14:06 +03:00
ImmanuelSegol
a1d9fa6648 fix - add missing go wrappers for all curves + add missing constants for curves (#130) 2023-08-14 13:15:56 +03:00
Vitalii Hnatyk
2f21ec4aa7 large_msm compilation hotfix (#131)
hotfix for missing parameter in large_msm
2023-07-27 10:05:45 +02:00
Jeremy Felder
5b504c44b7 Fix badges 2023-07-20 08:57:50 +03:00
Jeremy Felder
d13143506e writing .so file requires sudo 2023-07-19 21:44:12 +03:00
ImmanuelSegol
94c73e637c Fix/cudacodegoicile (#128)
* refactor

* refactor

* Refactor

* Refactor

* refactor: add sh script

* refactor

* refactor

* refactor: fix path
2023-07-19 21:30:20 +03:00
Jeremy Felder
b71b041561 Integrate msm performance improvements (#129) 2023-07-19 16:32:59 +03:00
ImmanuelSegol
a9d6ac0e27 move header file import (#127) 2023-07-19 08:33:05 +03:00
ImmanuelSegol
8a11a2f60e some minor changes (#125)
Co-authored-by: Jeremy Felder <jeremy.felder1@gmail.com>
2023-07-18 20:17:16 +03:00
ImmanuelSegol
ab69139ade Goicicle (#77) 2023-07-16 14:31:41 +03:00
Vitalii Hnatyk
7a8191bcb4 NTT improvements (shared mem + inplace) (#116)
Resolves #112
2023-07-16 13:56:20 +03:00
Jeremy Felder
e3f089f0f3 Fix: Docs, Curve generation script (#102)
Fix CUDA compilation docs and update new curve script to generate correct cu(h) files. Resolves #101
2023-06-28 15:55:33 +03:00
Jeremy Felder
cb61755c8b Add streams to poseidon (#105)
Adds streams to the poseidon implementation for BLS12-381. Resolves #91
2023-06-20 12:22:16 +03:00
Jeremy Felder
34a556ac85 Update build workflow (#104) 2023-06-19 14:21:46 +03:00
Jeremy Felder
2a3f5a258a Merge pull request #106 from gkigiermo/fix/cuda-test-suite
Resolves #103
2023-06-15 11:48:36 +03:00
Guillermo Oyarzun
9023daeb4f Add c++17 requirement to cmake 2023-06-15 10:06:02 +02:00
Guillermo Oyarzun
4d83ba101c Fix curve config location and link to some namespace struct members 2023-06-13 23:52:48 +02:00
ChickenLover
26f2f5c76c reduce memory consumption in hash_blocks (#100)
* reduce memory consumption in hash_blocks
2023-06-08 20:51:41 +07:00
Jeremy Felder
434ab70305 Fixed omega retrieval issue (#99)
Resolves #71
2023-06-08 11:47:41 +03:00
Jeremy Felder
cd7b81102c Merge pull request #95 from ingonyama-zk/dev-v2 2023-06-01 14:45:49 +03:00
Jeremy Felder
2917024ef7 Update build action to include dev branch (#96) 2023-06-01 13:56:30 +03:00
ChickenLover
24096b6cf7 Fix windows build (#97) 2023-06-01 13:19:51 +03:00
Jeremy Felder
20de60fd43 Add streams capability (#89) 2023-06-01 10:44:11 +03:00
DmytroTym
b5a24d8e4d G2 CUDA (#78)
G2 arithmetic and MSM implemented in CUDA
2023-06-01 10:44:05 +03:00
HadarIngonyama
9ebf3d4f34 MSM for large sizes (#88)
* bugs fixed

* bugs fixed

* remove short msm from extern call

* code cleaning
2023-06-01 10:40:13 +03:00
ChickenLover
43f8c01afe Feat/poseidon (#75)
Adds poseidon function to compute hashes over multiple preimages in parallel
2023-06-01 10:40:13 +03:00
Vitalii Hnatyk
a4c676e5b6 fix compilation - rust bench (#80)
Co-authored-by: Vitalii <vitalii@ingonyama.com>
2023-05-29 10:07:18 +03:00
guy-ingo
a467c9f56e Supporting Additional Curves (#72)
* init commit - changes for supporting new curves

* refactor + additional curve (bls12-377 works, bn254 - not yet)

* general refactor + curves script + fixing bn245

* revert unnecessary changes + refactor new curve script

* add README and fix limbs_p=limbs_q case in python script
2023-05-15 15:31:18 +03:00
Vitalii Hnatyk
20da109a92 Fix for local machines GoogleTest and CMake (#70)
GoogleTest fix, updated readme
2023-05-10 13:56:25 +03:00
334 changed files with 73096 additions and 16735 deletions

39
.clang-format Normal file
View File

@@ -0,0 +1,39 @@
Language: Cpp
AlignAfterOpenBracket: AlwaysBreak
AlignConsecutiveMacros: true
AlignTrailingComments: true
AllowAllParametersOfDeclarationOnNextLine: true
AllowShortBlocksOnASingleLine: true
AllowShortCaseLabelsOnASingleLine: false
AllowShortFunctionsOnASingleLine: All
AllowShortIfStatementsOnASingleLine: true
AlwaysBreakTemplateDeclarations: true
BinPackArguments: true
BinPackParameters: false
BreakBeforeBraces: Custom
BraceWrapping:
AfterClass: true
AfterFunction: true
BreakBeforeBinaryOperators: false
BreakBeforeTernaryOperators: true
ColumnLimit: 120
ContinuationIndentWidth: 2
Cpp11BracedListStyle: true
DisableFormat: false
IndentFunctionDeclarationAfterType: false
IndentWidth: 2
KeepEmptyLinesAtTheStartOfBlocks: false
MaxEmptyLinesToKeep: 1
NamespaceIndentation: All
PointerAlignment: Left
SortIncludes: false
SpaceBeforeAssignmentOperators: true
SpaceBeforeParens: ControlStatements
SpaceInEmptyParentheses: false
SpacesBeforeTrailingComments: 1
SpacesInAngles: false
SpacesInContainerLiterals: false
SpacesInCStyleCastParentheses: false
SpacesInParentheses: false
Standard: c++17
UseTab: Never

3
.codespellignore Normal file
View File

@@ -0,0 +1,3 @@
inout
crate
lmit

View File

@@ -2,7 +2,7 @@
name: ":bug: Bug Report"
about: Create a bug report to help us improve the repo
title: "[BUG]: "
labels: bug
labels: type:bug
---
## Description

View File

@@ -2,7 +2,7 @@
name: ":sparkles: Feature Request"
about: Request the inclusion of a new feature or functionality
title: "[FEAT]: "
labels: enhancement
labels: type:feature
---
## Description

13
.github/changed-files.yml vendored Normal file
View File

@@ -0,0 +1,13 @@
golang:
- goicicle/**/*.go'
- go.mod
rust:
- wrappers/rust
cpp:
- icicle/**/*.cu
- icicle/**/*.cuh
- icicle/**/*.cpp
- icicle/**/*.hpp
- icicle/**/*.c
- icicle/**/*.h
- icicle/CMakeLists.txt

View File

@@ -4,4 +4,4 @@ This PR...
## Linked Issues
Closes #
Resolves #

View File

@@ -1,46 +0,0 @@
name: Build
on:
pull_request:
branches: [ "main" ]
paths:
- "icicle/**"
- "src/**"
- "Cargo.toml"
- "build.rs"
env:
CARGO_TERM_COLOR: always
ARCH_TYPE: sm_70
jobs:
build-linux:
runs-on: ubuntu-latest
steps:
# Checkout code
- uses: actions/checkout@v3
# Download (or from cache) and install CUDA Toolkit 12.1.0
- uses: Jimver/cuda-toolkit@v0.2.9
id: cuda-toolkit
with:
cuda: '12.1.0'
use-github-cache: true
# Build from cargo - Rust utils are preinstalled on latest images
# https://github.com/actions/runner-images/blob/main/images/linux/Ubuntu2204-Readme.md#rust-tools
- name: Build
run: cargo build --release --verbose
build-windows:
runs-on: windows-latest
steps:
- uses: actions/checkout@v3
- uses: Jimver/cuda-toolkit@v0.2.9
id: cuda-toolkit
with:
cuda: '12.1.0'
use-github-cache: true
- name: Build
run: cargo build --release --verbose

20
.github/workflows/codespell.yml vendored Normal file
View File

@@ -0,0 +1,20 @@
name: Check Spelling
on:
pull_request:
branches:
- main
- dev
jobs:
spelling-checker:
name: Check Spelling
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@v3
- uses: codespell-project/actions-codespell@v2
with:
# https://github.com/codespell-project/actions-codespell?tab=readme-ov-file#parameter-skip
skip: ./**/target,./**/build
# https://github.com/codespell-project/actions-codespell?tab=readme-ov-file#parameter-ignore_words_file
ignore_words_file: .codespellignore

50
.github/workflows/examples.yml vendored Normal file
View File

@@ -0,0 +1,50 @@
# This workflow is a demo of how to run all examples in the Icicle repository.
# For each language directory (c++, Rust, etc.) the workflow
# (1) loops over all examples (msm, ntt, etc.) and
# (2) runs ./compile.sh and ./run.sh in each directory.
# The script ./compile.sh should compile the example and ./run.sh should run it.
# Each script should return 0 for success and 1 otherwise.
name: Examples
on:
pull_request:
branches:
- main
- dev
push:
branches:
- main
- dev
jobs:
test-examples:
runs-on: [self-hosted, Linux, X64, icicle] # ubuntu-latest
steps:
- name: Checkout
uses: actions/checkout@v2
- name: c++ examples
working-directory: ./examples/c++
run: |
# loop over all directories in the current directory
for dir in $(find . -mindepth 1 -maxdepth 1 -type d); do
if [ -d "$dir" ]; then
echo "Running command in $dir"
cd $dir
./compile.sh
./run.sh
cd -
fi
done
- name: Rust examples
working-directory: ./examples/rust
run: |
# loop over all directories in the current directory
for dir in $(find . -mindepth 1 -maxdepth 1 -type d); do
if [ -d "$dir" ]; then
echo "Running command in $dir"
cd $dir
cargo run --release
cd -
fi
done

115
.github/workflows/main-build.yml vendored Normal file
View File

@@ -0,0 +1,115 @@
name: Build
on:
pull_request:
branches:
- main
- dev
push:
branches:
- main
- dev
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
env:
CARGO_TERM_COLOR: always
ARCH_TYPE: native
jobs:
check-changed-files:
name: Check Changed Files
runs-on: ubuntu-22.04
outputs:
golang: ${{ steps.changed_files.outputs.golang }}
rust: ${{ steps.changed_files.outputs.rust }}
cpp_cuda: ${{ steps.changed_files.outputs.cpp_cuda }}
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Get all changed files
id: changed-files-yaml
uses: tj-actions/changed-files@v39
# https://github.com/tj-actions/changed-files#input_files_yaml_from_source_file
with:
files_yaml_from_source_file: .github/changed-files.yml
- name: Run Changed Files script
id: changed_files
# https://github.com/tj-actions/changed-files#outputs-
run: |
echo "golang=${{ steps.changed-files-yaml.outputs.golang_any_modified }}" >> "$GITHUB_OUTPUT"
echo "rust=${{ steps.changed-files-yaml.outputs.rust_any_modified }}" >> "$GITHUB_OUTPUT"
echo "cpp_cuda=${{ steps.changed-files-yaml.outputs.cpp_any_modified }}" >> "$GITHUB_OUTPUT"
build-rust-linux:
name: Build Rust on Linux
runs-on: [self-hosted, Linux, X64, icicle]
needs: check-changed-files
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Build Rust
working-directory: ./wrappers/rust
if: needs.check-changed-files.outputs.rust == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# Building from the root workspace will build all members of the workspace by default
run: cargo build --release --verbose
build-rust-windows:
name: Build Rust on Windows
runs-on: windows-2022
needs: check-changed-files
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Download and Install Cuda
if: needs.check-changed-files.outputs.rust == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
id: cuda-toolkit
uses: Jimver/cuda-toolkit@v0.2.11
with:
cuda: '12.0.0'
method: 'network'
# https://docs.nvidia.com/cuda/archive/12.0.0/cuda-installation-guide-microsoft-windows/index.html
sub-packages: '["cudart", "nvcc", "thrust", "visual_studio_integration"]'
- name: Build Rust Targets
working-directory: ./wrappers/rust
if: needs.check-changed-files.outputs.rust == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
env:
CUDA_PATH: ${{ steps.cuda-toolkit.outputs.CUDA_PATH }}
# Building from the root workspace will build all members of the workspace by default
run: cargo build --release --verbose
# TODO: Re-enable once Golang bindings for v1+ is finished
# build-golang-linux:
# name: Build Golang on Linux
# runs-on: [self-hosted, Linux, X64, icicle]
# needs: check-changed-files
# steps:
# - name: Checkout Repo
# uses: actions/checkout@v3
# - name: Build CUDA libs
# if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# run: make all
# working-directory: ./goicicle
# TODO: Add once Golang make file supports building for Windows
# build-golang-windows:
# name: Build Golang on Windows
# runs-on: windows-2022
# needs: check-changed-files
# steps:
# - name: Checkout Repo
# uses: actions/checkout@v3
# - name: Download and Install Cuda
# if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# uses: Jimver/cuda-toolkit@v0.2.11
# with:
# cuda: '12.0.0'
# method: 'network'
# # https://docs.nvidia.com/cuda/archive/12.0.0/cuda-installation-guide-microsoft-windows/index.html
# sub-packages: '["cudart", "nvcc", "thrust"]'
# - name: Build cpp libs
# if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# run: make all
# working-directory: ./goicicle

47
.github/workflows/main-format.yml vendored Normal file
View File

@@ -0,0 +1,47 @@
name: Format
on:
pull_request:
branches:
- main
- dev
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
jobs:
formatting-rust:
name: Check Rust Code Formatting
runs-on: ubuntu-22.04
steps:
- name: Checkout
uses: actions/checkout@v3
- name: Check rustfmt
working-directory: ./wrappers/rust
# "-name tagret -prune" removes searching in any directory named "target"
# Formatting by single file is necessary due to generated files not being present
# before building the project.
# e.g. icicle-cuda-runtime/src/bindings.rs is generated and icicle-cuda-runtime/src/lib.rs includes that module
# causing rustfmt to fail.
run: if [[ $(find . -name target -prune -o -iname *.rs -print | xargs cargo fmt --check --) ]]; then echo "Please run cargo fmt"; exit 1; fi
# - name: Check clippy
# run: cargo clippy --no-deps --all-features --all-targets
formatting-golang:
name: Check Golang Code Formatting
runs-on: ubuntu-22.04
steps:
- name: Checkout
uses: actions/checkout@v3
- name: Check gofmt
run: if [[ $(go list ./... | xargs go fmt) ]]; then echo "Please run go fmt"; exit 1; fi
formatting-cpp-cuda:
name: Check C++/CUDA Code Formatting
runs-on: ubuntu-22.04
steps:
- name: Checkout
uses: actions/checkout@v3
- name: Check clang-format
run: if [[ $(find ./ \( -path ./icicle/build -prune -o -path ./**/target -prune -o -path ./examples -prune \) -iname *.h -or -iname *.cuh -or -iname *.cu -or -iname *.c -or -iname *.cpp | xargs clang-format --dry-run -ferror-limit=1 -style=file 2>&1) ]]; then echo "Please run clang-format"; exit 1; fi

94
.github/workflows/main-test.yml vendored Normal file
View File

@@ -0,0 +1,94 @@
name: Test
on:
pull_request:
branches:
- main
- dev
concurrency:
group: ${{ github.workflow }}-${{ github.ref }}
cancel-in-progress: true
env:
CARGO_TERM_COLOR: always
ARCH_TYPE: native
jobs:
check-changed-files:
name: Check Changed Files
runs-on: ubuntu-22.04
outputs:
golang: ${{ steps.changed_files.outputs.golang }}
rust: ${{ steps.changed_files.outputs.rust }}
cpp_cuda: ${{ steps.changed_files.outputs.cpp_cuda }}
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Get all changed files
id: changed-files-yaml
uses: tj-actions/changed-files@v39
# https://github.com/tj-actions/changed-files#input_files_yaml_from_source_file
with:
files_yaml_from_source_file: .github/changed-files.yml
- name: Run Changed Files script
id: changed_files
# https://github.com/tj-actions/changed-files#outputs-
run: |
echo "golang=${{ steps.changed-files-yaml.outputs.golang_any_modified }}" >> "$GITHUB_OUTPUT"
echo "rust=${{ steps.changed-files-yaml.outputs.rust_any_modified }}" >> "$GITHUB_OUTPUT"
echo "cpp_cuda=${{ steps.changed-files-yaml.outputs.cpp_any_modified }}" >> "$GITHUB_OUTPUT"
test-rust-linux:
name: Test Rust on Linux
runs-on: [self-hosted, Linux, X64, icicle]
needs: check-changed-files
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Run Rust Tests
working-directory: ./wrappers/rust
if: needs.check-changed-files.outputs.rust == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# Running tests from the root workspace will run all workspace members' tests by default
# We need to limit the number of threads to avoid running out of memory on weaker machines
run: cargo test --release --verbose --features=g2 -- --test-threads=2
test-cpp-linux:
name: Test C++ on Linux
runs-on: [self-hosted, Linux, X64, icicle]
needs: check-changed-files
strategy:
matrix:
curve: [bn254, bls12_381, bls12_377, bw6_761]
steps:
- name: Checkout Repo
uses: actions/checkout@v3
- name: Build C++
working-directory: ./icicle
if: needs.check-changed-files.outputs.cpp_cuda == 'true'
run: |
mkdir -p build
cmake -DBUILD_TESTS=ON -DCMAKE_BUILD_TYPE=Release -DCURVE=${{ matrix.curve }} -S . -B build
cmake --build build
- name: Run C++ Tests
working-directory: ./icicle/build
if: needs.check-changed-files.outputs.cpp_cuda == 'true'
run: ctest
# TODO: Re-enable once Golang bindings for v1+ is finished
# test-golang-linux:
# name: Test Golang on Linux
# runs-on: [self-hosted, Linux, X64, icicle]
# needs: check-changed-files
# steps:
# - name: Checkout Repo
# uses: actions/checkout@v3
# - name: Build CUDA libs
# working-directory: ./goicicle
# if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# run: make libbn254.so
# - name: Run Golang Tests
# if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
# run: |
# export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$(pwd)/goicicle
# go test ./goicicle/curves/bn254 -count=1

5
.gitignore vendored
View File

@@ -5,6 +5,9 @@
*.cubin
*.bin
*.fatbin
*.so
*.nsys-rep
*.ncu-rep
**/target
**/.vscode
**/.*lock*csv#
@@ -12,3 +15,5 @@
**/.DS_Store
**/Cargo.lock
**/icicle/build/
**/wrappers/rust/icicle-cuda-runtime/src/bindings.rs
**/build

10
.rustfmt.toml Normal file
View File

@@ -0,0 +1,10 @@
# https://github.com/rust-lang/rustfmt/blob/master/Configurations.md
# Stable Configs
chain_width = 0
max_width = 120
merge_derives = true
use_field_init_shorthand = true
use_try_shorthand = true
# Unstable Configs

8
CITATION.cff Normal file
View File

@@ -0,0 +1,8 @@
cff-version: 1.2.0
message: "If you use this software, please cite it as below."
authors:
- family-names: "Ingonyama"
title: "ICICLE: GPU Library for ZK Acceleration"
version: 1.0.0
date-released: 2024-01-04
url: "https://github.com/ingonyama-zk/icicle"

View File

@@ -1,45 +0,0 @@
[package]
name = "icicle-utils"
version = "0.1.0"
edition = "2021"
authors = [ "Ingonyama" ]
description = "An implementation of the Ingonyama CUDA Library"
homepage = "https://www.ingonyama.com"
repository = "https://github.com/ingonyama-zk/icicle"
[[bench]]
name = "ntt"
path = "benches/ntt.rs"
harness = false
[[bench]]
name = "msm"
path = "benches/msm.rs"
harness = false
[dependencies]
hex = "*"
ark-std = "0.3.0"
ark-ff = "0.3.0"
ark-poly = "0.3.0"
ark-ec = { version = "0.3.0", features = [ "parallel" ] }
ark-bls12-381 = "0.3.0"
ark-bls12-377 = "0.3.0"
ark-bn254 = "0.3.0"
rustacuda = "0.1"
rustacuda_core = "0.1"
rustacuda_derive = "0.1"
rand = "*" #TODO: move rand and ark dependencies to dev once random scalar/point generation is done "natively"
[build-dependencies]
cc = { version = "1.0", features = ["parallel"] }
[dev-dependencies]
"criterion" = "0.4.0"
[features]
default = ["bls12_381"]
bls12_381 = ["ark-bls12-381/curve"]
g2 = []

28
Dockerfile Normal file
View File

@@ -0,0 +1,28 @@
# Use the specified base image
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
protobuf-compiler \
curl \
build-essential \
&& rm -rf /var/lib/apt/lists/*
# Install Rust
RUN curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
ENV PATH="/root/.cargo/bin:${PATH}"
# Install Golang
ENV GOLANG_VERSION 1.21.1
RUN curl -L https://golang.org/dl/go${GOLANG_VERSION}.linux-amd64.tar.gz | tar -xz -C /usr/local
ENV PATH="/usr/local/go/bin:${PATH}"
# Set the working directory in the container
WORKDIR /app
# Copy the content of the local directory to the working directory
COPY . .
# Specify the default command for the container
CMD ["/bin/bash"]

219
README.md
View File

@@ -1,150 +1,124 @@
# ICICLE
<div align="center">Icicle is a library for ZK acceleration using CUDA-enabled GPUs.</div>
![image (4)](https://user-images.githubusercontent.com/2446179/223707486-ed8eb5ab-0616-4601-8557-12050df8ccf7.png)
**<div align="center">ICICLE is a library for ZK acceleration using CUDA-enabled GPUs.</div>**
<p align="center">
<img alt="ICICLE" width="300" height="300" src="https://user-images.githubusercontent.com/2446179/223707486-ed8eb5ab-0616-4601-8557-12050df8ccf7.png"/>
</p>
<p align="center">
<a href="https://discord.gg/EVVXTdt6DF">
<img src="https://img.shields.io/discord/1063033227788423299?logo=discord" alt="Chat with us on Discord">
</a>
<a href="https://twitter.com/intent/follow?screen_name=Ingo_zk">
<img src="https://img.shields.io/twitter/follow/Ingo_zk?style=social&logo=twitter" alt="Follow us on Twitter">
</a>
<img src="https://img.shields.io/badge/Machines%20running%20ICICLE-544-lightblue" alt="Machines running ICICLE">
</p>
## Background
Zero Knowledge Proofs (ZKPs) are considered one of the greatest achievements of modern cryptography. Accordingly, ZKPs are expected to disrupt a number of industries and will usher in an era of trustless and privacy preserving services and infrastructure.
If we want ZK hardware today we have FPGAs or GPUs which are relatively inexpensive. However, the biggest selling point of GPUs is the software; we talk in particular about CUDA, which makes it easy to write code running on Nvidia GPUs, taking advantage of their highly parallel architecture. Together with the widespread availability of these devices, if we can get GPUs to work on ZK workloads, then we have made a giant step towards accessible and efficient ZK provers.
We believe GPUs are as important for ZK as for AI.
## Zero Knowledge on GPU
- GPUs are a perfect match for ZK compute - around 97% of ZK protocol runtime is parallel by nature.
- GPUs are simple for developers to use and scale compared to other hardware platforms.
- GPUs are extremely competitive in terms of power / performance and price (3x cheaper).
- GPUs are popular and readily available.
ICICLE is a CUDA implementation of general functions widely used in ZKP. ICICLE currently provides support for MSM, NTT, and ECNTT, with plans to support Hash functions soon.
## Getting Started
### Supported primitives
ICICLE is a CUDA implementation of general functions widely used in ZKP.
- Fields
- Scalars
- Points
- Projective: {x, y, z}
- Affine: {x, y}
- Curves
- [BLS12-381]
- [BLS12-377]
- [BN254]
> [!NOTE]
> Developers: We highly recommend reading our [documentation]
## Build and usage
> [!TIP]
> Try out ICICLE by running some [examples] using ICICLE in C++ and our Rust bindings
> NOTE: [NVCC] is a prerequisite for building.
### Prerequisites
1. Define or select a curve for your application; we've provided a [template][CRV_TEMPLATE] for defining a curve
2. Include the curve in [`curve_config.cuh`][CRV_CONFIG]
3. Now you can build the ICICLE library using nvcc
- [CUDA Toolkit](https://developer.nvidia.com/cuda-downloads) version 12.0 or newer.
- [CMake]((https://cmake.org/files/)), version 3.18 and above. Latest version is recommended.
- [GCC](https://gcc.gnu.org/install/download.html) version 9, latest version is recommended.
- Any Nvidia GPU (which supports CUDA Toolkit version 12.0 or above).
```sh
mkdir -p build
nvcc -o build/<ENTER_DIR_NAME> ./icicle/appUtils/ntt/ntt.cu ./icicle/appUtils/msm/msm.cu ./icicle/appUtils/vector_manipulation/ve_mod_mult.cu ./icicle/primitives/projective.cu -lib -arch=native
```
> [!NOTE]
> It is possible to use CUDA 11 for cards which dont support CUDA 12, however we dont officially support this version and in the future there may be issues.
### Testing the CUDA code
### Accessing Hardware
We are using [googletest] library for testing. To build and run [the test suite](./icicle/README.md) for finite field and elliptic curve arithmetic, run from the `icicle` folder:
If you don't have access to a Nvidia GPU we have some options for you.
```sh
mkdir -p build
cmake -S . -B build
cmake --build build
cd build && ctest
```
Checkout [Google Colab](https://colab.google/). Google Colab offers a free [T4 GPU](https://www.nvidia.com/en-us/data-center/tesla-t4/) instance and ICICLE can be used with it, reference this guide for setting up your [Google Colab workplace][GOOGLE-COLAB-ICICLE].
### Rust Bindings
If you require more compute and have an interesting research project, we have [bounty and grant programs][GRANT_PROGRAM].
For convenience, we also provide rust bindings to the ICICLE library for the following primitives:
- MSM
- NTT
- Forward NTT
- Inverse NTT
- ECNTT
- Forward ECNTT
- Inverse NTT
- Scalar Vector Multiplication
- Point Vector Multiplication
### Build systems
A custom [build script][B_SCRIPT] is used to compile and link the ICICLE library. The environement variable `ARCH_TYPE` is used to determine which GPU type the library should be compiled for and it defaults to `native` when it is not set allowing the compiler to detect the installed GPU type.
ICICLE has three build systems.
> NOTE: A GPU must be detectable and therefore installed if the `ARCH_TYPE` is not set.
- [ICICLE core][ICICLE-CORE], C++ and CUDA
- [ICICLE Rust][ICICLE-RUST] bindings, requires [Rust](https://www.rust-lang.org/) version 1.70 and above
- [ICICLE Golang][ICICLE-GO] bindings, requires [Go](https://go.dev/) version 1.20 and above
Once you have your parameters set, run:
ICICLE core always needs to be built as part of the other build systems as it contains the core ICICLE primitives implemented in CUDA. Reference these guides for the different build systems, [ICICLE core guide][ICICLE-CORE-README], [ICICLE Rust guide][ICICLE-RUST-README] and [ICICLE Golang guide][ICICLE-GO-README].
```sh
cargo build --release
```
### Compiling ICICLE
You'll find a release ready library at `target/release/libicicle_utils.rlib`.
Running ICICLE via Rust bindings is highly recommended and simple:
- Clone this repo
- go to our [Rust bindings][ICICLE-RUST]
- Enter a [curve](./wrappers/rust/icicle-curves) implementation
- run `cargo build --release` to build or `cargo test -- --test-threads=1` to build and execute tests
To benchmark and test the functionality available in RUST, run:
In any case you would want to compile and run core icicle c++ tests, just follow these setps:
- Clone this repo
- go to [ICICLE core][ICICLE-CORE]
- execute the small [script](https://github.com/ingonyama-zk/icicle/tree/main/icicle#running-tests) to compile via cmake and run c++ and cuda tests
## Docker
We offer a simple Docker container so you can simply run ICICLE without setting everything up locally.
```
cargo bench
cargo test -- --test-threads=1
docker build -t <name_of_your_choice> .
docker run --gpus all -it <name_of_your_choice> /bin/bash
```
The flag `--test-threads=1` is needed because currently some tests might interfere with one another inside the GPU.
### Example Usage
An example of using the Rust bindings library can be found in our [fast-danksharding implementation][FDI]
### Supporting Additional Curves
Supporting additional curves can be done as follows:
Create a JSON file with the curve parameters. The curve is defined by the following parameters:
- ``curve_name`` - e.g. ``bls12_381``.
- ``modolus_p`` - scalar field modolus (in decimal).
- ``bit_count_p`` - number of bits needed to represent `` modolus_p`` .
- ``limb_p`` - number of bytes needed to represent `` modolus_p`` (rounded).
- ``ntt_size`` - log of the maximal size subgroup of the scalar field.
- ``modolus_q`` - base field modulus (in decimal).
- ``bit_count_q`` - number of bits needed to represent `` modolus_q`` .
- ``limb_q`` number of bytes needed to represent `` modolus_p`` (rounded).
- ``weierstrass_b`` - Weierstrauss constant of the curve.
- ``gen_x`` - x-value of a generator element for the curve.
- ``gen_y`` - y-value of a generator element for the curve.
Here's an example for BLS12-381.
```
{
"curve_name" : "bls12_381",
"modolus_p" : 52435875175126190479447740508185965837690552500527637822603658699938581184513,
"bit_count_p" : 255,
"limb_p" : 8,
"ntt_size" : 32,
"modolus_q" : 4002409555221667393417789825735904156556882819939007885332058136124031650490837864442687629129015664037894272559787,
"bit_count_q" : 381,
"limb_q" : 12,
"weierstrass_b" : 4,
"gen_x" : 3685416753713387016781088315183077757961620795782546409894578378688607592378376318836054947676345821548104185464507,
"gen_y" : 1339506544944476473020471379941921221584933875938349620426543736416511423956333506472724655353366534992391756441569
}
```
Save the parameters JSON file in ``curve_parameters``.
Then run the Python script ``new_curve_script.py `` from the main icicle folder:
```
python3 ./curve_parameters/new_curve_script_rust.py ./curve_parameters/bls12_381.json
```
The script does the following:
- Creates a folder in ``icicle/curves`` with the curve name, which contains all of the files needed for the supported operations in cuda.
- Adds the curve exported operations to ``icicle/curves/index.cu``.
- Creates a file with the curve name in ``src/curves`` with the relevant objects for the curve.
- Creates a test file with the curve name in ``src``.
Testing the new curve could be done by running the tests in ``tests_curve_name`` (e.g. ``tests_bls12_381``).
## Contributions
Join our [Discord Server](https://discord.gg/Y4SkbDf2Ff) and find us on the icicle channel. We will be happy to work together to support your use case and talk features, bugs and design.
Join our [Discord Server][DISCORD] and find us on the icicle channel. We will be happy to work together to support your use case and talk features, bugs and design.
### Development Contributions
If you are changing code, please make sure to change your [git hooks path][HOOKS_DOCS] to the repo's [hooks directory][HOOKS_PATH] by running the following command:
```sh
git config core.hooksPath ./scripts/hooks
```
In case `clang-format` is missing on your system, you can install it using the following command:
```sh
sudo apt install clang-format
```
You will also need to install [codespell](https://github.com/codespell-project/codespell?tab=readme-ov-file#installation) to check for typos.
This will ensure our custom hooks are run and will make it easier to follow our coding guidelines.
### Hall of Fame
- [Robik](https://github.com/robik75), for his on-going support and mentorship
- [Robik](https://github.com/robik75), for his ongoing support and mentorship
- [liuxiao](https://github.com/liuxiaobleach), for being a top notch bug smasher
- [gkigiermo](https://github.com/gkigiermo), for making it intuitive to use ICICLE in Google Colab.
## Help & Support
For help and support talk to our devs in our discord channel ["ICICLE"](https://discord.gg/EVVXTdt6DF)
## License
@@ -153,13 +127,26 @@ ICICLE is distributed under the terms of the MIT License.
See [LICENSE-MIT][LMIT] for details.
<!-- Begin Links -->
[BLS12-381]: ./icicle/curves/bls12_381.cuh
[BLS12-381]: ./icicle/curves/
[BLS12-377]: ./icicle/curves/
[BN254]: ./icicle/curves/
[BW6-671]: ./icicle/curves/
[NVCC]: https://docs.nvidia.com/cuda/#installation-guides
[CRV_TEMPLATE]: ./icicle/curves/curve_template.cuh
[CRV_CONFIG]: ./icicle/curves/curve_config.cuh
[B_SCRIPT]: ./build.rs
[FDI]: https://github.com/ingonyama-zk/fast-danksharding
[LMIT]: ./LICENSE
[DISCORD]: https://discord.gg/Y4SkbDf2Ff
[googletest]: https://github.com/google/googletest/
[HOOKS_DOCS]: https://git-scm.com/docs/githooks
[HOOKS_PATH]: ./scripts/hooks/
[CMAKELISTS]: https://github.com/ingonyama-zk/icicle/blob/f0e6b465611227b858ec4590f4de5432e892748d/icicle/CMakeLists.txt#L28
[GOOGLE-COLAB-ICICLE]: https://dev.ingonyama.com/icicle/colab-instructions
[GRANT_PROGRAM]: https://medium.com/@ingonyama/icicle-for-researchers-grants-challenges-9be1f040998e
[ICICLE-CORE]: ./icicle/
[ICICLE-RUST]: ./wrappers/rust/
[ICICLE-GO]: ./goicicle/
[ICICLE-CORE-README]: ./icicle/README.md
[ICICLE-RUST-README]: ./wrappers/rust/README.md
[ICICLE-GO-README]: ./goicicle/README.md
[documentation]: https://dev.ingonyama.com/icicle/overview
[examples]: ./examples/
<!-- End Links -->

View File

@@ -1,50 +0,0 @@
extern crate criterion;
use criterion::{criterion_group, criterion_main, Criterion};
use icicle_utils::{set_up_scalars, generate_random_points, commit_batch, get_rng, field::BaseField};
#[cfg(feature = "g2")]
use icicle_utils::{commit_batch_g2, field::ExtensionField};
use rustacuda::prelude::*;
const LOG_MSM_SIZES: [usize; 1] = [12];
const BATCH_SIZES: [usize; 2] = [128, 256];
fn bench_msm(c: &mut Criterion) {
let mut group = c.benchmark_group("MSM");
for log_msm_size in LOG_MSM_SIZES {
for batch_size in BATCH_SIZES {
let msm_size = 1 << log_msm_size;
let (scalars, _, _) = set_up_scalars(msm_size, 0, false);
let batch_scalars = vec![scalars; batch_size].concat();
let mut d_scalars = DeviceBuffer::from_slice(&batch_scalars[..]).unwrap();
let points = generate_random_points::<BaseField>(msm_size, get_rng(None));
let batch_points = vec![points; batch_size].concat();
let mut d_points = DeviceBuffer::from_slice(&batch_points[..]).unwrap();
#[cfg(feature = "g2")]
let g2_points = generate_random_points::<ExtensionField>(msm_size, get_rng(None));
#[cfg(feature = "g2")]
let g2_batch_points = vec![g2_points; batch_size].concat();
#[cfg(feature = "g2")]
let mut d_g2_points = DeviceBuffer::from_slice(&g2_batch_points[..]).unwrap();
group.sample_size(30).bench_function(
&format!("MSM of size 2^{} in batch {}", log_msm_size, batch_size),
|b| b.iter(|| commit_batch(&mut d_points, &mut d_scalars, batch_size))
);
#[cfg(feature = "g2")]
group.sample_size(10).bench_function(
&format!("G2 MSM of size 2^{} in batch {}", log_msm_size, batch_size),
|b| b.iter(|| commit_batch_g2(&mut d_g2_points, &mut d_scalars, batch_size))
);
}
}
}
criterion_group!(msm_benches, bench_msm);
criterion_main!(msm_benches);

View File

@@ -1,33 +0,0 @@
extern crate criterion;
use criterion::{criterion_group, criterion_main, Criterion};
use icicle_utils::{interpolate_scalars_batch, interpolate_points_batch, set_up_scalars, set_up_points};
const LOG_NTT_SIZES: [usize; 1] = [15];
const BATCH_SIZES: [usize; 2] = [8, 16];
fn bench_ntt(c: &mut Criterion) {
let mut group = c.benchmark_group("NTT");
for log_ntt_size in LOG_NTT_SIZES {
for batch_size in BATCH_SIZES {
let ntt_size = 1 << log_ntt_size;
let (_, mut d_evals, mut d_domain) = set_up_scalars(ntt_size * batch_size, log_ntt_size, true);
let (_, mut d_points_evals, _) = set_up_points(ntt_size * batch_size, log_ntt_size, true);
group.sample_size(100).bench_function(
&format!("Scalar NTT of size 2^{} in batch {}", log_ntt_size, batch_size),
|b| b.iter(|| interpolate_scalars_batch(&mut d_evals, &mut d_domain, batch_size))
);
group.sample_size(10).bench_function(
&format!("EC NTT of size 2^{} in batch {}", log_ntt_size, batch_size),
|b| b.iter(|| interpolate_points_batch(&mut d_points_evals, &mut d_domain, batch_size))
);
}
}
}
criterion_group!(ntt_benches, bench_ntt);
criterion_main!(ntt_benches);

View File

@@ -1,29 +0,0 @@
use std::env;
fn main() {
//TODO: check cargo features selected
//TODO: can conflict/duplicate with make ?
println!("cargo:rerun-if-env-changed=CXXFLAGS");
println!("cargo:rerun-if-changed=./icicle");
let arch_type = env::var("ARCH_TYPE").unwrap_or(String::from("native"));
let mut arch = String::from("-arch=");
arch.push_str(&arch_type);
let mut nvcc = cc::Build::new();
println!("Compiling icicle library using arch: {}", &arch);
if cfg!(feature = "g2") {
nvcc.define("G2_DEFINED", None);
}
nvcc.cuda(true);
nvcc.debug(false);
nvcc.flag(&arch);
nvcc.files([
"./icicle/curves/index.cu",
]);
nvcc.compile("ingo_icicle"); //TODO: extension??
}

View File

@@ -1,13 +0,0 @@
{
"curve_name" : "bls12_377",
"modolus_p" : 8444461749428370424248824938781546531375899335154063827935233455917409239041,
"bit_count_p" : 253,
"limb_p" : 8,
"ntt_size" : 32,
"modolus_q" : 258664426012969094010652733694893533536393512754914660539884262666720468348340822774968888139573360124440321458177,
"bit_count_q" : 377,
"limb_q" : 12,
"weierstrass_b" : 1,
"gen_x" : 81937999373150964239938255573465948239988671502647976594219695644855304257327692006745978603320413799295628339695,
"gen_y" : 241266749859715473739788878240585681733927191168601896383759122102112907357779751001206799952863815012735208165030
}

View File

@@ -1,13 +0,0 @@
{
"curve_name" : "bls12_381",
"modolus_p" : 52435875175126190479447740508185965837690552500527637822603658699938581184513,
"bit_count_p" : 255,
"limb_p" : 8,
"ntt_size" : 32,
"modolus_q" : 4002409555221667393417789825735904156556882819939007885332058136124031650490837864442687629129015664037894272559787,
"bit_count_q" : 381,
"limb_q" : 12,
"weierstrass_b" : 4,
"gen_x" : 3685416753713387016781088315183077757961620795782546409894578378688607592378376318836054947676345821548104185464507,
"gen_y" : 1339506544944476473020471379941921221584933875938349620426543736416511423956333506472724655353366534992391756441569
}

View File

@@ -1,13 +0,0 @@
{
"curve_name" : "bn254",
"modolus_p" : 21888242871839275222246405745257275088548364400416034343698204186575808495617,
"bit_count_p" : 254,
"limb_p" : 8,
"ntt_size" : 16,
"modolus_q" : 21888242871839275222246405745257275088696311157297823662689037894645226208583,
"bit_count_q" : 254,
"limb_q" : 8,
"weierstrass_b" : 3,
"gen_x" : 1,
"gen_y" : 2
}

View File

@@ -1,203 +0,0 @@
import json
import math
import os
from sympy.ntheory import isprime, primitive_root
import subprocess
import random
import sys
data = None
with open(sys.argv[1]) as json_file:
data = json.load(json_file)
curve_name = data["curve_name"]
modolus_p = data["modolus_p"]
bit_count_p = data["bit_count_p"]
limb_p = data["limb_p"]
ntt_size = data["ntt_size"]
modolus_q = data["modolus_q"]
bit_count_q = data["bit_count_q"]
limb_q = data["limb_q"]
weierstrass_b = data["weierstrass_b"]
gen_x = data["gen_x"]
gen_y = data["gen_y"]
def to_hex(val, length):
x = str(hex(val))[2:]
if len(x) % 8 != 0:
x = "0" * (8-len(x) % 8) + x
if len(x) != length:
x = "0" * (length-len(x)) + x
n = 8
chunks = [x[i:i+n] for i in range(0, len(x), n)][::-1]
s = ""
for c in chunks:
s += "0x" + c + ", "
return s
def get_root_of_unity(order: int) -> int:
assert (modolus_p - 1) % order == 0
return pow(5, (modolus_p - 1) // order, modolus_p)
def create_field_parameters_struct(modulus, modulus_bits_count,limbs,ntt,size,name):
s = " struct "+name+"{\n"
s += " static constexpr unsigned limbs_count = " + str(limbs)+";\n"
s += " static constexpr storage<limbs_count> modulus = {"+to_hex(modulus,8*limbs)[:-2]+"};\n"
s += " static constexpr storage<limbs_count> modulus_2 = {"+to_hex(modulus*2,8*limbs)[:-2]+"};\n"
s += " static constexpr storage<limbs_count> modulus_4 = {"+to_hex(modulus*4,8*limbs)[:-2]+"};\n"
s += " static constexpr storage<2*limbs_count> modulus_wide = {"+to_hex(modulus,8*limbs*2)[:-2]+"};\n"
s += " static constexpr storage<2*limbs_count> modulus_sqared = {"+to_hex(modulus*modulus,8*limbs)[:-2]+"};\n"
s += " static constexpr storage<2*limbs_count> modulus_sqared_2 = {"+to_hex(modulus*modulus*2,8*limbs)[:-2]+"};\n"
s += " static constexpr storage<2*limbs_count> modulus_sqared_4 = {"+to_hex(modulus*modulus*2*2,8*limbs)[:-2]+"};\n"
s += " static constexpr unsigned modulus_bits_count = "+str(modulus_bits_count)+";\n"
m = int(math.floor(int(pow(2,2*modulus_bits_count) // modulus)))
s += " static constexpr storage<limbs_count> m = {"+ to_hex(m,8*limbs)[:-2] +"};\n"
s += " static constexpr storage<limbs_count> one = {"+ to_hex(1,8*limbs)[:-2] +"};\n"
s += " static constexpr storage<limbs_count> zero = {"+ to_hex(0,8*limbs)[:-2] +"};\n"
if ntt:
for k in range(size):
omega = get_root_of_unity(int(pow(2,k+1)))
s += " static constexpr storage<limbs_count> omega"+str(k+1)+"= {"+ to_hex(omega,8*limbs)[:-2]+"};\n"
for k in range(size):
omega = get_root_of_unity(int(pow(2,k+1)))
s += " static constexpr storage<limbs_count> omega_inv"+str(k+1)+"= {"+ to_hex(pow(omega, -1, modulus),8*limbs)[:-2]+"};\n"
for k in range(size):
s += " static constexpr storage<limbs_count> inv"+str(k+1)+"= {"+ to_hex(pow(int(pow(2,k+1)), -1, modulus),8*limbs)[:-2]+"};\n"
s+=" };\n"
return s
def create_gen():
s = " struct group_generator {\n"
s += " static constexpr storage<fq_config::limbs_count> generator_x = {"+to_hex(gen_x,8*limb_q)[:-2]+ "};\n"
s += " static constexpr storage<fq_config::limbs_count> generator_y = {"+to_hex(gen_y,8*limb_q)[:-2]+ "};\n"
s+=" };\n"
return s
def get_config_file_content(modolus_p, bit_count_p, limb_p, ntt_size, modolus_q, bit_count_q, limb_q, weierstrass_b):
file_content = ""
file_content += "#pragma once\n#include \"../../utils/storage.cuh\"\n"
file_content += "namespace PARAMS_"+curve_name.upper()+"{\n"
file_content += create_field_parameters_struct(modolus_p,bit_count_p,limb_p,True,ntt_size,"fp_config")
file_content += create_field_parameters_struct(modolus_q,bit_count_q,limb_q,False,0,"fq_config")
file_content += " static constexpr unsigned weierstrass_b = " + str(weierstrass_b)+ ";\n"
file_content += create_gen()
file_content+="}\n"
return file_content
# Create Cuda interface
newpath = "./icicle/curves/"+curve_name
if not os.path.exists(newpath):
os.makedirs(newpath)
fc = get_config_file_content(modolus_p, bit_count_p, limb_p, ntt_size, modolus_q, bit_count_q, limb_q, weierstrass_b)
text_file = open("./icicle/curves/"+curve_name+"/params.cuh", "w")
n = text_file.write(fc)
text_file.close()
with open("./icicle/curves/curve_template/lde.cu", "r") as lde_file:
content = lde_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
text_file = open("./icicle/curves/"+curve_name+"/lde.cu", "w")
n = text_file.write(content)
text_file.close()
with open("./icicle/curves/curve_template/msm.cu", "r") as msm_file:
content = msm_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
text_file = open("./icicle/curves/"+curve_name+"/msm.cu", "w")
n = text_file.write(content)
text_file.close()
with open("./icicle/curves/curve_template/ve_mod_mult.cu", "r") as ve_mod_mult_file:
content = ve_mod_mult_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
text_file = open("./icicle/curves/"+curve_name+"/ve_mod_mult.cu", "w")
n = text_file.write(content)
text_file.close()
namespace = '#include "params.cuh"\n'+'''namespace CURVE_NAME_U {
typedef Field<PARAMS_CURVE_NAME_U::fp_config> scalar_field_t;\
typedef scalar_field_t scalar_t;\
typedef Field<PARAMS_CURVE_NAME_U::fq_config> point_field_t;
typedef Projective<point_field_t, scalar_field_t, PARAMS_CURVE_NAME_U::group_generator, PARAMS_CURVE_NAME_U::weierstrass_b> projective_t;
typedef Affine<point_field_t> affine_t;
}'''
with open('./icicle/curves/'+curve_name+'/curve_config.cuh', 'w') as f:
f.write(namespace.replace("CURVE_NAME_U",curve_name.upper()))
eq = '''
#include <cuda.h>\n
#include "curve_config.cuh"\n
#include "../../primitives/projective.cuh"\n
extern "C" bool eq_CURVE_NAME_L(CURVE_NAME_U::projective_t *point1, CURVE_NAME_U::projective_t *point2)
{
return (*point1 == *point2);
}'''
with open('./icicle/curves/'+curve_name+'/projective.cu', 'w') as f:
f.write(eq.replace("CURVE_NAME_U",curve_name.upper()).replace("CURVE_NAME_L",curve_name.lower()))
supported_operations = '''
#include "projective.cu"
#include "lde.cu"
#include "msm.cu"
#include "ve_mod_mult.cu"
'''
with open('./icicle/curves/'+curve_name+'/supported_operations.cu', 'w') as f:
f.write(supported_operations.replace("CURVE_NAME_U",curve_name.upper()).replace("CURVE_NAME_L",curve_name.lower()))
with open('./icicle/curves/index.cu', 'a') as f:
f.write('\n#include "'+curve_name.lower()+'/supported_operations.cu"')
# Create Rust interface and tests
if limb_p == limb_q:
with open("./src/curve_templates/curve_same_limbs.rs", "r") as curve_file:
content = curve_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
content = content.replace("_limbs_p",str(limb_p * 8 * 4))
content = content.replace("limbs_p",str(limb_p))
text_file = open("./src/curves/"+curve_name+".rs", "w")
n = text_file.write(content)
text_file.close()
else:
with open("./src/curve_templates/curve_different_limbs.rs", "r") as curve_file:
content = curve_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
content = content.replace("_limbs_p",str(limb_p * 8 * 4))
content = content.replace("limbs_p",str(limb_p))
content = content.replace("_limbs_q",str(limb_q * 8 * 4))
content = content.replace("limbs_q",str(limb_q))
text_file = open("./src/curves/"+curve_name+".rs", "w")
n = text_file.write(content)
text_file.close()
with open("./src/curve_templates/test.rs", "r") as test_file:
content = test_file.read()
content = content.replace("CURVE_NAME_U",curve_name.upper())
content = content.replace("CURVE_NAME_L",curve_name.lower())
text_file = open("./src/test_"+curve_name+".rs", "w")
n = text_file.write(content)
text_file.close()
with open('./src/curves/mod.rs', 'a') as f:
f.write('\n pub mod ' + curve_name + ';')
with open('./src/lib.rs', 'a') as f:
f.write('\npub mod ' + curve_name + ';')

23
examples/ZKContainer.md Normal file
View File

@@ -0,0 +1,23 @@
# ZKContainer
We recommend using [ZKContainer](https://ingonyama.com/blog/Immanuel-ZKDC), where we have already preinstalled all the required dependencies, to run Icicle examples.
To use our containers you will need [Docker](https://www.docker.com/) and [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/index.html).
In each example directory, ZKContainer files are located in a subdirectory `.devcontainer`.
- File `Dockerfile` specifies how to build an image of a ZKContainer.
- File `devcontainer.json` enables running ZKContainer from Visual Studio Code.
## Running ZKContainer from shell
```sh
docker build -t icicle-example-poseidon -f .devcontainer/Dockerfile .
```
To run the example interactively, start the container
```sh
docker run -it --rm --gpus all -v .:/icicle-example icicle-example-poseidon
```
Inside the container, run the commands for building the library for whichever [build system](../README.md#build-systems) you choose to use.

View File

@@ -0,0 +1,25 @@
cmake_minimum_required(VERSION 3.18)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
if (${CMAKE_VERSION} VERSION_LESS "3.24.0")
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH})
else()
set(CMAKE_CUDA_ARCHITECTURES native) # on 3.24+, on earlier it is ignored, and the target is not passed
endif ()
project(icicle LANGUAGES CUDA CXX)
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr")
set(CMAKE_CUDA_FLAGS_RELEASE "")
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -g -G -O0")
# change the path to your Icicle location
include_directories("../../../icicle")
add_executable(
example
example.cu
)
find_library(NVML_LIBRARY nvidia-ml PATHS /usr/local/cuda-12.0/targets/x86_64-linux/lib/stubs/ )
target_link_libraries(example ${NVML_LIBRARY})
set_target_properties(example PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

View File

@@ -0,0 +1,72 @@
# Icicle example: build a Merkle tree using Poseidon hash
## Best-Practices
We recommend to run our examples in [ZK-containers](../../ZK-containers.md) to save your time and mental energy.
## Key-Takeaway
`Icicle` provides CUDA C++ template `poseidon_hash` to accelerate the popular [Poseidon hash function](https://www.poseidon-hash.info/).
## Concise Usage Explanation
```c++
#include "appUtils/poseidon/poseidon.cu"
...
poseidon_hash<scalar_t, arity+1>(input, output, n, constants, config);
```
**Parameters:**
- **`scalar_t`:** a scalar field of the selected curve.
You can think of field's elements as 32-byte integers modulo `p`, where `p` is a prime number, specific to this field.
- **arity:** number of elements in a hashed block.
- **n:** number of blocks we hash in parallel.
- **input, output:** `scalar_t` arrays of size $arity*n$ and $n$ respectively.
- **constants:** are defined as below
```c++
device_context::DeviceContext ctx= device_context::get_default_device_context();
PoseidonConstants<scalar_t> constants;
init_optimized_poseidon_constants<scalar_t>(ctx, &constants);
```
## What's in the example
1. Define the size of the example: the height of the full binary Merkle tree.
2. Hash blocks in parallel. The tree width determines the number of blocks to hash.
3. Build a Merkle tree from the hashes.
4. Use the tree to generate a membership proof for one of computed hashes.
5. Validate the hash membership.
6. Tamper the hash.
7. Invalidate the membership of the tempered hash.
## Details
### Merkle tree structure
Our Merkle tree is a **full binary tree** stored in a 1D array.
The tree nodes are stored following a level-first traversal of the binary tree.
For a given level, we use offset to number elements from left to right. The node numbers on the figure below correspond to their locations in the array.
```text
Tree Level
0 0
/ \
1 2 1
/ \ / \
3 4 5 6 2
1D array representation: {0, 1, 2, 3, 4, 5, 6}
```
### Membership proof structure
We use two arrays:
- position (left/right) of the node along the path toward the root
- hash of a second node with the same parent

View File

@@ -0,0 +1,9 @@
#!/bin/bash
# Exit immediately on error
set -e
rm -rf build
mkdir -p build
cmake -S . -B build
cmake --build build

View File

@@ -0,0 +1,152 @@
#include <chrono>
#include <fstream>
#include <iostream>
// select the curve
#define CURVE_ID 2
// include Poseidon template
#include "appUtils/poseidon/poseidon.cu"
using namespace poseidon;
using namespace curve_config;
device_context::DeviceContext ctx= device_context::get_default_device_context();
// location of a tree node in the array for a given level and offset
inline uint32_t tree_index(uint32_t level, uint32_t offset) { return (1 << level) - 1 + offset; }
// We assume the tree has leaves already set, compute all other levels
void build_tree(
const uint32_t tree_height, scalar_t* tree, PoseidonConstants<scalar_t> * constants, PoseidonConfig config)
{
for (uint32_t level = tree_height - 1; level > 0; level--) {
const uint32_t next_level = level - 1;
const uint32_t next_level_width = 1 << next_level;
poseidon_hash<scalar_t, 2+1>(
&tree[tree_index(level, 0)], &tree[tree_index(next_level, 0)], next_level_width, *constants, config);
}
}
// linear search leaves for a given hash, return offset
uint32_t query_membership(scalar_t query, scalar_t* tree, const uint32_t tree_height)
{
const uint32_t tree_width = (1 << (tree_height - 1));
for (uint32_t i = 0; i < tree_width; i++) {
const scalar_t leaf = tree[tree_index(tree_height - 1, i)];
if (leaf == query) {
return i; // found the hash
}
}
return tree_height; // hash not found
}
void generate_proof(
uint32_t position,
scalar_t* tree,
const uint32_t tree_height,
uint32_t* proof_lr,
scalar_t* proof_hash)
{
uint32_t level_index = position;
for (uint32_t level = tree_height - 1; level > 0; level--) {
uint32_t lr;
uint32_t neighbour_index;
lr = level_index % 2;
if (lr == 0) {
// left
neighbour_index = level_index + 1;
} else {
// right
neighbour_index = level_index - 1;
}
proof_lr[level] = lr;
proof_hash[level] = tree[tree_index(level, neighbour_index)];
level_index /= 2;
}
// the proof must match this:
proof_hash[0] = tree[tree_index(0, 0)];
}
uint32_t validate_proof(
const scalar_t hash,
const uint32_t tree_height,
const uint32_t* proof_lr,
const scalar_t* proof_hash,
PoseidonConstants<scalar_t> * constants,
PoseidonConfig config)
{
scalar_t hashes_in[2], hash_out[1], level_hash;
level_hash = hash;
for (uint32_t level = tree_height - 1; level > 0; level--) {
if (proof_lr[level] == 0) {
hashes_in[0] = level_hash;
hashes_in[1] = proof_hash[level];
} else {
hashes_in[0] = proof_hash[level];
hashes_in[1] = level_hash;
}
// next level hash
poseidon_hash<scalar_t, 2+1>(hashes_in, hash_out, 1, *constants, config);
level_hash = hash_out[0];
}
return proof_hash[0] == level_hash;
}
int main(int argc, char* argv[])
{
std::cout << "1. Defining the size of the example: height of the full binary Merkle tree" << std::endl;
const uint32_t tree_height = 21;
std::cout << "Tree height: " << tree_height << std::endl;
const uint32_t tree_arity = 2;
const uint32_t leaf_level = tree_height - 1;
const uint32_t tree_width = 1 << leaf_level;
std::cout << "Tree width: " << tree_width << std::endl;
const uint32_t tree_size = (1 << tree_height) - 1;
std::cout << "Tree size: " << tree_size << std::endl;
scalar_t* tree = static_cast<scalar_t*>(malloc(tree_size * sizeof(scalar_t)));
std::cout << "2. Hashing blocks in parallel" << std::endl;
const uint32_t data_arity = 4;
std::cout << "Block size (arity): " << data_arity << std::endl;
std::cout << "Initializing blocks..." << std::endl;
scalar_t d = scalar_t::zero();
scalar_t* data = static_cast<scalar_t*>(malloc(tree_width * data_arity * sizeof(scalar_t)));
for (uint32_t i = 0; i < tree_width * data_arity; i++) {
data[i] = d;
d = d + scalar_t::one();
}
std::cout << "Hashing blocks into tree leaves..." << std::endl;
PoseidonConstants<scalar_t> constants;
init_optimized_poseidon_constants<scalar_t>(data_arity, ctx, &constants);
PoseidonConfig config = default_poseidon_config<scalar_t>(data_arity+1);
poseidon_hash<curve_config::scalar_t, data_arity+1>(data, &tree[tree_index(leaf_level, 0)], tree_width, constants, config);
std::cout << "3. Building Merkle tree" << std::endl;
PoseidonConstants<scalar_t> tree_constants;
init_optimized_poseidon_constants<scalar_t>(tree_arity, ctx, &tree_constants);
PoseidonConfig tree_config = default_poseidon_config<scalar_t>(tree_arity+1);
build_tree(tree_height, tree, &tree_constants, tree_config);
std::cout << "4. Generate membership proof" << std::endl;
uint32_t position = tree_width - 1;
std::cout << "Using the hash for block: " << position << std::endl;
scalar_t query = tree[tree_index(leaf_level, position)];
uint32_t query_position = query_membership(query, tree, tree_height);
// allocate arrays for the proof
uint32_t* proof_lr = static_cast<uint32_t*>(malloc(tree_height * sizeof(uint32_t)));
scalar_t* proof_hash = static_cast<scalar_t*>(malloc(tree_height * sizeof(scalar_t)));
generate_proof(query_position, tree, tree_height, proof_lr, proof_hash);
std::cout << "5. Validate the hash membership" << std::endl;
uint32_t validated;
const scalar_t hash = tree[tree_index(leaf_level, query_position)];
validated = validate_proof(hash, tree_height, proof_lr, proof_hash, &tree_constants, tree_config);
std::cout << "Validated: " << validated << std::endl;
std::cout << "6. Tamper the hash" << std::endl;
const scalar_t tampered_hash = hash + scalar_t::one();
validated = validate_proof(tampered_hash, tree_height, proof_lr, proof_hash, &tree_constants, tree_config);
std::cout << "7. Invalidate tamper hash membership" << std::endl;
std::cout << "Validated: " << validated << std::endl;
return 0;
}

View File

@@ -0,0 +1,2 @@
#!/bin/bash
./build/example

View File

@@ -0,0 +1,25 @@
# Make sure NVIDIA Container Toolkit is installed on your host
# Use the specified base image
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
curl \
build-essential \
git \
libboost-all-dev \
&& rm -rf /var/lib/apt/lists/*
# Clone Icicle from a GitHub repository
RUN git clone https://github.com/ingonyama-zk/icicle.git /opt/icicle
# Set the working directory in the container
WORKDIR /icicle-example
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,21 @@
{
"name": "Icicle Examples: msm",
"build": {
"dockerfile": "Dockerfile"
},
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-python.python"
]
}
}
}

View File

@@ -0,0 +1,25 @@
cmake_minimum_required(VERSION 3.18)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
if (${CMAKE_VERSION} VERSION_LESS "3.24.0")
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH})
else()
set(CMAKE_CUDA_ARCHITECTURES native) # on 3.24+, on earlier it is ignored, and the target is not passed
endif ()
project(icicle LANGUAGES CUDA CXX)
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr")
set(CMAKE_CUDA_FLAGS_RELEASE "")
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -g -G -O0")
# change the path to your Icicle location
include_directories("../../../icicle")
add_executable(
example
example.cu
)
find_library(NVML_LIBRARY nvidia-ml PATHS /usr/local/cuda-12.0/targets/x86_64-linux/lib/stubs/ )
target_link_libraries(example ${NVML_LIBRARY})
set_target_properties(example PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

View File

@@ -0,0 +1,52 @@
# Icicle example: Muli-Scalar Multiplication (MSM)
## Best-Practices
We recommend to run our examples in [ZK-containers](../../ZK-containers.md) to save your time and mental energy.
## Key-Takeaway
`Icicle` provides CUDA C++ template function `MSM` to accelerate [Multi-Scalar Multiplication](https://github.com/ingonyama-zk/ingopedia/blob/master/src/msm.md).
## Concise Usage Explanation
1. Select the curve
2. Include an MSM template
3. Configure MSM
4. Call the template
```c++
#define CURVE_ID 1
#include "icicle/appUtils/msm/msm.cu"
...
msm::MSMConfig config = {...};
...
msm::MSM<scalar_t, affine_t, projective_t>(scalars, points, size, config, &result);
```
In this example we use `BN254` curve (`CURVE_ID=1`). The function computes $result = \sum_{i=0}^{size-1} scalars[i] \cdot points[i]$, where input `points[]` use affine coordinates, and `result` uses projective coordinates.
**Parameters:**
The configuration is passed to the kernel as a structure of type `msm::MSMConfig`. Some of the most important fields are listed below:
- `are_scalars_on_device`, `are_points_on_device`, `are_results_on_device`: location of the data
- `is_async`: blocking vs. non-blocking kernel call
- `large_bucket_factor`: distinguishes between large bucket and normal bucket sizes. If there is a scalar distribution that is skewed heavily to a few values we can operate on those separately from the rest of the values. The ideal value here can vary by circuit (based on the distribution of scalars) but start with 10 and adjust to see if it improves performance.
## Running the example
- `cd` to your example directory
- compile with `./compile.sh`
- run with `./run.sh`
## What's in the example
1. Define the parameters of MSM
2. Generate random inputs on-host
3. Configure and execute MSM using on-host data
4. Copy inputs on-device
5. Configure and execute MSM using on-device data
6. Repeat the above steps for G2 points

9
examples/c++/msm/compile.sh Executable file
View File

@@ -0,0 +1,9 @@
#!/bin/bash
# Exit immediately on error
set -e
rm -rf build
mkdir -p build
cmake -S . -B build
cmake --build build

180
examples/c++/msm/example.cu Normal file
View File

@@ -0,0 +1,180 @@
#include <fstream>
#include <iostream>
#include <iomanip>
#define G2_DEFINED
#define CURVE_ID 1
// include MSM template
#include "appUtils/msm/msm.cu"
using namespace curve_config;
int main(int argc, char* argv[])
{
std::cout << "Icicle example: Muli-Scalar Multiplication (MSM)" << std::endl;
std::cout << "Example parameters" << std::endl;
int batch_size = 1;
std::cout << "Batch size: " << batch_size << std::endl;
unsigned msm_size = 1048576;
std::cout << "MSM size: " << msm_size << std::endl;
int N = batch_size * msm_size;
std::cout << "Part I: use G1 points" << std::endl;
std::cout << "Generating random inputs on-host" << std::endl;
scalar_t* scalars = new scalar_t[N];
affine_t* points = new affine_t[N];
projective_t result;
scalar_t::RandHostMany(scalars, N);
projective_t::RandHostManyAffine(points, N);
std::cout << "Using default MSM configuration with on-host inputs" << std::endl;
// auto config = msm::DefaultMSMConfig();
device_context::DeviceContext ctx = device_context::get_default_device_context();
msm::MSMConfig config = {
ctx, // ctx
0, // points_size
1, // precompute_factor
0, // c
0, // bitsize
10, // large_bucket_factor
1, // batch_size
false, // are_scalars_on_device
false, // are_scalars_montgomery_form
false, // are_points_on_device
false, // are_points_montgomery_form
false, // are_results_on_device
false, // is_big_triangle
false, // is_async
};
config.batch_size = batch_size;
std::cout << "Running MSM kernel with on-host inputs" << std::endl;
// Create two events to time the MSM kernel
cudaStream_t stream = config.ctx.stream;
cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// Record the start event on the stream
cudaEventRecord(start, stream);
// Execute the MSM kernel
msm::MSM<scalar_t, affine_t, projective_t>(scalars, points, msm_size, config, &result);
// Record the stop event on the stream
cudaEventRecord(stop, stream);
// Wait for the stop event to complete
cudaEventSynchronize(stop);
// Calculate the elapsed time between the start and stop events
cudaEventElapsedTime(&time, start, stop);
// Destroy the events
cudaEventDestroy(start);
cudaEventDestroy(stop);
// Print the elapsed time
std::cout << "Kernel runtime: " << std::fixed << std::setprecision(3) << time * 1e-3 << " sec." << std::endl;
// Print the result
std::cout << projective_t::to_affine(result) << std::endl;
std::cout << "Copying inputs on-device" << std::endl;
scalar_t* scalars_d;
affine_t* points_d;
projective_t* result_d;
cudaMalloc(&scalars_d, sizeof(scalar_t) * N);
cudaMalloc(&points_d, sizeof(affine_t) * N);
cudaMalloc(&result_d, sizeof(projective_t));
cudaMemcpy(scalars_d, scalars, sizeof(scalar_t) * N, cudaMemcpyHostToDevice);
cudaMemcpy(points_d, points, sizeof(affine_t) * N, cudaMemcpyHostToDevice);
std::cout << "Reconfiguring MSM to use on-device inputs" << std::endl;
config.are_results_on_device = true;
config.are_scalars_on_device = true;
config.are_points_on_device = true;
std::cout << "Running MSM kernel with on-device inputs" << std::endl;
// Create two events to time the MSM kernel
cudaEventCreate(&start);
cudaEventCreate(&stop);
// Record the start event on the stream
cudaEventRecord(start, stream);
// Execute the MSM kernel
msm::MSM<scalar_t, affine_t, projective_t>(scalars_d, points_d, msm_size, config, result_d);
// Record the stop event on the stream
cudaEventRecord(stop, stream);
// Wait for the stop event to complete
cudaEventSynchronize(stop);
// Calculate the elapsed time between the start and stop events
cudaEventElapsedTime(&time, start, stop);
// Destroy the events
cudaEventDestroy(start);
cudaEventDestroy(stop);
// Print the elapsed time
std::cout << "Kernel runtime: " << std::fixed << std::setprecision(3) << time * 1e-3 << " sec." << std::endl;
// Copy the result back to the host
cudaMemcpy(&result, result_d, sizeof(projective_t), cudaMemcpyDeviceToHost);
// Print the result
std::cout << projective_t::to_affine(result) << std::endl;
// Free the device memory
cudaFree(scalars_d);
cudaFree(points_d);
cudaFree(result_d);
// Free the host memory, keep scalars for G2 example
delete[] points;
std::cout << "Part II: use G2 points" << std::endl;
std::cout << "Generating random inputs on-host" << std::endl;
// use the same scalars
g2_affine_t* g2_points = new g2_affine_t[N];
g2_projective_t::RandHostManyAffine(g2_points, N);
std::cout << "Reconfiguring MSM to use on-host inputs" << std::endl;
config.are_results_on_device = false;
config.are_scalars_on_device = false;
config.are_points_on_device = false;
g2_projective_t g2_result;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
msm::MSM<scalar_t, g2_affine_t, g2_projective_t>(scalars, g2_points, msm_size, config, &g2_result);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
std::cout << "Kernel runtime: " << std::fixed << std::setprecision(3) << time * 1e-3 << " sec." << std::endl;
std::cout << g2_projective_t::to_affine(g2_result) << std::endl;
std::cout << "Copying inputs on-device" << std::endl;
g2_affine_t* g2_points_d;
g2_projective_t* g2_result_d;
cudaMalloc(&scalars_d, sizeof(scalar_t) * N);
cudaMalloc(&g2_points_d, sizeof(g2_affine_t) * N);
cudaMalloc(&g2_result_d, sizeof(g2_projective_t));
cudaMemcpy(scalars_d, scalars, sizeof(scalar_t) * N, cudaMemcpyHostToDevice);
cudaMemcpy(g2_points_d, g2_points, sizeof(g2_affine_t) * N, cudaMemcpyHostToDevice);
std::cout << "Reconfiguring MSM to use on-device inputs" << std::endl;
config.are_results_on_device = true;
config.are_scalars_on_device = true;
config.are_points_on_device = true;
std::cout << "Running MSM kernel with on-device inputs" << std::endl;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, stream);
msm::MSM<scalar_t, g2_affine_t, g2_projective_t>(scalars_d, g2_points_d, msm_size, config, g2_result_d);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
std::cout << "Kernel runtime: " << std::fixed << std::setprecision(3) << time * 1e-3 << " sec." << std::endl;
cudaMemcpy(&g2_result, g2_result_d, sizeof(g2_projective_t), cudaMemcpyDeviceToHost);
std::cout << g2_projective_t::to_affine(g2_result) << std::endl;
cudaFree(scalars_d);
cudaFree(g2_points_d);
cudaFree(g2_result_d);
delete[] g2_points;
delete[] scalars;
cudaStreamDestroy(stream);
return 0;
}

2
examples/c++/msm/run.sh Executable file
View File

@@ -0,0 +1,2 @@
#!/bin/bash
./build/example

View File

@@ -0,0 +1,23 @@
# Make sure NVIDIA Container Toolkit is installed on your host
# Use NVIDIA base image
FROM nvidia/cuda:12.2.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
nsight-systems-12.2 \
cmake \
protobuf-compiler \
curl \
build-essential \
git \
&& rm -rf /var/lib/apt/lists/*
# Clone Icicle from a GitHub repository
RUN git clone https://github.com/ingonyama-zk/icicle.git /icicle
# Set the working directory in the container
WORKDIR /icicle-example
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,24 @@
{
"name": "Icicle Examples - Multiply",
"build": {
"dockerfile": "Dockerfile"
},
"workspaceMount": "source=${localWorkspaceFolder}/.,target=/icicle-example,type=bind",
"workspaceFolder": "/icicle-example",
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-azuretools.vscode-docker",
"ms-vscode.cpptools-extension-pack"
]
}
}
}

View File

@@ -0,0 +1,25 @@
cmake_minimum_required(VERSION 3.18)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
if (${CMAKE_VERSION} VERSION_LESS "3.24.0")
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH})
else()
set(CMAKE_CUDA_ARCHITECTURES native) # on 3.24+, on earlier it is ignored, and the target is not passed
endif ()
project(icicle LANGUAGES CUDA CXX)
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr")
set(CMAKE_CUDA_FLAGS_RELEASE "")
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -g -G -O0")
# change the path to your Icicle location
include_directories("../../../icicle")
add_executable(
example
example.cu
)
find_library(NVML_LIBRARY nvidia-ml PATHS /usr/local/cuda/targets/x86_64-linux/lib/stubs/ )
target_link_libraries(example ${NVML_LIBRARY})
set_target_properties(example PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

View File

@@ -0,0 +1,41 @@
# Icicle example: Multiplication
## Best-Practices
We recommend to run our examples in [ZK-containers](../../ZK-containers.md) to save your time and mental energy.
## Key-Takeaway
`Icicle` accelerates multiplication operation `*` using [Karatsuba algorithm](https://en.wikipedia.org/wiki/Karatsuba_algorithm)
## Concise Usage Explanation
Define a `CURVE_ID` and include curve configuration header:
```c++
#define CURVE_ID 1
#include "curves/curve_config.cuh"
```
The values of `CURVE_ID` for different curves are in the above header. Multiplication is accelerated both for field scalars and point fields.
```c++
using namespace curve_config;
scalar_t a;
point_field_t b;
```
## Running the example
- `cd` to your example directory
- compile with `./compile.sh`
- run with `./run.sh`
## What's in the example
1. Define the parameters for the example such as vector size
2. Generate random vectors on-host
3. Copy them on-device
4. Execute element-wise vector multiplication on-device
5. Copy results on-host

View File

@@ -0,0 +1,9 @@
#!/bin/bash
# Exit immediately on error
set -e
rm -rf build
mkdir -p build
cmake -S . -B build
cmake --build build

View File

@@ -0,0 +1,163 @@
#include <iostream>
#include <iomanip>
#include <chrono>
#include <nvml.h>
#define CURVE_ID 1
#include "curves/curve_config.cuh"
#include "utils/device_context.cuh"
#include "utils/vec_ops.cu"
using namespace curve_config;
// select scalar or point field
//typedef scalar_t T;
typedef point_field_t T;
int vector_mult(T* vec_b, T* vec_a, T* vec_result, size_t n_elments, device_context::DeviceContext ctx)
{
const bool is_on_device = true;
const bool is_montgomery = false;
cudaError_t err = vec_ops::Mul<T,T>(vec_a, vec_b, n_elments, is_on_device, is_montgomery, ctx, vec_result);
if (err != cudaSuccess) {
std::cerr << "Failed to multiply vectors - " << cudaGetErrorString(err) << std::endl;
return 0;
}
return 0;
}
int main(int argc, char** argv)
{
const unsigned vector_size = 1 << 15;
const unsigned repetitions = 1 << 15;
cudaError_t err;
nvmlInit();
nvmlDevice_t device;
nvmlDeviceGetHandleByIndex(0, &device); // for GPU 0
std::cout << "Icicle-Examples: vector multiplications" << std::endl;
char name[NVML_DEVICE_NAME_BUFFER_SIZE];
if (nvmlDeviceGetName(device, name, NVML_DEVICE_NAME_BUFFER_SIZE) == NVML_SUCCESS) {
std::cout << "GPU Model: " << name << std::endl;
} else {
std::cerr << "Failed to get GPU model name." << std::endl;
}
unsigned power_limit;
nvmlDeviceGetPowerManagementLimit(device, &power_limit);
std::cout << "Vector size: " << vector_size << std::endl;
std::cout << "Repetitions: " << repetitions << std::endl;
std::cout << "Power limit: " << std::fixed << std::setprecision(3) << 1.0e-3 * power_limit << " W" << std::endl;
unsigned int baseline_power;
nvmlDeviceGetPowerUsage(device, &baseline_power);
std::cout << "Baseline power: " << std::fixed << std::setprecision(3) << 1.0e-3 * baseline_power << " W" << std::endl;
unsigned baseline_temperature;
if (nvmlDeviceGetTemperature(device, NVML_TEMPERATURE_GPU, &baseline_temperature) == NVML_SUCCESS) {
std::cout << "Baseline GPU Temperature: " << baseline_temperature << " C" << std::endl;
} else {
std::cerr << "Failed to get GPU temperature." << std::endl;
}
// host data
T* host_in1 = (T*)malloc(vector_size * sizeof(T));
T* host_in2 = (T*)malloc(vector_size * sizeof(T));
std::cout << "Initializing vectors with random data" << std::endl;
T::RandHostMany(host_in1, vector_size);
T::RandHostMany(host_in2, vector_size);
// device data
device_context::DeviceContext ctx = device_context::get_default_device_context();
T* device_in1;
T* device_in2;
T* device_out;
err = cudaMalloc((void**)&device_in1, vector_size * sizeof(T));
if (err != cudaSuccess) {
std::cerr << "Failed to allocate device memory - " << cudaGetErrorString(err) << std::endl;
return 0;
}
err = cudaMalloc((void**)&device_in2, vector_size * sizeof(T));
if (err != cudaSuccess) {
std::cerr << "Failed to allocate device memory - " << cudaGetErrorString(err) << std::endl;
return 0;
}
err = cudaMalloc((void**)&device_out, vector_size * sizeof(T));
if (err != cudaSuccess) {
std::cerr << "Failed to allocate device memory - " << cudaGetErrorString(err) << std::endl;
return 0;
}
// copy from host to device
err = cudaMemcpy(device_in1, host_in1, vector_size * sizeof(T), cudaMemcpyHostToDevice);
if (err != cudaSuccess) {
std::cerr << "Failed to copy data from host to device - " << cudaGetErrorString(err) << std::endl;
return 0;
}
err = cudaMemcpy(device_in2, host_in2, vector_size * sizeof(T), cudaMemcpyHostToDevice);
if (err != cudaSuccess) {
std::cerr << "Failed to copy data from host to device - " << cudaGetErrorString(err) << std::endl;
return 0;
}
std::cout << "Starting warm-up" << std::endl;
// Warm-up loop
for (int i = 0; i < repetitions; i++) {
vector_mult(device_in1, device_in2, device_out, vector_size, ctx);
}
std::cout << "Starting benchmarking" << std::endl;
unsigned power_before;
nvmlDeviceGetPowerUsage(device, &power_before);
std::cout << "Power before: " << std::fixed << std::setprecision(3) << 1.0e-3 * power_before << " W" << std::endl;
std::cout << "Power utilization: " << std::fixed << std::setprecision(1) << (float)100.0 * power_before / power_limit
<< " %" << std::endl;
unsigned temperature_before;
if (nvmlDeviceGetTemperature(device, NVML_TEMPERATURE_GPU, &temperature_before) == NVML_SUCCESS) {
std::cout << "GPU Temperature before: " << temperature_before << " C" << std::endl;
} else {
std::cerr << "Failed to get GPU temperature." << std::endl;
}
auto start_time = std::chrono::high_resolution_clock::now();
// Benchmark loop
for (int i = 0; i < repetitions; i++) {
vector_mult(device_in1, device_in2, device_out, vector_size, ctx);
}
auto end_time = std::chrono::high_resolution_clock::now();
auto duration = std::chrono::duration_cast<std::chrono::microseconds>(end_time - start_time);
std::cout << "Elapsed time: " << duration.count() << " microseconds" << std::endl;
unsigned power_after;
nvmlDeviceGetPowerUsage(device, &power_after);
std::cout << "Power after: " << std::fixed << std::setprecision(3) << 1.0e-3 * power_after << " W" << std::endl;
std::cout << "Power utilization: " << std::fixed << std::setprecision(1) << (float)100.0 * power_after / power_limit
<< " %" << std::endl;
unsigned temperature_after;
if (nvmlDeviceGetTemperature(device, NVML_TEMPERATURE_GPU, &temperature_after) == NVML_SUCCESS) {
std::cout << "GPU Temperature after: " << temperature_after << " C" << std::endl;
} else {
std::cerr << "Failed to get GPU temperature." << std::endl;
}
// Report performance in GMPS: Giga Multiplications Per Second
double GMPS = 1.0e-9 * repetitions * vector_size / (1.0e-6 * duration.count());
std::cout << "Performance: " << GMPS << " Giga Multiplications Per Second" << std::endl;
// Optional: validate multiplication
T* host_out = (T*)malloc(vector_size * sizeof(T));
cudaMemcpy(host_out, device_out, vector_size * sizeof(T), cudaMemcpyDeviceToHost);
// validate multiplication here...
// clean up and exit
free(host_in1);
free(host_in2);
free(host_out);
cudaFree(device_in1);
cudaFree(device_in2);
cudaFree(device_out);
nvmlShutdown();
return 0;
}

2
examples/c++/multiply/run.sh Executable file
View File

@@ -0,0 +1,2 @@
#!/bin/bash
./build/example

View File

@@ -0,0 +1,25 @@
# Make sure NVIDIA Container Toolkit is installed on your host
# Use the specified base image
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
curl \
build-essential \
git \
libboost-all-dev \
&& rm -rf /var/lib/apt/lists/*
# Clone Icicle from a GitHub repository
RUN git clone https://github.com/ingonyama-zk/icicle.git /icicle
# Set the working directory in the container
WORKDIR /icicle-example
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,22 @@
{
"name": "Icicle Examples: ntt",
"build": {
"dockerfile": "Dockerfile"
},
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-python.python",
"ms-vscode.cpptools"
]
}
}
}

View File

@@ -0,0 +1,26 @@
cmake_minimum_required(VERSION 3.18)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED TRUE)
set(CMAKE_CXX_STANDARD_REQUIRED TRUE)
if (${CMAKE_VERSION} VERSION_LESS "3.24.0")
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCH})
else()
set(CMAKE_CUDA_ARCHITECTURES native) # on 3.24+, on earlier it is ignored, and the target is not passed
endif ()
project(icicle LANGUAGES CUDA CXX)
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr")
set(CMAKE_CUDA_FLAGS_RELEASE "")
set(CMAKE_CUDA_FLAGS_DEBUG "${CMAKE_CUDA_FLAGS_DEBUG} -g -G -O0")
# change the path to your Icicle location
include_directories("../../../icicle")
add_executable(
example
example.cu
)
find_library(NVML_LIBRARY nvidia-ml PATHS /usr/local/cuda-12.0/targets/x86_64-linux/lib/stubs/ )
target_link_libraries(example ${NVML_LIBRARY})
set_target_properties(example PROPERTIES CUDA_SEPARABLE_COMPILATION ON)

View File

@@ -0,0 +1,32 @@
# Icicle example: Number-Theoretical Transform (NTT)
## Best-Practices
We recommend to run our examples in [ZK-containers](../../ZK-containers.md) to save your time and mental energy.
## Key-Takeaway
`Icicle` provides CUDA C++ template function NTT for [Number Theoretical Transform](https://github.com/ingonyama-zk/ingopedia/blob/master/src/fft.md), also known as Discrete Fourier Transform.
## Concise Usage Explanation
```c++
// Select the curve
#define CURVE_ID 1
// Include NTT template
#include "appUtils/ntt/ntt.cu"
using namespace curve_config;
// Configure NTT
ntt::NTTConfig<S> config=ntt::DefaultNTTConfig<S>();
// Call NTT
ntt::NTT<S, E>(input, ntt_size, ntt::NTTDir::kForward, config, output);
```
## Running the example
- `cd` to your example directory
- compile with `./compile.sh`
- run with `./run.sh`

11
examples/c++/ntt/compile.sh Executable file
View File

@@ -0,0 +1,11 @@
#!/bin/bash
# Exit immediately on error
set -e
rm -rf build
mkdir -p build
cmake -S . -B build
cmake --build build

View File

@@ -0,0 +1,99 @@
#include <chrono>
#include <iostream>
// select the curve
#define CURVE_ID 1
// include NTT template
#include "appUtils/ntt/ntt.cu"
using namespace curve_config;
// Operate on scalars
typedef scalar_t S;
typedef scalar_t E;
void print_elements(const unsigned n, E * elements ) {
for (unsigned i = 0; i < n; i++) {
std::cout << i << ": " << elements[i] << std::endl;
}
}
void initialize_input(const unsigned ntt_size, const unsigned nof_ntts, E * elements ) {
// Lowest Harmonics
for (unsigned i = 0; i < ntt_size; i=i+1) {
elements[i] = E::one();
}
// print_elements(ntt_size, elements );
// Highest Harmonics
for (unsigned i = 1*ntt_size; i < 2*ntt_size; i=i+2) {
elements[i] = E::one();
elements[i+1] = E::neg(scalar_t::one());
}
// print_elements(ntt_size, &elements[1*ntt_size] );
}
int validate_output(const unsigned ntt_size, const unsigned nof_ntts, E* elements)
{
int nof_errors = 0;
E amplitude = E::from((uint32_t) ntt_size);
// std::cout << "Amplitude: " << amplitude << std::endl;
// Lowest Harmonics
if (elements[0] != amplitude) {
++nof_errors;
std::cout << "Error in lowest harmonics 0! " << std::endl;
// print_elements(ntt_size, elements );
} else {
std::cout << "Validated lowest harmonics" << std::endl;
}
// Highest Harmonics
if (elements[1*ntt_size+ntt_size/2] != amplitude) {
++nof_errors;
std::cout << "Error in highest harmonics! " << std::endl;
// print_elements(ntt_size, &elements[1*ntt_size] );
} else {
std::cout << "Validated highest harmonics" << std::endl;
}
return nof_errors;
}
int main(int argc, char* argv[])
{
std::cout << "Icicle Examples: Number Theoretical Transform (NTT)" << std::endl;
std::cout << "Example parameters" << std::endl;
const unsigned log_ntt_size = 20;
std::cout << "Log2(NTT size): " << log_ntt_size << std::endl;
const unsigned ntt_size = 1 << log_ntt_size;
std::cout << "NTT size: " << ntt_size << std::endl;
const unsigned nof_ntts = 2;
std::cout << "Number of NTTs: " << nof_ntts << std::endl;
const unsigned batch_size = nof_ntts * ntt_size;
std::cout << "Generating input data for lowest and highest harmonics" << std::endl;
E* input;
input = (E*) malloc(sizeof(E) * batch_size);
initialize_input(ntt_size, nof_ntts, input );
E* output;
output = (E*) malloc(sizeof(E) * batch_size);
std::cout << "Running NTT with on-host data" << std::endl;
cudaStream_t stream;
cudaStreamCreate(&stream);
// Create a device context
auto ctx = device_context::get_default_device_context();
// the next line is valid only for CURVE_ID 1 (will add support for other curves soon)
S rou = S{ {0x53337857, 0x53422da9, 0xdbed349f, 0xac616632, 0x6d1e303, 0x27508aba, 0xa0ed063, 0x26125da1} };
ntt::InitDomain(rou, ctx);
// Create an NTTConfig instance
ntt::NTTConfig<S> config=ntt::DefaultNTTConfig<S>();
config.batch_size = nof_ntts;
config.ctx.stream = stream;
auto begin0 = std::chrono::high_resolution_clock::now();
cudaError_t err = ntt::NTT<S, E>(input, ntt_size, ntt::NTTDir::kForward, config, output);
auto end0 = std::chrono::high_resolution_clock::now();
auto elapsed0 = std::chrono::duration_cast<std::chrono::nanoseconds>(end0 - begin0);
printf("On-device runtime: %.3f seconds\n", elapsed0.count() * 1e-9);
validate_output(ntt_size, nof_ntts, output );
cudaStreamDestroy(stream);
free(input);
free(output);
return 0;
}

2
examples/c++/ntt/run.sh Executable file
View File

@@ -0,0 +1,2 @@
#!/bin/bash
./build/example

View File

@@ -0,0 +1,27 @@
# Use the specified base image
#FROM nvidia/cuda:12.2.0-devel-ubuntu22.04
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
protobuf-compiler \
curl \
build-essential \
git \
llvm \
clang \
&& rm -rf /var/lib/apt/lists/*
# Install Rust
RUN curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
ENV PATH="/root/.cargo/bin:${PATH}"
# Set the working directory in the container
WORKDIR /icicle-example
# Copy the content of the local directory to the working directory
COPY . .
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,23 @@
{
"name": "Icicle Examples: rust msm",
"build": {
"dockerfile": "Dockerfile"
},
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-azuretools.vscode-docker",
"rust-lang.rust-analyzer",
"vadimcn.vscode-lldb"
]
}
}
}

View File

@@ -0,0 +1,19 @@
[package]
name = "msm"
version = "1.0.0"
edition = "2018"
[dependencies]
icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" }
icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" }
icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0", features = [ "g2" ] }
icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" }
ark-bn254 = { version = "0.4.0", optional = true}
ark-bls12-377 = { version = "0.4.0", optional = true}
ark-ec = { version = "0.4.0", optional = true}
clap = { version = "4.4.12", features = ["derive"] }
[features]
arkworks = ["ark-bn254", "ark-bls12-377", "ark-ec", "icicle-core/arkworks", "icicle-bn254/arkworks", "icicle-bls12-377/arkworks"]
profile = []
g2 = []

View File

@@ -0,0 +1,56 @@
# ICICLE example: MultiScalar Multiplication (MSM) in Rust
`ICICLE` provides Rust bindings to CUDA-accelerated C++ implementation of [Multi-Scalar Multiplication](https://github.com/ingonyama-zk/ingopedia/blob/master/src/msm.md).
## Best Practices
In order to save time and setting up prerequisites manually, we recommend running this example in our [ZKContainer](../../ZKContainer.md).
## Usage
```rust
msm(
/* Scalars input vector */ &scalars,
/* Points input vector */ &points,
/* MSMConfig reference */ &cfg,
/* Projective point result */ &mut msm_results.as_slice()
).unwrap();
```
In this example we use `BN254` curve. The function computes $result = \sum_{i=0}^{size-1} scalars[i] \cdot points[i]$, where input `points[]` uses affine coordinates, and `result` uses projective coordinates.
## What's in the example
1. Define the size of MSM.
2. Generate random inputs on-device
3. Configure MSM
4. Execute MSM on-device
5. Move the result on host
Running the example:
```sh
cargo run --release
```
You can add the `--feature arkworks,profile` flag to measure times of both ICICLE and arkworks.
> [!NOTE]
> The default sizes are 2^19 - 2^23. You can change this by passing the `--lower_bound_log_size <size> --upper_bound_log_size <size>` options. To change the size range to 2^21 - 2^24, run the example like this:
> ```sh
> cargo run --release -- -l 21 -u 24
> ```
## Benchmarks
These benchmarks were run on a 16 core 24 thread i9-12900k CPU and an RTX 3090 Ti GPU
### Single BN254 MSM
| Library\Size | 2^19 | 2^20 | 2^21 | 2^22 | 2^23 |
|--------------|------|------|------|------|------|
| ICICLE | 10 ms | 11 ms | 21 ms | 39 ms | 77 ms |
| Arkworks | 284 ms | 540 ms | 1,152 ms | 2,320 ms | 4,491 ms |
### Single BLS12377 MSM
| Library\Size | 2^19 | 2^20 | 2^21 | 2^22 | 2^23 |
|--------------|------|------|------|------|------|
| ICICLE | 9 ms | 14 ms | 25 ms | 48 ms | 93 ms |
| Arkworks | 490 ms | 918 ms | 1,861 ms | 3,624 ms | 7,191 ms |

View File

@@ -0,0 +1,192 @@
use icicle_bn254::curve::{
CurveCfg,
ScalarCfg,
G1Projective,
G2CurveCfg,
G2Projective
};
use icicle_bls12_377::curve::{
CurveCfg as BLS12377CurveCfg,
ScalarCfg as BLS12377ScalarCfg,
G1Projective as BLS12377G1Projective
};
use icicle_cuda_runtime::{
stream::CudaStream,
memory::HostOrDeviceSlice
};
use icicle_core::{
msm,
curve::Curve,
traits::GenerateRandom
};
#[cfg(feature = "arkworks")]
use icicle_core::traits::ArkConvertible;
#[cfg(feature = "arkworks")]
use ark_bn254::{
G1Projective as Bn254ArkG1Projective,
G1Affine as Bn254G1Affine,
Fr as Bn254Fr
};
#[cfg(feature = "arkworks")]
use ark_bls12_377::{
G1Projective as Bls12377ArkG1Projective,
G1Affine as Bls12377G1Affine,
Fr as Bls12377Fr
};
#[cfg(feature = "arkworks")]
use ark_ec::scalar_mul::variable_base::VariableBaseMSM;
#[cfg(feature = "profile")]
use std::time::Instant;
use clap::Parser;
#[derive(Parser, Debug)]
struct Args {
/// Lower bound (inclusive) of MSM sizes to run for
#[arg(short, long, default_value_t = 19)]
lower_bound_log_size: u8,
/// Upper bound of MSM sizes to run for
#[arg(short, long, default_value_t = 23)]
upper_bound_log_size: u8,
}
fn main() {
let args = Args::parse();
let lower_bound = args.lower_bound_log_size;
let upper_bound = args.upper_bound_log_size;
println!("Running Icicle Examples: Rust MSM");
let upper_size = 1 << (upper_bound);
println!("Generating random inputs on host for bn254...");
let upper_points = CurveCfg::generate_random_affine_points(upper_size);
let g2_upper_points = G2CurveCfg::generate_random_affine_points(upper_size);
let upper_scalars = ScalarCfg::generate_random(upper_size);
println!("Generating random inputs on host for bls12377...");
let upper_points_bls12377 = BLS12377CurveCfg::generate_random_affine_points(upper_size);
let upper_scalars_bls12377 = BLS12377ScalarCfg::generate_random(upper_size);
for i in lower_bound..=upper_bound {
let log_size = i;
let size = 1 << log_size;
println!("---------------------- MSM size 2^{}={} ------------------------", log_size, size);
// Setting Bn254 points and scalars
let points = HostOrDeviceSlice::Host(upper_points[..size].to_vec());
let g2_points = HostOrDeviceSlice::Host(g2_upper_points[..size].to_vec());
let scalars = HostOrDeviceSlice::Host(upper_scalars[..size].to_vec());
// Setting bls12377 points and scalars
// let points_bls12377 = &upper_points_bls12377[..size];
let points_bls12377 = HostOrDeviceSlice::Host(upper_points_bls12377[..size].to_vec()); // &upper_points_bls12377[..size];
let scalars_bls12377 = HostOrDeviceSlice::Host(upper_scalars_bls12377[..size].to_vec());
println!("Configuring bn254 MSM...");
let mut msm_results: HostOrDeviceSlice<'_, G1Projective> = HostOrDeviceSlice::cuda_malloc(1).unwrap();
let mut g2_msm_results: HostOrDeviceSlice<'_, G2Projective> = HostOrDeviceSlice::cuda_malloc(1).unwrap();
let stream = CudaStream::create().unwrap();
let g2_stream = CudaStream::create().unwrap();
let mut cfg = msm::get_default_msm_config::<CurveCfg>();
let mut g2_cfg = msm::get_default_msm_config::<G2CurveCfg>();
cfg.ctx.stream = &stream;
g2_cfg.ctx.stream = &g2_stream;
cfg.is_async = true;
g2_cfg.is_async = true;
println!("Configuring bls12377 MSM...");
let mut msm_results_bls12377: HostOrDeviceSlice<'_, BLS12377G1Projective> = HostOrDeviceSlice::cuda_malloc(1).unwrap();
let stream_bls12377 = CudaStream::create().unwrap();
let mut cfg_bls12377 = msm::get_default_msm_config::<BLS12377CurveCfg>();
cfg_bls12377.ctx.stream = &stream_bls12377;
cfg_bls12377.is_async = true;
println!("Executing bn254 MSM on device...");
#[cfg(feature = "profile")]
let start = Instant::now();
msm::msm(&scalars, &points, &cfg, &mut msm_results).unwrap();
#[cfg(feature = "profile")]
println!("ICICLE BN254 MSM on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
msm::msm(&scalars, &g2_points, &g2_cfg, &mut g2_msm_results).unwrap();
println!("Executing bls12377 MSM on device...");
#[cfg(feature = "profile")]
let start = Instant::now();
msm::msm(&scalars_bls12377, &points_bls12377, &cfg_bls12377, &mut msm_results_bls12377 ).unwrap();
#[cfg(feature = "profile")]
println!("ICICLE BLS12377 MSM on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
println!("Moving results to host..");
let mut msm_host_result = vec![G1Projective::zero(); 1];
let mut g2_msm_host_result = vec![G2Projective::zero(); 1];
let mut msm_host_result_bls12377 = vec![BLS12377G1Projective::zero(); 1];
stream
.synchronize()
.unwrap();
g2_stream
.synchronize()
.unwrap();
msm_results
.copy_to_host(&mut msm_host_result[..])
.unwrap();
g2_msm_results
.copy_to_host(&mut g2_msm_host_result[..])
.unwrap();
println!("bn254 result: {:#?}", msm_host_result);
println!("G2 bn254 result: {:#?}", g2_msm_host_result);
stream_bls12377
.synchronize()
.unwrap();
msm_results_bls12377
.copy_to_host(&mut msm_host_result_bls12377[..])
.unwrap();
println!("bls12377 result: {:#?}", msm_host_result_bls12377);
#[cfg(feature = "arkworks")]
{
println!("Checking against arkworks...");
let ark_points: Vec<Bn254G1Affine> = points.as_slice().iter().map(|&point| point.to_ark()).collect();
let ark_scalars: Vec<Bn254Fr> = scalars.as_slice().iter().map(|scalar| scalar.to_ark()).collect();
let ark_points_bls12377: Vec<Bls12377G1Affine> = points_bls12377.as_slice().iter().map(|point| point.to_ark()).collect();
let ark_scalars_bls12377: Vec<Bls12377Fr> = scalars_bls12377.as_slice().iter().map(|scalar| scalar.to_ark()).collect();
#[cfg(feature = "profile")]
let start = Instant::now();
let bn254_ark_msm_res = Bn254ArkG1Projective::msm(&ark_points, &ark_scalars).unwrap();
println!("Arkworks Bn254 result: {:#?}", bn254_ark_msm_res);
#[cfg(feature = "profile")]
println!("Ark BN254 MSM on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
#[cfg(feature = "profile")]
let start = Instant::now();
let bls12377_ark_msm_res = Bls12377ArkG1Projective::msm(&ark_points_bls12377, &ark_scalars_bls12377).unwrap();
println!("Arkworks Bls12377 result: {:#?}", bls12377_ark_msm_res);
#[cfg(feature = "profile")]
println!("Ark BLS12377 MSM on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
let bn254_icicle_msm_res_as_ark = msm_host_result[0].to_ark();
let bls12377_icicle_msm_res_as_ark = msm_host_result_bls12377[0].to_ark();
println!("Bn254 MSM is correct: {}", bn254_ark_msm_res.eq(&bn254_icicle_msm_res_as_ark));
println!("Bls12377 MSM is correct: {}", bls12377_ark_msm_res.eq(&bls12377_icicle_msm_res_as_ark));
}
println!("Cleaning up bn254...");
stream
.destroy()
.unwrap();
println!("Cleaning up bls12377...");
stream_bls12377
.destroy()
.unwrap();
println!("");
}
}

View File

@@ -0,0 +1,27 @@
# Use the specified base image
#FROM nvidia/cuda:12.2.0-devel-ubuntu22.04
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
protobuf-compiler \
curl \
build-essential \
git \
llvm \
clang \
&& rm -rf /var/lib/apt/lists/*
# Install Rust
RUN curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
ENV PATH="/root/.cargo/bin:${PATH}"
# Set the working directory in the container
WORKDIR /icicle-example
# Copy the content of the local directory to the working directory
COPY . .
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,23 @@
{
"name": "Icicle Examples: rust ntt",
"build": {
"dockerfile": "Dockerfile"
},
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-azuretools.vscode-docker",
"rust-lang.rust-analyzer",
"vadimcn.vscode-lldb"
]
}
}
}

View File

@@ -0,0 +1,20 @@
[package]
name = "ntt"
version = "1.0.0"
edition = "2018"
[dependencies]
icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0" }
icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0", features = ["arkworks"] }
icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0", features = ["arkworks"] }
icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.1.0", features = ["arkworks"] }
ark-ff = { version = "0.4.0" }
ark-poly = "0.4.0"
ark-std = "0.4.0"
ark-bn254 = { version = "0.4.0" }
ark-bls12-377 = { version = "0.4.0" }
clap = { version = "4.4.12", features = ["derive"] }
[features]
profile = []

View File

@@ -0,0 +1,65 @@
# ICICLE example: Number Theoretic Transform (NTT) in Rust
## Key-Takeaway
`ICICLE` provides Rust bindings to CUDA-accelerated C++ implementation of [Number Theoretic Transform](https://github.com/ingonyama-zk/ingopedia/blob/master/src/fft.md).
## Best Practices
In order to save time and setting up prerequisites manually, we recommend running this example in our [ZKContainer](../../ZKContainer.md).
## Usage
```rust
ntt::ntt(
/* input slice */ scalars.as_slice(),
/* NTT Direction */ ntt::NTTDir::kForward,
/* NTT Configuration */ &cfg,
/* output slice */ ntt_results.as_slice()
).unwrap();
```
In this example we use the `BN254` and `BLS12377` fields.
## What's in this example
1. Define the size of NTT.
2. Generate random inputs on-host
3. Set up the domain.
4. Configure NTT
5. Execute NTT on-device
6. Move the result on host
7. Compare results with arkworks
Running the example:
```sh
cargo run --release
```
You can add the `--feature profile` flag to measure times of both ICICLE and arkworks.
> [!NOTE]
> The default size is 2^20. You can change this by passing the `--size <size>` option. To change the size to 2^23, run the example like this:
```sh
cargo run --release -- -s 23
```
## Benchmarks
These benchmarks were run on a 16 core 24 thread i9-12900k CPU and an RTX 3090 Ti GPU
### Single BN254 NTT
| Library\Size | 2^19 | 2^20 | 2^21 | 2^22 | 2^23 |
|--------------|------|------|------|------|------|
| ICICLE | 1.263 ms | 2.986 ms | 4.651 ms | 9.308 ms | 18.618 ms |
| Arkworks | 138 ms | 290 ms | 611 ms | 1,295 ms | 2,715 ms |
### Single BLS12377 NTT
| Library\Size | 2^19 | 2^20 | 2^21 | 2^22 | 2^23 |
|--------------|------|------|------|------|------|
| ICICLE | 1.272 ms | 2.893 ms | 4.728 ms | 9.211 ms | 18.319 ms |
| Arkworks | 135 ms | 286 ms | 605 ms | 1,279 ms | 2,682 ms |

View File

@@ -0,0 +1,157 @@
use icicle_bn254::curve::{
ScalarCfg,
ScalarField,
};
use icicle_bls12_377::curve::{
ScalarCfg as BLS12377ScalarCfg,
ScalarField as BLS12377ScalarField
};
use icicle_cuda_runtime::{
stream::CudaStream,
memory::HostOrDeviceSlice,
device_context::get_default_device_context
};
use icicle_core::{
ntt::{self, NTT},
traits::{GenerateRandom, FieldImpl}
};
use icicle_core::traits::ArkConvertible;
use ark_bn254::Fr as Bn254Fr;
use ark_bls12_377::Fr as Bls12377Fr;
use ark_ff::FftField;
use ark_poly::{EvaluationDomain, Radix2EvaluationDomain};
use ark_std::cmp::{Ord, Ordering};
use std::convert::TryInto;
#[cfg(feature = "profile")]
use std::time::Instant;
use clap::Parser;
#[derive(Parser, Debug)]
struct Args {
/// Size of NTT to run (20 for 2^20)
#[arg(short, long, default_value_t = 20)]
size: u8,
}
fn main() {
let args = Args::parse();
println!("Running Icicle Examples: Rust NTT");
let log_size = args.size;
let size = 1 << log_size;
println!("---------------------- NTT size 2^{}={} ------------------------", log_size, size);
// Setting Bn254 points and scalars
println!("Generating random inputs on host for bn254...");
let scalars = HostOrDeviceSlice::Host(ScalarCfg::generate_random(size));
let mut ntt_results: HostOrDeviceSlice<'_, ScalarField> = HostOrDeviceSlice::cuda_malloc(size).unwrap();
// Setting bls12377 points and scalars
println!("Generating random inputs on host for bls12377...");
let scalars_bls12377 = HostOrDeviceSlice::Host(BLS12377ScalarCfg::generate_random(size));
let mut ntt_results_bls12377: HostOrDeviceSlice<'_, BLS12377ScalarField> = HostOrDeviceSlice::cuda_malloc(size).unwrap();
println!("Setting up bn254 Domain...");
let icicle_omega = <Bn254Fr as FftField>::get_root_of_unity(size.try_into().unwrap()).unwrap();
let ctx = get_default_device_context();
ScalarCfg::initialize_domain(ScalarField::from_ark(icicle_omega), &ctx).unwrap();
println!("Configuring bn254 NTT...");
let stream = CudaStream::create().unwrap();
let mut cfg = ntt::get_default_ntt_config::<ScalarField>();
cfg.ctx.stream = &stream;
cfg.is_async = true;
println!("Setting up bls12377 Domain...");
let icicle_omega = <Bls12377Fr as FftField>::get_root_of_unity(size.try_into().unwrap()).unwrap();
// reusing ctx from above
BLS12377ScalarCfg::initialize_domain(BLS12377ScalarField::from_ark(icicle_omega), &ctx).unwrap();
println!("Configuring bls12377 NTT...");
let stream_bls12377 = CudaStream::create().unwrap();
let mut cfg_bls12377 = ntt::get_default_ntt_config::<BLS12377ScalarField>();
cfg_bls12377.ctx.stream = &stream_bls12377;
cfg_bls12377.is_async = true;
println!("Executing bn254 NTT on device...");
#[cfg(feature = "profile")]
let start = Instant::now();
ntt::ntt(&scalars, ntt::NTTDir::kForward, &cfg, &mut ntt_results).unwrap();
#[cfg(feature = "profile")]
println!("ICICLE BN254 NTT on size 2^{log_size} took: {} μs", start.elapsed().as_micros());
println!("Executing bls12377 NTT on device...");
#[cfg(feature = "profile")]
let start = Instant::now();
ntt::ntt(&scalars_bls12377, ntt::NTTDir::kForward, &cfg_bls12377, &mut ntt_results_bls12377).unwrap();
#[cfg(feature = "profile")]
println!("ICICLE BLS12377 NTT on size 2^{log_size} took: {} μs", start.elapsed().as_micros());
println!("Moving results to host..");
stream
.synchronize()
.unwrap();
let mut host_bn254_results = vec![ScalarField::zero(); size];
ntt_results
.copy_to_host(&mut host_bn254_results[..])
.unwrap();
stream_bls12377
.synchronize()
.unwrap();
let mut host_bls12377_results = vec![BLS12377ScalarField::zero(); size];
ntt_results_bls12377
.copy_to_host(&mut host_bls12377_results[..])
.unwrap();
println!("Checking against arkworks...");
let mut ark_scalars: Vec<Bn254Fr> = scalars.as_slice().iter().map(|scalar| scalar.to_ark()).collect();
let bn254_domain = <Radix2EvaluationDomain<Bn254Fr> as EvaluationDomain<Bn254Fr>>::new(size).unwrap();
let mut ark_scalars_bls12377: Vec<Bls12377Fr> = scalars_bls12377.as_slice().iter().map(|scalar| scalar.to_ark()).collect();
let bls12_377_domain = <Radix2EvaluationDomain<Bls12377Fr> as EvaluationDomain<Bls12377Fr>>::new(size).unwrap();
#[cfg(feature = "profile")]
let start = Instant::now();
bn254_domain.fft_in_place(&mut ark_scalars);
#[cfg(feature = "profile")]
println!("Ark BN254 NTT on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
#[cfg(feature = "profile")]
let start = Instant::now();
bls12_377_domain.fft_in_place(&mut ark_scalars_bls12377);
#[cfg(feature = "profile")]
println!("Ark BLS12377 NTT on size 2^{log_size} took: {} ms", start.elapsed().as_millis());
host_bn254_results
.iter()
.zip(ark_scalars.iter())
.for_each(|(icicle_scalar, &ark_scalar)| {
assert_eq!(ark_scalar.cmp(&icicle_scalar.to_ark()), Ordering::Equal);
});
println!("Bn254 NTT is correct");
host_bls12377_results
.iter()
.zip(ark_scalars_bls12377.iter())
.for_each(|(icicle_scalar, &ark_scalar)| {
assert_eq!(ark_scalar.cmp(&icicle_scalar.to_ark()), Ordering::Equal);
});
println!("Bls12377 NTT is correct");
println!("Cleaning up bn254...");
stream
.destroy()
.unwrap();
println!("Cleaning up bls12377...");
stream_bls12377
.destroy()
.unwrap();
println!("");
}

View File

@@ -0,0 +1,27 @@
# Use the specified base image
#FROM nvidia/cuda:12.2.0-devel-ubuntu22.04
FROM nvidia/cuda:12.0.0-devel-ubuntu22.04
# Update and install dependencies
RUN apt-get update && apt-get install -y \
cmake \
protobuf-compiler \
curl \
build-essential \
git \
llvm \
clang \
&& rm -rf /var/lib/apt/lists/*
# Install Rust
RUN curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
ENV PATH="/root/.cargo/bin:${PATH}"
# Set the working directory in the container
WORKDIR /icicle-example
# Copy the content of the local directory to the working directory
COPY . .
# Specify the default command for the container
CMD ["/bin/bash"]

View File

@@ -0,0 +1,23 @@
{
"name": "Icicle Examples: rust poseidon hash",
"build": {
"dockerfile": "Dockerfile"
},
"runArgs": [
"--gpus",
"all"
],
"postCreateCommand": [
"nvidia-smi"
],
"customizations": {
"vscode": {
"extensions": [
"ms-vscode.cmake-tools",
"ms-azuretools.vscode-docker",
"rust-lang.rust-analyzer",
"vadimcn.vscode-lldb"
]
}
}
}

View File

@@ -0,0 +1,14 @@
[package]
name = "posedion"
version = "1.0.0"
edition = "2018"
[dependencies]
icicle-cuda-runtime = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.2.0" }
icicle-core = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.2.0", features = ["arkworks"] }
icicle-bn254 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.2.0", features = ["arkworks"] }
icicle-bls12-377 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.2.0", features = ["arkworks"] }
icicle-bls12-381 = { git = "https://github.com/ingonyama-zk/icicle.git", tag = "v1.2.0", features = ["arkworks"] }
[features]
profile = []

View File

@@ -0,0 +1,50 @@
# ICICLE example: Poseidon hash in Rust
## Key-Takeaway
`ICICLE` provides Rust bindings to CUDA-accelerated C++ implementation of [Poseidon hash](https://github.com/ingonyama-zk/ingopedia/blob/9f602aae051100ee4c60791db5c6fa23d01e1f79/src/hashzk.md?plain=1#L30).
## Best Practices
In order to save time and setting up prerequisites manually, we recommend running this example in our [ZKContainer](../../ZKContainer.md).
## Usage
```rust
poseidon::poseidon_hash_many<F>(
input: &mut HostOrDeviceSlice<F>, // a pointer to a vector of input data
output: &mut HostOrDeviceSlice<F>, // a pointer to a vector of output data,
number_of_states: u32, // number of input blocks of size `arity`
arity: u32, // the arity of the hash function
constants: &PoseidonConstants<F>, // Poseidon constants
config: &PoseidonConfig, // config used to specify extra arguments of the Poseidon
) -> IcicleResult<()>
```
In this example we use the `BN254`, `BLS12377` and `BLS12381` fields.
## What's in this example
1. Load optimized Poseidon hash constants.
2. Generate custom Poseidon hash constants.
3. Configure Poseidon hash to use inputs and outputs on device
4. Execute Poseidon Hash on-device
Running the example:
```sh
cargo run --release
```
You can add the `--feature profile` flag to measure times of both ICICLE and arkworks.
> [!NOTE]
> The default size is 2^20. You can change this by passing the `--size <size>` option. To change the size to 2^23, run the example like this:
```sh
cargo run --release -- -s 23
```
## Benchmarks
TODO

View File

@@ -0,0 +1,50 @@
use icicle_bls12_381::poseidon;
use icicle_cuda_runtime::device_context::get_default_device_context;
use icicle_core::poseidon::{load_optimized_poseidon_constants, poseidon_hash_many, PoseidonConfig};
#[cfg(feature = "profile")]
use std::time::Instant;
use clap::Parser;
#[derive(Parser, Debug)]
struct Args {
/// Size of Poseidon input to run (20 for 2^20)
#[arg(short, long, default_value_t = 20)]
size: u8,
}
fn main() {
let size = args.size;
let test_size = 1 << size;
println!("Running Icicle Examples: Rust Poseidon Hash");
let arity = 2u32;
println!("---------------------- Loading optimized Poseidon constants for arity={} ------------------------", arity);
let ctx = get_default_device_context();
let constants = load_optimized_poseidon_constants::<F>(arity, &ctx).unwrap();
let config = PoseidonConfig::default();
println!("---------------------- Input size 2^{}={} ------------------------", size, test_size);
let inputs = vec![F::one(); test_size * arity as usize];
let outputs = vec![F::zero(); test_size];
let mut input_slice = HostOrDeviceSlice::on_host(inputs);
let mut output_slice = HostOrDeviceSlice::on_host(outputs);
println!("Executing BLS12-381 Poseidon Hash on device...");
#[cfg(feature = "profile")]
let start = Instant::now();
poseidon_hash_many::<F>(
&mut input_slice,
&mut output_slice,
test_size as u32,
arity as u32,
&constants,
&config,
)
.unwrap();
#[cfg(feature = "profile")]
println!("ICICLE BLS12-381 Poseidon Hash on size 2^{size} took: {} μs", start.elapsed().as_micros());
}

17
go.mod Normal file
View File

@@ -0,0 +1,17 @@
module github.com/ingonyama-zk/icicle
go 1.20
require (
github.com/davecgh/go-spew v1.1.1 // indirect
github.com/kr/pretty v0.1.0 // indirect
github.com/pmezard/go-difflib v1.0.0 // indirect
gopkg.in/check.v1 v1.0.0-20180628173108-788fd7840127 // indirect
gopkg.in/yaml.v3 v3.0.1 // indirect
)
require (
github.com/consensys/bavard v0.1.13
github.com/stretchr/testify v1.8.3
rsc.io/tmplfunc v0.0.3 // indirect
)

20
go.sum Normal file
View File

@@ -0,0 +1,20 @@
github.com/consensys/bavard v0.1.13 h1:oLhMLOFGTLdlda/kma4VOJazblc7IM5y5QPd2A/YjhQ=
github.com/consensys/bavard v0.1.13/go.mod h1:9ItSMtA/dXMAiL7BG6bqW2m3NdSEObYWoH223nGHukI=
github.com/davecgh/go-spew v1.1.1 h1:vj9j/u1bqnvCEfJOwUhtlOARqs3+rkHYY13jYWTU97c=
github.com/davecgh/go-spew v1.1.1/go.mod h1:J7Y8YcW2NihsgmVo/mv3lAwl/skON4iLHjSsI+c5H38=
github.com/kr/pretty v0.1.0 h1:L/CwN0zerZDmRFUapSPitk6f+Q3+0za1rQkzVuMiMFI=
github.com/kr/pretty v0.1.0/go.mod h1:dAy3ld7l9f0ibDNOQOHHMYYIIbhfbHSm3C4ZsoJORNo=
github.com/kr/pty v1.1.1/go.mod h1:pFQYn66WHrOpPYNljwOMqo10TkYh1fy3cYio2l3bCsQ=
github.com/kr/text v0.1.0 h1:45sCR5RtlFHMR4UwH9sdQ5TC8v0qDQCHnXt+kaKSTVE=
github.com/kr/text v0.1.0/go.mod h1:4Jbv+DJW3UT/LiOwJeYQe1efqtUx/iVham/4vfdArNI=
github.com/pmezard/go-difflib v1.0.0 h1:4DBwDE0NGyQoBHbLQYPwSUPoCMWR5BEzIk/f1lZbAQM=
github.com/pmezard/go-difflib v1.0.0/go.mod h1:iKH77koFhYxTK1pcRnkKkqfTogsbg7gZNVY4sRDYZ/4=
github.com/stretchr/testify v1.8.3 h1:RP3t2pwF7cMEbC1dqtB6poj3niw/9gnV4Cjg5oW5gtY=
github.com/stretchr/testify v1.8.3/go.mod h1:sz/lmYIOXD/1dqDmKjjqLyZ2RngseejIcXlSw2iwfAo=
gopkg.in/check.v1 v0.0.0-20161208181325-20d25e280405/go.mod h1:Co6ibVJAznAaIkqp8huTwlJQCZ016jof/cbN4VW5Yz0=
gopkg.in/check.v1 v1.0.0-20180628173108-788fd7840127 h1:qIbj1fsPNlZgppZ+VLlY7N33q108Sa+fhmuc+sWQYwY=
gopkg.in/check.v1 v1.0.0-20180628173108-788fd7840127/go.mod h1:Co6ibVJAznAaIkqp8huTwlJQCZ016jof/cbN4VW5Yz0=
gopkg.in/yaml.v3 v3.0.1 h1:fxVm/GzAzEWqLHuvctI91KS9hhNmmWOoWu0XTYJS7CA=
gopkg.in/yaml.v3 v3.0.1/go.mod h1:K4uyk7z7BCEPqu6E+C64Yfv1cQ7kz7rIZviUmN+EgEM=
rsc.io/tmplfunc v0.0.3 h1:53XFQh69AfOa8Tw0Jm7t+GV7KZhOi6jzsCzTtKbMvzU=
rsc.io/tmplfunc v0.0.3/go.mod h1:AG3sTPzElb1Io3Yg4voV9AGZJuleGAwaVRxL9M49PhA=

34
goicicle/Makefile Normal file
View File

@@ -0,0 +1,34 @@
CUDA_ROOT_DIR = /usr/local/cuda
NVCC = $(CUDA_ROOT_DIR)/bin/nvcc
CFLAGS = -Xcompiler -fPIC -std=c++17
LDFLAGS = -shared
FEATURES = -DG2_DEFINED
TARGET_BN254 = libbn254.so
TARGET_BW6761 = libbw6761.so
TARGET_BLS12_381 = libbls12_381.so
TARGET_BLS12_377 = libbls12_377.so
VPATH = ../icicle/curves/bn254:../icicle/curves/bls12_377:../icicle/curves/bls12_381:../icicle/curves/bw6_761
SRCS_BN254 = lde.cu msm.cu projective.cu ve_mod_mult.cu
SRCS_BW6761 = lde.cu msm.cu projective.cu ve_mod_mult.cu
SRCS_BLS12_381 = lde.cu msm.cu projective.cu ve_mod_mult.cu poseidon.cu
SRCS_BLS12_377 = lde.cu msm.cu projective.cu ve_mod_mult.cu
all: $(TARGET_BN254) $(TARGET_BLS12_381) $(TARGET_BLS12_377) $(TARGET_BW6761)
$(TARGET_BN254):
$(NVCC) $(FEATURES) $(CFLAGS) $(LDFLAGS) $(addprefix ../icicle/curves/bn254/, $(SRCS_BN254)) -o $@
$(TARGET_BW6761):
$(NVCC) $(FEATURES) $(CFLAGS) $(LDFLAGS) $(addprefix ../icicle/curves/bw6_761/, $(SRCS_BW6761)) -o $@
$(TARGET_BLS12_381):
$(NVCC) $(FEATURES) $(CFLAGS) $(LDFLAGS) $(addprefix ../icicle/curves/bls12_381/, $(SRCS_BLS12_381)) -o $@
$(TARGET_BLS12_377):
$(NVCC) $(FEATURES) $(CFLAGS) $(LDFLAGS) $(addprefix ../icicle/curves/bls12_377/, $(SRCS_BLS12_377)) -o $@
clean:
rm -f $(TARGET_BN254) $(TARGET_BLS12_381) $(TARGET_BLS12_377) $(TARGET_BW6761)

82
goicicle/README.md Normal file
View File

@@ -0,0 +1,82 @@
# Golang Bindings
To build the shared library:
To build shared libraries for all supported curves.
```
make all
```
If you wish to build for a specific curve, for example bn254.
```
make libbn254.so
```
The current supported options are `libbn254.so`, `libbls12_381.so`, `libbls12_377.so` and `libbw6_671.so`. The resulting `.so` files are the compiled shared libraries for each curve.
Finally to allow your system to find the shared libraries
```
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH/<path_to_shared_libs>
```
## Running golang tests
To run the tests for curve bn254.
```
go test ./goicicle/curves/bn254 -count=1
```
## Cleaning up
If you want to remove the compiled files
```
make clean
```
This will remove all shared libraries generated from the `make` file.
# How do Golang bindings work?
The shared libraries produced from the CUDA code compilation are used to bind Golang to ICICLE's CUDA code.
1. These shared libraries (`libbn254.so`, `libbls12_381.so`, `libbls12_377.so`, `libbw6_671.so`) can be imported in your Go project to leverage the GPU accelerated functionalities provided by ICICLE.
2. In your Go project, you can use `cgo` to link these shared libraries. Here's a basic example on how you can use `cgo` to link these libraries:
```go
/*
#cgo LDFLAGS: -L/path/to/shared/libs -lbn254 -lbls12_381 -lbls12_377 -lbw6_671
#include "icicle.h" // make sure you use the correct header file(s)
*/
import "C"
func main() {
// Now you can call the C functions from the ICICLE libraries.
// Note that C function calls are prefixed with 'C.' in Go code.
}
```
Replace `/path/to/shared/libs` with the actual path where the shared libraries are located on your system.
# Common issues
### Cannot find shared library
In some cases you may encounter the following error, despite exporting the correct `LD_LIBRARY_PATH`.
```
/usr/local/go/pkg/tool/linux_amd64/link: running gcc failed: exit status 1
/usr/bin/ld: cannot find -lbn254: No such file or directory
/usr/bin/ld: cannot find -lbn254: No such file or directory
/usr/bin/ld: cannot find -lbn254: No such file or directory
/usr/bin/ld: cannot find -lbn254: No such file or directory
/usr/bin/ld: cannot find -lbn254: No such file or directory
collect2: error: ld returned 1 exit status
```
This is normally fixed by exporting the path to the shared library location in the following way: `export CGO_LDFLAGS="-L/<path_to_shared_lib>/"`

View File

@@ -0,0 +1,328 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"unsafe"
"encoding/binary"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_377
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
const SCALAR_SIZE = 8
const BASE_SIZE = 12
type G1ScalarField struct {
S [SCALAR_SIZE]uint32
}
type G1BaseField struct {
S [BASE_SIZE]uint32
}
/*
* BaseField Constructors
*/
func (f *G1BaseField) SetZero() *G1BaseField {
var S [BASE_SIZE]uint32
f.S = S
return f
}
func (f *G1BaseField) SetOne() *G1BaseField {
var S [BASE_SIZE]uint32
S[0] = 1
f.S = S
return f
}
func (p *G1ProjectivePoint) FromAffine(affine *G1PointAffine) *G1ProjectivePoint {
out := (*C.BLS12_377_projective_t)(unsafe.Pointer(p))
in := (*C.BLS12_377_affine_t)(unsafe.Pointer(affine))
C.projective_from_affine_bls12_377(out, in)
return p
}
func (f *G1BaseField) FromLimbs(limbs [BASE_SIZE]uint32) *G1BaseField {
copy(f.S[:], limbs[:])
return f
}
/*
* BaseField methods
*/
func (f *G1BaseField) Limbs() [BASE_SIZE]uint32 {
return f.S
}
func (f *G1BaseField) ToBytesLe() []byte {
bytes := make([]byte, len(f.S)*4)
for i, v := range f.S {
binary.LittleEndian.PutUint32(bytes[i*4:], v)
}
return bytes
}
/*
* ScalarField methods
*/
func (p *G1ScalarField) Random() *G1ScalarField {
outC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(p))
C.random_scalar_bls12_377(outC)
return p
}
func (f *G1ScalarField) SetZero() *G1ScalarField {
var S [SCALAR_SIZE]uint32
f.S = S
return f
}
func (f *G1ScalarField) SetOne() *G1ScalarField {
var S [SCALAR_SIZE]uint32
S[0] = 1
f.S = S
return f
}
func (a *G1ScalarField) Eq(b *G1ScalarField) bool {
for i, v := range a.S {
if b.S[i] != v {
return false
}
}
return true
}
/*
* ScalarField methods
*/
func (f *G1ScalarField) Limbs() [SCALAR_SIZE]uint32 {
return f.S
}
func (f *G1ScalarField) ToBytesLe() []byte {
bytes := make([]byte, len(f.S)*4)
for i, v := range f.S {
binary.LittleEndian.PutUint32(bytes[i*4:], v)
}
return bytes
}
/*
* PointBLS12_377
*/
type G1ProjectivePoint struct {
X, Y, Z G1BaseField
}
func (f *G1ProjectivePoint) SetZero() *G1ProjectivePoint {
var yOne G1BaseField
yOne.SetOne()
var xZero G1BaseField
xZero.SetZero()
var zZero G1BaseField
zZero.SetZero()
f.X = xZero
f.Y = yOne
f.Z = zZero
return f
}
func (p *G1ProjectivePoint) Eq(pCompare *G1ProjectivePoint) bool {
// Cast *PointBLS12_377 to *C.BLS12_377_projective_t
// The unsafe.Pointer cast is necessary because Go doesn't allow direct casts
// between different pointer types.
// It'S your responsibility to ensure that the types are compatible.
pC := (*C.BLS12_377_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BLS12_377_projective_t)(unsafe.Pointer(pCompare))
// Call the C function
// The C function doesn't keep any references to the data,
// so it'S fine if the Go garbage collector moves or deletes the data later.
return bool(C.eq_bls12_377(pC, pCompareC))
}
func (p *G1ProjectivePoint) IsOnCurve() bool {
point := (*C.BLS12_377_projective_t)(unsafe.Pointer(p))
res := C.projective_is_on_curve_bls12_377(point)
return bool(res)
}
func (p *G1ProjectivePoint) Random() *G1ProjectivePoint {
outC := (*C.BLS12_377_projective_t)(unsafe.Pointer(p))
C.random_projective_bls12_377(outC)
return p
}
func (p *G1ProjectivePoint) StripZ() *G1PointAffine {
return &G1PointAffine{
X: p.X,
Y: p.Y,
}
}
func (p *G1ProjectivePoint) FromLimbs(x, y, z *[]uint32) *G1ProjectivePoint {
var _x G1BaseField
var _y G1BaseField
var _z G1BaseField
_x.FromLimbs(GetFixedLimbs(x))
_y.FromLimbs(GetFixedLimbs(y))
_z.FromLimbs(GetFixedLimbs(z))
p.X = _x
p.Y = _y
p.Z = _z
return p
}
/*
* PointAffineNoInfinityBLS12_377
*/
type G1PointAffine struct {
X, Y G1BaseField
}
func (p *G1PointAffine) FromProjective(projective *G1ProjectivePoint) *G1PointAffine {
in := (*C.BLS12_377_projective_t)(unsafe.Pointer(projective))
out := (*C.BLS12_377_affine_t)(unsafe.Pointer(p))
C.projective_to_affine_bls12_377(out, in)
return p
}
func (p *G1PointAffine) ToProjective() *G1ProjectivePoint {
var Z G1BaseField
Z.SetOne()
return &G1ProjectivePoint{
X: p.X,
Y: p.Y,
Z: Z,
}
}
func (p *G1PointAffine) FromLimbs(X, Y *[]uint32) *G1PointAffine {
var _x G1BaseField
var _y G1BaseField
_x.FromLimbs(GetFixedLimbs(X))
_y.FromLimbs(GetFixedLimbs(Y))
p.X = _x
p.Y = _y
return p
}
/*
* Multiplication
*/
func MultiplyVec(a []G1ProjectivePoint, b []G1ScalarField, deviceID int) {
if len(a) != len(b) {
panic("a and b have different lengths")
}
pointsC := (*C.BLS12_377_projective_t)(unsafe.Pointer(&a[0]))
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_point_bls12_377(pointsC, scalarsC, nElementsC, deviceIdC)
}
func MultiplyScalar(a []G1ScalarField, b []G1ScalarField, deviceID int) {
if len(a) != len(b) {
panic("a and b have different lengths")
}
aC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_scalar_bls12_377(aC, bC, nElementsC, deviceIdC)
}
// Multiply a matrix by a scalar:
//
// `a` - flattenned matrix;
// `b` - vector to multiply `a` by;
func MultiplyMatrix(a []G1ScalarField, b []G1ScalarField, deviceID int) {
c := make([]G1ScalarField, len(b))
for i := range c {
var p G1ScalarField
p.SetZero()
c[i] = p
}
aC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&b[0]))
cC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&c[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.matrix_vec_mod_mult_bls12_377(aC, bC, cC, nElementsC, deviceIdC)
}
/*
* Utils
*/
func GetFixedLimbs(slice *[]uint32) [BASE_SIZE]uint32 {
if len(*slice) <= BASE_SIZE {
limbs := [BASE_SIZE]uint32{}
copy(limbs[:len(*slice)], *slice)
return limbs
}
panic("slice has too many elements")
}

View File

@@ -0,0 +1,198 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"encoding/binary"
"testing"
"github.com/stretchr/testify/assert"
)
func TestNewFieldBLS12_377One(t *testing.T) {
var oneField G1BaseField
oneField.SetOne()
rawOneField := [8]uint32([8]uint32{0x1, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0})
assert.Equal(t, oneField.S, rawOneField)
}
func TestNewFieldBLS12_377Zero(t *testing.T) {
var zeroField G1BaseField
zeroField.SetZero()
rawZeroField := [8]uint32([8]uint32{0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0})
assert.Equal(t, zeroField.S, rawZeroField)
}
func TestFieldBLS12_377ToBytesLe(t *testing.T) {
var p G1ProjectivePoint
p.Random()
expected := make([]byte, len(p.X.S)*4) // each uint32 takes 4 bytes
for i, v := range p.X.S {
binary.LittleEndian.PutUint32(expected[i*4:], v)
}
assert.Equal(t, p.X.ToBytesLe(), expected)
assert.Equal(t, len(p.X.ToBytesLe()), 32)
}
func TestNewPointBLS12_377Zero(t *testing.T) {
var pointZero G1ProjectivePoint
pointZero.SetZero()
var baseOne G1BaseField
baseOne.SetOne()
var zeroSanity G1BaseField
zeroSanity.SetZero()
assert.Equal(t, pointZero.X, zeroSanity)
assert.Equal(t, pointZero.Y, baseOne)
assert.Equal(t, pointZero.Z, zeroSanity)
}
func TestFromProjectiveToAffine(t *testing.T) {
var projective G1ProjectivePoint
var affine G1PointAffine
projective.Random()
affine.FromProjective(&projective)
var projective2 G1ProjectivePoint
projective2.FromAffine(&affine)
assert.True(t, projective.IsOnCurve())
assert.True(t, projective2.IsOnCurve())
assert.True(t, projective.Eq(&projective2))
}
func TestBLS12_377Eq(t *testing.T) {
var p1 G1ProjectivePoint
p1.Random()
var p2 G1ProjectivePoint
p2.Random()
assert.Equal(t, p1.Eq(&p1), true)
assert.Equal(t, p1.Eq(&p2), false)
}
func TestBLS12_377StripZ(t *testing.T) {
var p1 G1ProjectivePoint
p1.Random()
p2ZLess := p1.StripZ()
assert.IsType(t, G1PointAffine{}, *p2ZLess)
assert.Equal(t, p1.X, p2ZLess.X)
assert.Equal(t, p1.Y, p2ZLess.Y)
}
func TestPointBLS12_377fromLimbs(t *testing.T) {
var p G1ProjectivePoint
p.Random()
x := p.X.Limbs()
y := p.Y.Limbs()
z := p.Z.Limbs()
xSlice := x[:]
ySlice := y[:]
zSlice := z[:]
var pFromLimbs G1ProjectivePoint
pFromLimbs.FromLimbs(&xSlice, &ySlice, &zSlice)
assert.Equal(t, pFromLimbs, p)
}
func TestNewPointAffineNoInfinityBLS12_377Zero(t *testing.T) {
var zeroP G1PointAffine
var zeroSanity G1BaseField
zeroSanity.SetZero()
assert.Equal(t, zeroP.X, zeroSanity)
assert.Equal(t, zeroP.Y, zeroSanity)
}
func TestPointAffineNoInfinityBLS12_377FromLimbs(t *testing.T) {
// Initialize your test values
x := [12]uint32{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}
y := [12]uint32{9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20}
xSlice := x[:]
ySlice := y[:]
// Execute your function
var result G1PointAffine
result.FromLimbs(&xSlice, &ySlice)
var xBase G1BaseField
var yBase G1BaseField
xBase.FromLimbs(x)
yBase.FromLimbs(y)
// Define your expected result
expected := G1PointAffine{
X: xBase,
Y: yBase,
}
// Test if result is as expected
assert.Equal(t, expected, result)
}
func TestGetFixedLimbs(t *testing.T) {
t.Run("case of valid input of length less than 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7}
expected := [8]uint32{1, 2, 3, 4, 5, 6, 7, 0}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of valid input of length 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7, 8}
expected := [8]uint32{1, 2, 3, 4, 5, 6, 7, 8}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of empty input", func(t *testing.T) {
slice := []uint32{}
expected := [8]uint32{0, 0, 0, 0, 0, 0, 0, 0}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of input length greater than 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7, 8, 9}
defer func() {
if r := recover(); r == nil {
t.Errorf("the code did not panic")
}
}()
GetFixedLimbs(&slice)
})
}

View File

@@ -0,0 +1,102 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"encoding/binary"
"unsafe"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_377
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
// G2 extension field
type G2Element [6]uint64
type ExtentionField struct {
A0, A1 G2Element
}
type G2PointAffine struct {
X, Y ExtentionField
}
type G2Point struct {
X, Y, Z ExtentionField
}
func (p *G2Point) Random() *G2Point {
outC := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(p))
C.random_g2_projective_bls12_377(outC)
return p
}
func (p *G2Point) FromAffine(affine *G2PointAffine) *G2Point {
out := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(p))
in := (*C.BLS12_377_g2_affine_t)(unsafe.Pointer(affine))
C.g2_projective_from_affine_bls12_377(out, in)
return p
}
func (p *G2Point) Eq(pCompare *G2Point) bool {
// Cast *PointBLS12_377 to *C.BLS12_377_projective_t
// The unsafe.Pointer cast is necessary because Go doesn't allow direct casts
// between different pointer types.
// It's your responsibility to ensure that the types are compatible.
pC := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(pCompare))
// Call the C function
// The C function doesn't keep any references to the data,
// so it's fine if the Go garbage collector moves or deletes the data later.
return bool(C.eq_g2_bls12_377(pC, pCompareC))
}
func (f *G2Element) ToBytesLe() []byte {
var bytes []byte
for _, val := range f {
buf := make([]byte, 8) // 8 bytes because uint64 is 64-bit
binary.LittleEndian.PutUint64(buf, val)
bytes = append(bytes, buf...)
}
return bytes
}
func (p *G2PointAffine) FromProjective(projective *G2Point) *G2PointAffine {
out := (*C.BLS12_377_g2_affine_t)(unsafe.Pointer(p))
in := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(projective))
C.g2_projective_to_affine_bls12_377(out, in)
return p
}
func (p *G2Point) IsOnCurve() bool {
// Directly copy memory from the C struct to the Go struct
point := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(p))
res := C.g2_projective_is_on_curve_bls12_377(point)
return bool(res)
}

View File

@@ -0,0 +1,79 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"fmt"
"testing"
"github.com/stretchr/testify/assert"
)
func TestG2Eqg2(t *testing.T) {
var point G2Point
point.Random()
assert.True(t, point.Eq(&point))
}
func TestG2FromProjectiveToAffine(t *testing.T) {
var projective G2Point
projective.Random()
var affine G2PointAffine
affine.FromProjective(&projective)
var projective2 G2Point
projective2.FromAffine(&affine)
assert.True(t, projective.IsOnCurve())
assert.True(t, projective2.IsOnCurve())
assert.True(t, projective.Eq(&projective2))
}
func TestG2Eqg2NotEqual(t *testing.T) {
var point G2Point
point.Random()
var point2 G2Point
point2.Random()
assert.False(t, point.Eq(&point2))
}
func TestG2ToBytes(t *testing.T) {
element := G2Element{0x6546098ea84b6298, 0x4a384533d1f68aca, 0xaa0666972d771336, 0x1569e4a34321993}
bytes := element.ToBytesLe()
assert.Equal(t, bytes, []byte{0x98, 0x62, 0x4b, 0xa8, 0x8e, 0x9, 0x46, 0x65, 0xca, 0x8a, 0xf6, 0xd1, 0x33, 0x45, 0x38, 0x4a, 0x36, 0x13, 0x77, 0x2d, 0x97, 0x66, 0x6, 0xaa, 0x93, 0x19, 0x32, 0x34, 0x4a, 0x9e, 0x56, 0x1})
}
func TestG2ShouldConvertToProjective(t *testing.T) {
fmt.Print() // this prevents the test from hanging. TODO: figure out why
var pointProjective G2Point
pointProjective.Random()
var pointAffine G2PointAffine
pointAffine.FromProjective(&pointProjective)
var proj G2Point
proj.FromAffine(&pointAffine)
assert.True(t, proj.IsOnCurve())
assert.True(t, pointProjective.Eq(&proj))
}

View File

@@ -0,0 +1,98 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdbool.h>
// msm.h
#ifndef _BLS12_377_MSM_H
#define _BLS12_377_MSM_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BLS12_377 projective and affine structs
typedef struct BLS12_377_projective_t BLS12_377_projective_t;
typedef struct BLS12_377_g2_projective_t BLS12_377_g2_projective_t;
typedef struct BLS12_377_affine_t BLS12_377_affine_t;
typedef struct BLS12_377_g2_affine_t BLS12_377_g2_affine_t;
typedef struct BLS12_377_scalar_t BLS12_377_scalar_t;
typedef cudaStream_t CudaStream_t;
int msm_cuda_bls12_377(
BLS12_377_projective_t* out, BLS12_377_affine_t* points, BLS12_377_scalar_t* scalars, size_t count, size_t device_id);
int msm_batch_cuda_bls12_377(
BLS12_377_projective_t* out,
BLS12_377_affine_t* points,
BLS12_377_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_scalar_t* d_scalars,
BLS12_377_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_scalar_t* d_scalars,
BLS12_377_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id);
int msm_g2_cuda_bls12_377(
BLS12_377_g2_projective_t* out,
BLS12_377_g2_affine_t* points,
BLS12_377_scalar_t* scalars,
size_t count,
size_t device_id);
int msm_batch_g2_cuda_bls12_377(
BLS12_377_g2_projective_t* out,
BLS12_377_g2_affine_t* points,
BLS12_377_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_g2_cuda_bls12_377(
BLS12_377_g2_projective_t* d_out,
BLS12_377_scalar_t* d_scalars,
BLS12_377_g2_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_g2_cuda_bls12_377(
BLS12_377_g2_projective_t* d_out,
BLS12_377_scalar_t* d_scalars,
BLS12_377_g2_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id,
cudaStream_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_377_MSM_H */

View File

@@ -0,0 +1,195 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// ntt.h
#ifndef _BLS12_377_NTT_H
#define _BLS12_377_NTT_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BLS12_377 projective and affine structs
typedef struct BLS12_377_projective_t BLS12_377_projective_t;
typedef struct BLS12_377_affine_t BLS12_377_affine_t;
typedef struct BLS12_377_scalar_t BLS12_377_scalar_t;
typedef struct BLS12_377_g2_projective_t BLS12_377_g2_projective_t;
typedef struct BLS12_377_g2_affine_t BLS12_377_g2_affine_t;
int ntt_cuda_bls12_377(BLS12_377_scalar_t* arr, uint32_t n, bool inverse, size_t device_id);
int ntt_batch_cuda_bls12_377(
BLS12_377_scalar_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
int ecntt_cuda_bls12_377(BLS12_377_projective_t* arr, uint32_t n, bool inverse, size_t device_id);
int ecntt_batch_cuda_bls12_377(
BLS12_377_projective_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
BLS12_377_scalar_t*
build_domain_cuda_bls12_377(uint32_t domain_size, uint32_t logn, bool inverse, size_t device_id, size_t stream);
int interpolate_scalars_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
unsigned device_id,
size_t stream);
int interpolate_scalars_batch_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_points_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
size_t device_id,
size_t stream);
int interpolate_points_batch_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_scalars_on_coset_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
BLS12_377_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int interpolate_scalars_batch_on_coset_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_evaluations,
BLS12_377_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
BLS12_377_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_scalars_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned device_id,
size_t stream);
int evaluate_scalars_batch_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int evaluate_points_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
size_t device_id,
size_t stream);
int evaluate_points_batch_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int evaluate_scalars_on_coset_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BLS12_377_scalar_t* coset_powers,
unsigned device_id,
size_t stream);
int evaluate_scalars_on_coset_batch_cuda_bls12_377(
BLS12_377_scalar_t* d_out,
BLS12_377_scalar_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BLS12_377_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BLS12_377_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_batch_cuda_bls12_377(
BLS12_377_projective_t* d_out,
BLS12_377_projective_t* d_coefficients,
BLS12_377_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BLS12_377_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int reverse_order_scalars_cuda_bls12_377(BLS12_377_scalar_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_scalars_batch_cuda_bls12_377(
BLS12_377_scalar_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int reverse_order_points_cuda_bls12_377(BLS12_377_projective_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_points_batch_cuda_bls12_377(
BLS12_377_projective_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int add_scalars_cuda_bls12_377(
BLS12_377_scalar_t* d_out, BLS12_377_scalar_t* d_in1, BLS12_377_scalar_t* d_in2, unsigned n, size_t stream);
int sub_scalars_cuda_bls12_377(
BLS12_377_scalar_t* d_out, BLS12_377_scalar_t* d_in1, BLS12_377_scalar_t* d_in2, unsigned n, size_t stream);
int to_montgomery_scalars_cuda_bls12_377(BLS12_377_scalar_t* d_inout, unsigned n, size_t stream);
int from_montgomery_scalars_cuda_bls12_377(BLS12_377_scalar_t* d_inout, unsigned n, size_t stream);
// points g1
int to_montgomery_proj_points_cuda_bls12_377(BLS12_377_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_cuda_bls12_377(BLS12_377_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_cuda_bls12_377(BLS12_377_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_cuda_bls12_377(BLS12_377_affine_t* d_inout, unsigned n, size_t stream);
// points g2
int to_montgomery_proj_points_g2_cuda_bls12_377(BLS12_377_g2_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_g2_cuda_bls12_377(BLS12_377_g2_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_g2_cuda_bls12_377(BLS12_377_g2_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_g2_cuda_bls12_377(BLS12_377_g2_affine_t* d_inout, unsigned n, size_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_377_NTT_H */

View File

@@ -0,0 +1,50 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// projective.h
#ifdef __cplusplus
extern "C" {
#endif
typedef struct BLS12_377_projective_t BLS12_377_projective_t;
typedef struct BLS12_377_g2_projective_t BLS12_377_g2_projective_t;
typedef struct BLS12_377_affine_t BLS12_377_affine_t;
typedef struct BLS12_377_g2_affine_t BLS12_377_g2_affine_t;
typedef struct BLS12_377_scalar_t BLS12_377_scalar_t;
bool projective_is_on_curve_bls12_377(BLS12_377_projective_t* point1);
int random_scalar_bls12_377(BLS12_377_scalar_t* out);
int random_projective_bls12_377(BLS12_377_projective_t* out);
BLS12_377_projective_t* projective_zero_bls12_377();
int projective_to_affine_bls12_377(BLS12_377_affine_t* out, BLS12_377_projective_t* point1);
int projective_from_affine_bls12_377(BLS12_377_projective_t* out, BLS12_377_affine_t* point1);
int random_g2_projective_bls12_377(BLS12_377_g2_projective_t* out);
int g2_projective_to_affine_bls12_377(BLS12_377_g2_affine_t* out, BLS12_377_g2_projective_t* point1);
int g2_projective_from_affine_bls12_377(BLS12_377_g2_projective_t* out, BLS12_377_g2_affine_t* point1);
bool g2_projective_is_on_curve_bls12_377(BLS12_377_g2_projective_t* point1);
bool eq_bls12_377(BLS12_377_projective_t* point1, BLS12_377_projective_t* point2);
bool eq_g2_bls12_377(BLS12_377_g2_projective_t* point1, BLS12_377_g2_projective_t* point2);
#ifdef __cplusplus
}
#endif

View File

@@ -0,0 +1,49 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// ve_mod_mult.h
#ifndef _BLS12_377_VEC_MULT_H
#define _BLS12_377_VEC_MULT_H
#ifdef __cplusplus
extern "C" {
#endif
typedef struct BLS12_377_projective_t BLS12_377_projective_t;
typedef struct BLS12_377_scalar_t BLS12_377_scalar_t;
int32_t vec_mod_mult_point_bls12_377(
BLS12_377_projective_t* inout, BLS12_377_scalar_t* scalar_vec, size_t n_elments, size_t device_id);
int32_t vec_mod_mult_scalar_bls12_377(
BLS12_377_scalar_t* inout, BLS12_377_scalar_t* scalar_vec, size_t n_elments, size_t device_id);
int32_t vec_mod_mult_device_scalar_bls12_377(
BLS12_377_scalar_t* inout, BLS12_377_scalar_t* scalar_vec, size_t n_elements, size_t device_id);
int32_t matrix_vec_mod_mult_bls12_377(
BLS12_377_scalar_t* matrix_flattened,
BLS12_377_scalar_t* input,
BLS12_377_scalar_t* output,
size_t n_elments,
size_t device_id);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_377_VEC_MULT_H */

View File

@@ -0,0 +1,209 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"errors"
"fmt"
"unsafe"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_377
// #include "msm.h"
import "C"
func Msm(out *G1ProjectivePoint, points []G1PointAffine, scalars []G1ScalarField, device_id int) (*G1ProjectivePoint, error) {
if len(points) != len(scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
pointsC := (*C.BLS12_377_affine_t)(unsafe.Pointer(&points[0]))
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&scalars[0]))
outC := (*C.BLS12_377_projective_t)(unsafe.Pointer(out))
ret := C.msm_cuda_bls12_377(outC, pointsC, scalarsC, C.size_t(len(points)), C.size_t(device_id))
if ret != 0 {
return nil, fmt.Errorf("msm_cuda_bls12_377 returned error code: %d", ret)
}
return out, nil
}
func MsmG2(out *G2Point, points []G2PointAffine, scalars []G1ScalarField, device_id int) (*G2Point, error) {
if len(points) != len(scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
pointsC := (*C.BLS12_377_g2_affine_t)(unsafe.Pointer(&points[0]))
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&scalars[0]))
outC := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(out))
ret := C.msm_g2_cuda_bls12_377(outC, pointsC, scalarsC, C.size_t(len(points)), C.size_t(device_id))
if ret != 0 {
return nil, fmt.Errorf("msm_g2_cuda_bls12_377 returned error code: %d", ret)
}
return out, nil
}
func MsmBatch(points *[]G1PointAffine, scalars *[]G1ScalarField, batchSize, deviceId int) ([]G1ProjectivePoint, error) {
// Check for nil pointers
if points == nil || scalars == nil {
return nil, errors.New("points or scalars is nil")
}
if len(*points) != len(*scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
// Check for empty slices
if len(*points) == 0 || len(*scalars) == 0 {
return nil, errors.New("points or scalars is empty")
}
// Check for zero batchSize
if batchSize <= 0 {
return nil, errors.New("error on: batchSize must be greater than zero")
}
out := make([]G1ProjectivePoint, batchSize)
for i := 0; i < len(out); i++ {
var p G1ProjectivePoint
p.SetZero()
out[i] = p
}
outC := (*C.BLS12_377_projective_t)(unsafe.Pointer(&out[0]))
pointsC := (*C.BLS12_377_affine_t)(unsafe.Pointer(&(*points)[0]))
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
msmSizeC := C.size_t(len(*points) / batchSize)
deviceIdC := C.size_t(deviceId)
batchSizeC := C.size_t(batchSize)
ret := C.msm_batch_cuda_bls12_377(outC, pointsC, scalarsC, batchSizeC, msmSizeC, deviceIdC)
if ret != 0 {
return nil, fmt.Errorf("msm_batch_cuda_bls12_377 returned error code: %d", ret)
}
return out, nil
}
func MsmG2Batch(points *[]G2PointAffine, scalars *[]G1ScalarField, batchSize, deviceId int) ([]G2Point, error) {
// Check for nil pointers
if points == nil || scalars == nil {
return nil, errors.New("points or scalars is nil")
}
if len(*points) != len(*scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
// Check for empty slices
if len(*points) == 0 || len(*scalars) == 0 {
return nil, errors.New("points or scalars is empty")
}
// Check for zero batchSize
if batchSize <= 0 {
return nil, errors.New("error on: batchSize must be greater than zero")
}
out := make([]G2Point, batchSize)
outC := (*C.BLS12_377_g2_projective_t)(unsafe.Pointer(&out[0]))
pointsC := (*C.BLS12_377_g2_affine_t)(unsafe.Pointer(&(*points)[0]))
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
msmSizeC := C.size_t(len(*points) / batchSize)
deviceIdC := C.size_t(deviceId)
batchSizeC := C.size_t(batchSize)
ret := C.msm_batch_g2_cuda_bls12_377(outC, pointsC, scalarsC, batchSizeC, msmSizeC, deviceIdC)
if ret != 0 {
return nil, fmt.Errorf("msm_batch_cuda_bls12_377 returned error code: %d", ret)
}
return out, nil
}
func Commit(d_out, d_scalars, d_points unsafe.Pointer, count, bucketFactor int) int {
d_outC := (*C.BLS12_377_projective_t)(d_out)
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
pointsC := (*C.BLS12_377_affine_t)(d_points)
countC := (C.size_t)(count)
largeBucketFactorC := C.uint(bucketFactor)
ret := C.commit_cuda_bls12_377(d_outC, scalarsC, pointsC, countC, largeBucketFactorC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitG2(d_out, d_scalars, d_points unsafe.Pointer, count, bucketFactor int) int {
d_outC := (*C.BLS12_377_g2_projective_t)(d_out)
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
pointsC := (*C.BLS12_377_g2_affine_t)(d_points)
countC := (C.size_t)(count)
largeBucketFactorC := C.uint(bucketFactor)
ret := C.commit_g2_cuda_bls12_377(d_outC, scalarsC, pointsC, countC, largeBucketFactorC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitBatch(d_out, d_scalars, d_points unsafe.Pointer, count, batch_size int) int {
d_outC := (*C.BLS12_377_projective_t)(d_out)
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
pointsC := (*C.BLS12_377_affine_t)(d_points)
countC := (C.size_t)(count)
batch_sizeC := (C.size_t)(batch_size)
ret := C.commit_batch_cuda_bls12_377(d_outC, scalarsC, pointsC, countC, batch_sizeC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitG2Batch(d_out, d_scalars, d_points unsafe.Pointer, count, batch_size int) int {
d_outC := (*C.BLS12_377_g2_projective_t)(d_out)
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
pointsC := (*C.BLS12_377_g2_affine_t)(d_points)
countC := (C.size_t)(count)
batch_sizeC := (C.size_t)(batch_size)
ret := C.msm_batch_g2_cuda_bls12_377(d_outC, pointsC, scalarsC, countC, batch_sizeC, 0)
if ret != 0 {
return -1
}
return 0
}

View File

@@ -0,0 +1,360 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"fmt"
"math"
"testing"
"time"
"unsafe"
"github.com/ingonyama-zk/icicle/goicicle"
"github.com/stretchr/testify/assert"
)
func GeneratePoints(count int) []G1PointAffine {
// Declare a slice of integers
var points []G1PointAffine
// populate the slice
for i := 0; i < 10; i++ {
var pointProjective G1ProjectivePoint
pointProjective.Random()
var pointAffine G1PointAffine
pointAffine.FromProjective(&pointProjective)
points = append(points, pointAffine)
}
log2_10 := math.Log2(10)
log2Count := math.Log2(float64(count))
log2Size := int(math.Ceil(log2Count - log2_10))
for i := 0; i < log2Size; i++ {
points = append(points, points...)
}
return points[:count]
}
func GeneratePointsProj(count int) []G1ProjectivePoint {
// Declare a slice of integers
var points []G1ProjectivePoint
// Use a loop to populate the slice
for i := 0; i < count; i++ {
var p G1ProjectivePoint
p.Random()
points = append(points, p)
}
return points
}
func GenerateScalars(count int, skewed bool) []G1ScalarField {
// Declare a slice of integers
var scalars []G1ScalarField
var rand G1ScalarField
var zero G1ScalarField
var one G1ScalarField
var randLarge G1ScalarField
zero.SetZero()
one.SetOne()
randLarge.Random()
if skewed && count > 1_200_000 {
for i := 0; i < count-1_200_000; i++ {
rand.Random()
scalars = append(scalars, rand)
}
for i := 0; i < 600_000; i++ {
scalars = append(scalars, randLarge)
}
for i := 0; i < 400_000; i++ {
scalars = append(scalars, zero)
}
for i := 0; i < 200_000; i++ {
scalars = append(scalars, one)
}
} else {
for i := 0; i < count; i++ {
rand.Random()
scalars = append(scalars, rand)
}
}
return scalars[:count]
}
func TestMSM(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GeneratePoints(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out := new(G1ProjectivePoint)
startTime := time.Now()
_, e := Msm(out, points, scalars, 0) // non mont
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
assert.Equal(t, e, nil, "error should be nil")
assert.True(t, out.IsOnCurve())
}
}
func TestCommitMSM(t *testing.T) {
for _, v := range []int{8} {
count := 1<<v - 1
points := GeneratePoints(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out_d, _ := goicicle.CudaMalloc(96)
pointsBytes := count * 64
points_d, _ := goicicle.CudaMalloc(pointsBytes)
goicicle.CudaMemCpyHtoD[G1PointAffine](points_d, points, pointsBytes)
scalarBytes := count * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
startTime := time.Now()
e := Commit(out_d, scalars_d, points_d, count, 10)
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
outHost := make([]G1ProjectivePoint, 1)
goicicle.CudaMemCpyDtoH[G1ProjectivePoint](outHost, out_d, 96)
assert.Equal(t, e, 0, "error should be 0")
assert.True(t, outHost[0].IsOnCurve())
}
}
func BenchmarkCommit(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GeneratePoints(msmSize)
scalars := GenerateScalars(msmSize, false)
out_d, _ := goicicle.CudaMalloc(96)
pointsBytes := msmSize * 64
points_d, _ := goicicle.CudaMalloc(pointsBytes)
goicicle.CudaMemCpyHtoD[G1PointAffine](points_d, points, pointsBytes)
scalarBytes := msmSize * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
b.Run(fmt.Sprintf("MSM %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
e := Commit(out_d, scalars_d, points_d, msmSize, 10)
if e != 0 {
panic("Error occurred")
}
}
})
}
}
func TestBatchMSM(t *testing.T) {
for _, batchPow2 := range []int{2, 4} {
for _, pow2 := range []int{4, 6} {
msmSize := 1 << pow2
batchSize := 1 << batchPow2
count := msmSize * batchSize
points := GeneratePoints(count)
scalars := GenerateScalars(count, false)
pointsResults, e := MsmBatch(&points, &scalars, batchSize, 0)
if e != nil {
t.Errorf("MsmBatchBLS12_377 returned an error: %v", e)
}
if len(pointsResults) != batchSize {
t.Errorf("Expected length %d, but got %d", batchSize, len(pointsResults))
}
for _, s := range pointsResults {
assert.True(t, s.IsOnCurve())
}
}
}
}
func BenchmarkMSM(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GeneratePoints(msmSize)
scalars := GenerateScalars(msmSize, false)
b.Run(fmt.Sprintf("MSM %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
out := new(G1ProjectivePoint)
_, e := Msm(out, points, scalars, 0)
if e != nil {
panic("Error occurred")
}
}
})
}
}
// G2
func GenerateG2Points(count int) []G2PointAffine {
// Declare a slice of integers
var points []G2PointAffine
// populate the slice
for i := 0; i < 10; i++ {
fmt.Print() // this prevents the test from hanging. TODO: figure out why
var p G2Point
p.Random()
var affine G2PointAffine
affine.FromProjective(&p)
points = append(points, affine)
}
log2_10 := math.Log2(10)
log2Count := math.Log2(float64(count))
log2Size := int(math.Ceil(log2Count - log2_10))
for i := 0; i < log2Size; i++ {
points = append(points, points...)
}
return points[:count]
}
func TestMsmG2BLS12_377(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GenerateG2Points(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out := new(G2Point)
_, e := MsmG2(out, points, scalars, 0)
assert.Equal(t, e, nil, "error should be nil")
assert.True(t, out.IsOnCurve())
}
}
func BenchmarkMsmG2BLS12_377(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GenerateG2Points(msmSize)
scalars := GenerateScalars(msmSize, false)
b.Run(fmt.Sprintf("MSM G2 %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
out := new(G2Point)
_, e := MsmG2(out, points, scalars, 0)
if e != nil {
panic("Error occurred")
}
}
})
}
}
func TestCommitG2MSM(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GenerateG2Points(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
var sizeCheckG2PointAffine G2PointAffine
inputPointsBytes := count * int(unsafe.Sizeof(sizeCheckG2PointAffine))
var sizeCheckG2Point G2Point
out_d, _ := goicicle.CudaMalloc(int(unsafe.Sizeof(sizeCheckG2Point)))
points_d, _ := goicicle.CudaMalloc(inputPointsBytes)
goicicle.CudaMemCpyHtoD[G2PointAffine](points_d, points, inputPointsBytes)
scalarBytes := count * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
startTime := time.Now()
e := CommitG2(out_d, scalars_d, points_d, count, 10)
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
outHost := make([]G2Point, 1)
goicicle.CudaMemCpyDtoH[G2Point](outHost, out_d, int(unsafe.Sizeof(sizeCheckG2Point)))
assert.Equal(t, e, 0, "error should be 0")
assert.Equal(t, len(outHost), 1)
result := outHost[0]
assert.True(t, result.IsOnCurve())
}
}
func TestBatchG2MSM(t *testing.T) {
for _, batchPow2 := range []int{2, 4} {
for _, pow2 := range []int{4, 6} {
msmSize := 1 << pow2
batchSize := 1 << batchPow2
count := msmSize * batchSize
points := GenerateG2Points(count)
scalars := GenerateScalars(count, false)
pointsResults, e := MsmG2Batch(&points, &scalars, batchSize, 0)
if e != nil {
t.Errorf("MsmBatchBLS12_377 returned an error: %v", e)
}
if len(pointsResults) != batchSize {
t.Errorf("Expected length %d, but got %d", batchSize, len(pointsResults))
}
for _, s := range pointsResults {
assert.True(t, s.IsOnCurve())
}
}
}
}

View File

@@ -0,0 +1,222 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"errors"
"fmt"
"unsafe"
"github.com/ingonyama-zk/icicle/goicicle"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_377
// #include "ntt.h"
import "C"
const (
NONE = 0
DIF = 1
DIT = 2
)
func Ntt(scalars *[]G1ScalarField, isInverse bool, deviceId int) uint64 {
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
ret := C.ntt_cuda_bls12_377(scalarsC, C.uint32_t(len(*scalars)), C.bool(isInverse), C.size_t(deviceId))
return uint64(ret)
}
func NttBatch(scalars *[]G1ScalarField, isInverse bool, batchSize, deviceId int) uint64 {
scalarsC := (*C.BLS12_377_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
isInverseC := C.bool(isInverse)
batchSizeC := C.uint32_t(batchSize)
deviceIdC := C.size_t(deviceId)
ret := C.ntt_batch_cuda_bls12_377(scalarsC, C.uint32_t(len(*scalars)), batchSizeC, isInverseC, deviceIdC)
return uint64(ret)
}
func EcNtt(values *[]G1ProjectivePoint, isInverse bool, deviceId int) uint64 {
valuesC := (*C.BLS12_377_projective_t)(unsafe.Pointer(&(*values)[0]))
deviceIdC := C.size_t(deviceId)
isInverseC := C.bool(isInverse)
n := C.uint32_t(len(*values))
ret := C.ecntt_cuda_bls12_377(valuesC, n, isInverseC, deviceIdC)
return uint64(ret)
}
func EcNttBatch(values *[]G1ProjectivePoint, isInverse bool, batchSize, deviceId int) uint64 {
valuesC := (*C.BLS12_377_projective_t)(unsafe.Pointer(&(*values)[0]))
deviceIdC := C.size_t(deviceId)
isInverseC := C.bool(isInverse)
n := C.uint32_t(len(*values))
batchSizeC := C.uint32_t(batchSize)
ret := C.ecntt_batch_cuda_bls12_377(valuesC, n, batchSizeC, isInverseC, deviceIdC)
return uint64(ret)
}
func GenerateTwiddles(d_size int, log_d_size int, inverse bool) (up unsafe.Pointer, err error) {
domain_size := C.uint32_t(d_size)
logn := C.uint32_t(log_d_size)
is_inverse := C.bool(inverse)
dp := C.build_domain_cuda_bls12_377(domain_size, logn, is_inverse, 0, 0)
if dp == nil {
err = errors.New("nullptr returned from generating twiddles")
return unsafe.Pointer(nil), err
}
return unsafe.Pointer(dp), nil
}
// Reverses d_scalars in-place
func ReverseScalars(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
lenC := C.int(len)
if success := C.reverse_order_scalars_cuda_bls12_377(scalarsC, lenC, 0, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func Interpolate(scalars, twiddles, cosetPowers unsafe.Pointer, size int, isCoset bool) unsafe.Pointer {
size_d := size * 32
dp, err := goicicle.CudaMalloc(size_d)
if err != nil {
return nil
}
d_out := (*C.BLS12_377_scalar_t)(dp)
scalarsC := (*C.BLS12_377_scalar_t)(scalars)
twiddlesC := (*C.BLS12_377_scalar_t)(twiddles)
cosetPowersC := (*C.BLS12_377_scalar_t)(cosetPowers)
sizeC := C.uint(size)
var ret C.int
if isCoset {
ret = C.interpolate_scalars_on_coset_cuda_bls12_377(d_out, scalarsC, twiddlesC, sizeC, cosetPowersC, 0, 0)
} else {
ret = C.interpolate_scalars_cuda_bls12_377(d_out, scalarsC, twiddlesC, sizeC, 0, 0)
}
if ret != 0 {
fmt.Print("error interpolating")
}
return unsafe.Pointer(d_out)
}
func Evaluate(scalars_out, scalars, twiddles, coset_powers unsafe.Pointer, scalars_size, twiddles_size int, isCoset bool) int {
scalars_outC := (*C.BLS12_377_scalar_t)(scalars_out)
scalarsC := (*C.BLS12_377_scalar_t)(scalars)
twiddlesC := (*C.BLS12_377_scalar_t)(twiddles)
coset_powersC := (*C.BLS12_377_scalar_t)(coset_powers)
sizeC := C.uint(scalars_size)
twiddlesC_size := C.uint(twiddles_size)
var ret C.int
if isCoset {
ret = C.evaluate_scalars_on_coset_cuda_bls12_377(scalars_outC, scalarsC, twiddlesC, twiddlesC_size, sizeC, coset_powersC, 0, 0)
} else {
ret = C.evaluate_scalars_cuda_bls12_377(scalars_outC, scalarsC, twiddlesC, twiddlesC_size, sizeC, 0, 0)
}
if ret != 0 {
fmt.Print("error interpolating")
return -1
}
return 0
}
func VecScalarAdd(in1_d, in2_d unsafe.Pointer, size int) int {
in1_dC := (*C.BLS12_377_scalar_t)(in1_d)
in2_dC := (*C.BLS12_377_scalar_t)(in2_d)
sizeC := C.uint(size)
ret := C.add_scalars_cuda_bls12_377(in1_dC, in1_dC, in2_dC, sizeC, 0)
if ret != 0 {
fmt.Print("error adding scalar vectors")
return -1
}
return 0
}
func VecScalarSub(in1_d, in2_d unsafe.Pointer, size int) int {
in1_dC := (*C.BLS12_377_scalar_t)(in1_d)
in2_dC := (*C.BLS12_377_scalar_t)(in2_d)
sizeC := C.uint(size)
ret := C.sub_scalars_cuda_bls12_377(in1_dC, in1_dC, in2_dC, sizeC, 0)
if ret != 0 {
fmt.Print("error subtracting scalar vectors")
return -1
}
return 0
}
func ToMontgomery(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
lenC := C.uint(len)
if success := C.to_montgomery_scalars_cuda_bls12_377(scalarsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func FromMontgomery(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_377_scalar_t)(d_scalars)
lenC := C.uint(len)
if success := C.from_montgomery_scalars_cuda_bls12_377(scalarsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func AffinePointFromMontgomery(d_points unsafe.Pointer, len int) (int, error) {
pointsC := (*C.BLS12_377_affine_t)(d_points)
lenC := C.uint(len)
if success := C.from_montgomery_aff_points_cuda_bls12_377(pointsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func G2AffinePointFromMontgomery(d_points unsafe.Pointer, len int) (int, error) {
pointsC := (*C.BLS12_377_g2_affine_t)(d_points)
lenC := C.uint(len)
if success := C.from_montgomery_aff_points_g2_cuda_bls12_377(pointsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}

View File

@@ -0,0 +1,148 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
import (
"fmt"
"github.com/stretchr/testify/assert"
"reflect"
"testing"
)
func TestNttBLS12_377Batch(t *testing.T) {
count := 1 << 20
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
NttBatch(&nttResult, false, count, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestNttBLS12_377CompareToGnarkDIF(t *testing.T) {
count := 1 << 2
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, false, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestINttBLS12_377CompareToGnarkDIT(t *testing.T) {
count := 1 << 3
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, true, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestNttBLS12_377(t *testing.T) {
count := 1 << 3
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, false, 0)
assert.NotEqual(t, nttResult, scalars)
inttResult := make([]G1ScalarField, len(nttResult))
copy(inttResult, nttResult)
assert.Equal(t, inttResult, nttResult)
Ntt(&inttResult, true, 0)
assert.Equal(t, inttResult, scalars)
}
func TestNttBatchBLS12_377(t *testing.T) {
count := 1 << 5
batches := 4
scalars := GenerateScalars(count*batches, false)
var scalarVecOfVec [][]G1ScalarField = make([][]G1ScalarField, 0)
for i := 0; i < batches; i++ {
start := i * count
end := (i + 1) * count
batch := make([]G1ScalarField, len(scalars[start:end]))
copy(batch, scalars[start:end])
scalarVecOfVec = append(scalarVecOfVec, batch)
}
nttBatchResult := make([]G1ScalarField, len(scalars))
copy(nttBatchResult, scalars)
NttBatch(&nttBatchResult, false, count, 0)
var nttResultVecOfVec [][]G1ScalarField
for i := 0; i < batches; i++ {
// Clone the slice
clone := make([]G1ScalarField, len(scalarVecOfVec[i]))
copy(clone, scalarVecOfVec[i])
// Add it to the result vector of vectors
nttResultVecOfVec = append(nttResultVecOfVec, clone)
// Call the ntt_bls12_377 function
Ntt(&nttResultVecOfVec[i], false, 0)
}
assert.NotEqual(t, nttBatchResult, scalars)
// Check that the ntt of each vec of scalars is equal to the intt of the specific batch
for i := 0; i < batches; i++ {
if !reflect.DeepEqual(nttResultVecOfVec[i], nttBatchResult[i*count:((i+1)*count)]) {
t.Errorf("ntt of vec of scalars not equal to intt of specific batch")
}
}
}
func BenchmarkNTT(b *testing.B) {
LOG_NTT_SIZES := []int{12, 15, 20, 21, 22, 23, 24, 25, 26}
for _, logNTTSize := range LOG_NTT_SIZES {
nttSize := 1 << logNTTSize
b.Run(fmt.Sprintf("NTT %d", logNTTSize), func(b *testing.B) {
scalars := GenerateScalars(nttSize, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
for n := 0; n < b.N; n++ {
Ntt(&nttResult, false, 0)
}
})
}
}

View File

@@ -0,0 +1,38 @@
package bls12377
import "encoding/binary"
// Function to convert [8]uint32 to [4]uint64
func ConvertUint32ArrToUint64Arr(arr32 [8]uint32) [4]uint64 {
var arr64 [4]uint64
for i := 0; i < len(arr32); i += 2 {
arr64[i/2] = (uint64(arr32[i]) << 32) | uint64(arr32[i+1])
}
return arr64
}
func ConvertUint64ArrToUint32Arr4(arr64 [4]uint64) [8]uint32 {
var arr32 [8]uint32
for i, v := range arr64 {
b := make([]byte, 8)
binary.LittleEndian.PutUint64(b, v)
arr32[i*2] = binary.LittleEndian.Uint32(b[0:4])
arr32[i*2+1] = binary.LittleEndian.Uint32(b[4:8])
}
return arr32
}
func ConvertUint64ArrToUint32Arr6(arr64 [6]uint64) [12]uint32 {
var arr32 [12]uint32
for i, v := range arr64 {
b := make([]byte, 8)
binary.LittleEndian.PutUint64(b, v)
arr32[i*2] = binary.LittleEndian.Uint32(b[0:4])
arr32[i*2+1] = binary.LittleEndian.Uint32(b[4:8])
}
return arr32
}

View File

@@ -0,0 +1,42 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12377
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_377
// #include "ve_mod_mult.h"
import "C"
import (
"fmt"
"unsafe"
)
func VecScalarMulMod(scalarVec1, scalarVec2 unsafe.Pointer, size int) int {
scalarVec1C := (*C.BLS12_377_scalar_t)(scalarVec1)
scalarVec2C := (*C.BLS12_377_scalar_t)(scalarVec2)
sizeC := C.size_t(size)
ret := C.vec_mod_mult_device_scalar_bls12_377(scalarVec1C, scalarVec2C, sizeC, 0)
if ret != 0 {
fmt.Print("error multiplying scalar vectors")
return -1
}
return 0
}

View File

@@ -0,0 +1,328 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"unsafe"
"encoding/binary"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_381
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
const SCALAR_SIZE = 8
const BASE_SIZE = 12
type G1ScalarField struct {
S [SCALAR_SIZE]uint32
}
type G1BaseField struct {
S [BASE_SIZE]uint32
}
/*
* BaseField Constructors
*/
func (f *G1BaseField) SetZero() *G1BaseField {
var S [BASE_SIZE]uint32
f.S = S
return f
}
func (f *G1BaseField) SetOne() *G1BaseField {
var S [BASE_SIZE]uint32
S[0] = 1
f.S = S
return f
}
func (p *G1ProjectivePoint) FromAffine(affine *G1PointAffine) *G1ProjectivePoint {
out := (*C.BLS12_381_projective_t)(unsafe.Pointer(p))
in := (*C.BLS12_381_affine_t)(unsafe.Pointer(affine))
C.projective_from_affine_bls12_381(out, in)
return p
}
func (f *G1BaseField) FromLimbs(limbs [BASE_SIZE]uint32) *G1BaseField {
copy(f.S[:], limbs[:])
return f
}
/*
* BaseField methods
*/
func (f *G1BaseField) Limbs() [BASE_SIZE]uint32 {
return f.S
}
func (f *G1BaseField) ToBytesLe() []byte {
bytes := make([]byte, len(f.S)*4)
for i, v := range f.S {
binary.LittleEndian.PutUint32(bytes[i*4:], v)
}
return bytes
}
/*
* ScalarField methods
*/
func (p *G1ScalarField) Random() *G1ScalarField {
outC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(p))
C.random_scalar_bls12_381(outC)
return p
}
func (f *G1ScalarField) SetZero() *G1ScalarField {
var S [SCALAR_SIZE]uint32
f.S = S
return f
}
func (f *G1ScalarField) SetOne() *G1ScalarField {
var S [SCALAR_SIZE]uint32
S[0] = 1
f.S = S
return f
}
func (a *G1ScalarField) Eq(b *G1ScalarField) bool {
for i, v := range a.S {
if b.S[i] != v {
return false
}
}
return true
}
/*
* ScalarField methods
*/
func (f *G1ScalarField) Limbs() [SCALAR_SIZE]uint32 {
return f.S
}
func (f *G1ScalarField) ToBytesLe() []byte {
bytes := make([]byte, len(f.S)*4)
for i, v := range f.S {
binary.LittleEndian.PutUint32(bytes[i*4:], v)
}
return bytes
}
/*
* PointBLS12_381
*/
type G1ProjectivePoint struct {
X, Y, Z G1BaseField
}
func (f *G1ProjectivePoint) SetZero() *G1ProjectivePoint {
var yOne G1BaseField
yOne.SetOne()
var xZero G1BaseField
xZero.SetZero()
var zZero G1BaseField
zZero.SetZero()
f.X = xZero
f.Y = yOne
f.Z = zZero
return f
}
func (p *G1ProjectivePoint) Eq(pCompare *G1ProjectivePoint) bool {
// Cast *PointBLS12_381 to *C.BLS12_381_projective_t
// The unsafe.Pointer cast is necessary because Go doesn't allow direct casts
// between different pointer types.
// It'S your responsibility to ensure that the types are compatible.
pC := (*C.BLS12_381_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BLS12_381_projective_t)(unsafe.Pointer(pCompare))
// Call the C function
// The C function doesn't keep any references to the data,
// so it'S fine if the Go garbage collector moves or deletes the data later.
return bool(C.eq_bls12_381(pC, pCompareC))
}
func (p *G1ProjectivePoint) IsOnCurve() bool {
point := (*C.BLS12_381_projective_t)(unsafe.Pointer(p))
res := C.projective_is_on_curve_bls12_381(point)
return bool(res)
}
func (p *G1ProjectivePoint) Random() *G1ProjectivePoint {
outC := (*C.BLS12_381_projective_t)(unsafe.Pointer(p))
C.random_projective_bls12_381(outC)
return p
}
func (p *G1ProjectivePoint) StripZ() *G1PointAffine {
return &G1PointAffine{
X: p.X,
Y: p.Y,
}
}
func (p *G1ProjectivePoint) FromLimbs(x, y, z *[]uint32) *G1ProjectivePoint {
var _x G1BaseField
var _y G1BaseField
var _z G1BaseField
_x.FromLimbs(GetFixedLimbs(x))
_y.FromLimbs(GetFixedLimbs(y))
_z.FromLimbs(GetFixedLimbs(z))
p.X = _x
p.Y = _y
p.Z = _z
return p
}
/*
* PointAffineNoInfinityBLS12_381
*/
type G1PointAffine struct {
X, Y G1BaseField
}
func (p *G1PointAffine) FromProjective(projective *G1ProjectivePoint) *G1PointAffine {
in := (*C.BLS12_381_projective_t)(unsafe.Pointer(projective))
out := (*C.BLS12_381_affine_t)(unsafe.Pointer(p))
C.projective_to_affine_bls12_381(out, in)
return p
}
func (p *G1PointAffine) ToProjective() *G1ProjectivePoint {
var Z G1BaseField
Z.SetOne()
return &G1ProjectivePoint{
X: p.X,
Y: p.Y,
Z: Z,
}
}
func (p *G1PointAffine) FromLimbs(X, Y *[]uint32) *G1PointAffine {
var _x G1BaseField
var _y G1BaseField
_x.FromLimbs(GetFixedLimbs(X))
_y.FromLimbs(GetFixedLimbs(Y))
p.X = _x
p.Y = _y
return p
}
/*
* Multiplication
*/
func MultiplyVec(a []G1ProjectivePoint, b []G1ScalarField, deviceID int) {
if len(a) != len(b) {
panic("a and b have different lengths")
}
pointsC := (*C.BLS12_381_projective_t)(unsafe.Pointer(&a[0]))
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_point_bls12_381(pointsC, scalarsC, nElementsC, deviceIdC)
}
func MultiplyScalar(a []G1ScalarField, b []G1ScalarField, deviceID int) {
if len(a) != len(b) {
panic("a and b have different lengths")
}
aC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_scalar_bls12_381(aC, bC, nElementsC, deviceIdC)
}
// Multiply a matrix by a scalar:
//
// `a` - flattenned matrix;
// `b` - vector to multiply `a` by;
func MultiplyMatrix(a []G1ScalarField, b []G1ScalarField, deviceID int) {
c := make([]G1ScalarField, len(b))
for i := range c {
var p G1ScalarField
p.SetZero()
c[i] = p
}
aC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&b[0]))
cC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&c[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.matrix_vec_mod_mult_bls12_381(aC, bC, cC, nElementsC, deviceIdC)
}
/*
* Utils
*/
func GetFixedLimbs(slice *[]uint32) [BASE_SIZE]uint32 {
if len(*slice) <= BASE_SIZE {
limbs := [BASE_SIZE]uint32{}
copy(limbs[:len(*slice)], *slice)
return limbs
}
panic("slice has too many elements")
}

View File

@@ -0,0 +1,198 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"encoding/binary"
"testing"
"github.com/stretchr/testify/assert"
)
func TestNewFieldBLS12_381One(t *testing.T) {
var oneField G1BaseField
oneField.SetOne()
rawOneField := [8]uint32([8]uint32{0x1, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0})
assert.Equal(t, oneField.S, rawOneField)
}
func TestNewFieldBLS12_381Zero(t *testing.T) {
var zeroField G1BaseField
zeroField.SetZero()
rawZeroField := [8]uint32([8]uint32{0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0, 0x0})
assert.Equal(t, zeroField.S, rawZeroField)
}
func TestFieldBLS12_381ToBytesLe(t *testing.T) {
var p G1ProjectivePoint
p.Random()
expected := make([]byte, len(p.X.S)*4) // each uint32 takes 4 bytes
for i, v := range p.X.S {
binary.LittleEndian.PutUint32(expected[i*4:], v)
}
assert.Equal(t, p.X.ToBytesLe(), expected)
assert.Equal(t, len(p.X.ToBytesLe()), 32)
}
func TestNewPointBLS12_381Zero(t *testing.T) {
var pointZero G1ProjectivePoint
pointZero.SetZero()
var baseOne G1BaseField
baseOne.SetOne()
var zeroSanity G1BaseField
zeroSanity.SetZero()
assert.Equal(t, pointZero.X, zeroSanity)
assert.Equal(t, pointZero.Y, baseOne)
assert.Equal(t, pointZero.Z, zeroSanity)
}
func TestFromProjectiveToAffine(t *testing.T) {
var projective G1ProjectivePoint
var affine G1PointAffine
projective.Random()
affine.FromProjective(&projective)
var projective2 G1ProjectivePoint
projective2.FromAffine(&affine)
assert.True(t, projective.IsOnCurve())
assert.True(t, projective2.IsOnCurve())
assert.True(t, projective.Eq(&projective2))
}
func TestBLS12_381Eq(t *testing.T) {
var p1 G1ProjectivePoint
p1.Random()
var p2 G1ProjectivePoint
p2.Random()
assert.Equal(t, p1.Eq(&p1), true)
assert.Equal(t, p1.Eq(&p2), false)
}
func TestBLS12_381StripZ(t *testing.T) {
var p1 G1ProjectivePoint
p1.Random()
p2ZLess := p1.StripZ()
assert.IsType(t, G1PointAffine{}, *p2ZLess)
assert.Equal(t, p1.X, p2ZLess.X)
assert.Equal(t, p1.Y, p2ZLess.Y)
}
func TestPointBLS12_381fromLimbs(t *testing.T) {
var p G1ProjectivePoint
p.Random()
x := p.X.Limbs()
y := p.Y.Limbs()
z := p.Z.Limbs()
xSlice := x[:]
ySlice := y[:]
zSlice := z[:]
var pFromLimbs G1ProjectivePoint
pFromLimbs.FromLimbs(&xSlice, &ySlice, &zSlice)
assert.Equal(t, pFromLimbs, p)
}
func TestNewPointAffineNoInfinityBLS12_381Zero(t *testing.T) {
var zeroP G1PointAffine
var zeroSanity G1BaseField
zeroSanity.SetZero()
assert.Equal(t, zeroP.X, zeroSanity)
assert.Equal(t, zeroP.Y, zeroSanity)
}
func TestPointAffineNoInfinityBLS12_381FromLimbs(t *testing.T) {
// Initialize your test values
x := [12]uint32{1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12}
y := [12]uint32{9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20}
xSlice := x[:]
ySlice := y[:]
// Execute your function
var result G1PointAffine
result.FromLimbs(&xSlice, &ySlice)
var xBase G1BaseField
var yBase G1BaseField
xBase.FromLimbs(x)
yBase.FromLimbs(y)
// Define your expected result
expected := G1PointAffine{
X: xBase,
Y: yBase,
}
// Test if result is as expected
assert.Equal(t, expected, result)
}
func TestGetFixedLimbs(t *testing.T) {
t.Run("case of valid input of length less than 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7}
expected := [8]uint32{1, 2, 3, 4, 5, 6, 7, 0}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of valid input of length 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7, 8}
expected := [8]uint32{1, 2, 3, 4, 5, 6, 7, 8}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of empty input", func(t *testing.T) {
slice := []uint32{}
expected := [8]uint32{0, 0, 0, 0, 0, 0, 0, 0}
result := GetFixedLimbs(&slice)
assert.Equal(t, result, expected)
})
t.Run("case of input length greater than 8", func(t *testing.T) {
slice := []uint32{1, 2, 3, 4, 5, 6, 7, 8, 9}
defer func() {
if r := recover(); r == nil {
t.Errorf("the code did not panic")
}
}()
GetFixedLimbs(&slice)
})
}

View File

@@ -0,0 +1,102 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"encoding/binary"
"unsafe"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_381
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
// G2 extension field
type G2Element [6]uint64
type ExtentionField struct {
A0, A1 G2Element
}
type G2PointAffine struct {
X, Y ExtentionField
}
type G2Point struct {
X, Y, Z ExtentionField
}
func (p *G2Point) Random() *G2Point {
outC := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(p))
C.random_g2_projective_bls12_381(outC)
return p
}
func (p *G2Point) FromAffine(affine *G2PointAffine) *G2Point {
out := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(p))
in := (*C.BLS12_381_g2_affine_t)(unsafe.Pointer(affine))
C.g2_projective_from_affine_bls12_381(out, in)
return p
}
func (p *G2Point) Eq(pCompare *G2Point) bool {
// Cast *PointBLS12_381 to *C.BLS12_381_projective_t
// The unsafe.Pointer cast is necessary because Go doesn't allow direct casts
// between different pointer types.
// It's your responsibility to ensure that the types are compatible.
pC := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(pCompare))
// Call the C function
// The C function doesn't keep any references to the data,
// so it's fine if the Go garbage collector moves or deletes the data later.
return bool(C.eq_g2_bls12_381(pC, pCompareC))
}
func (f *G2Element) ToBytesLe() []byte {
var bytes []byte
for _, val := range f {
buf := make([]byte, 8) // 8 bytes because uint64 is 64-bit
binary.LittleEndian.PutUint64(buf, val)
bytes = append(bytes, buf...)
}
return bytes
}
func (p *G2PointAffine) FromProjective(projective *G2Point) *G2PointAffine {
out := (*C.BLS12_381_g2_affine_t)(unsafe.Pointer(p))
in := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(projective))
C.g2_projective_to_affine_bls12_381(out, in)
return p
}
func (p *G2Point) IsOnCurve() bool {
// Directly copy memory from the C struct to the Go struct
point := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(p))
res := C.g2_projective_is_on_curve_bls12_381(point)
return bool(res)
}

View File

@@ -0,0 +1,79 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"fmt"
"testing"
"github.com/stretchr/testify/assert"
)
func TestG2Eqg2(t *testing.T) {
var point G2Point
point.Random()
assert.True(t, point.Eq(&point))
}
func TestG2FromProjectiveToAffine(t *testing.T) {
var projective G2Point
projective.Random()
var affine G2PointAffine
affine.FromProjective(&projective)
var projective2 G2Point
projective2.FromAffine(&affine)
assert.True(t, projective.IsOnCurve())
assert.True(t, projective2.IsOnCurve())
assert.True(t, projective.Eq(&projective2))
}
func TestG2Eqg2NotEqual(t *testing.T) {
var point G2Point
point.Random()
var point2 G2Point
point2.Random()
assert.False(t, point.Eq(&point2))
}
func TestG2ToBytes(t *testing.T) {
element := G2Element{0x6546098ea84b6298, 0x4a384533d1f68aca, 0xaa0666972d771336, 0x1569e4a34321993}
bytes := element.ToBytesLe()
assert.Equal(t, bytes, []byte{0x98, 0x62, 0x4b, 0xa8, 0x8e, 0x9, 0x46, 0x65, 0xca, 0x8a, 0xf6, 0xd1, 0x33, 0x45, 0x38, 0x4a, 0x36, 0x13, 0x77, 0x2d, 0x97, 0x66, 0x6, 0xaa, 0x93, 0x19, 0x32, 0x34, 0x4a, 0x9e, 0x56, 0x1})
}
func TestG2ShouldConvertToProjective(t *testing.T) {
fmt.Print() // this prevents the test from hanging. TODO: figure out why
var pointProjective G2Point
pointProjective.Random()
var pointAffine G2PointAffine
pointAffine.FromProjective(&pointProjective)
var proj G2Point
proj.FromAffine(&pointAffine)
assert.True(t, proj.IsOnCurve())
assert.True(t, pointProjective.Eq(&proj))
}

View File

@@ -0,0 +1,98 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdbool.h>
// msm.h
#ifndef _BLS12_381_MSM_H
#define _BLS12_381_MSM_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BLS12_381 projective and affine structs
typedef struct BLS12_381_projective_t BLS12_381_projective_t;
typedef struct BLS12_381_g2_projective_t BLS12_381_g2_projective_t;
typedef struct BLS12_381_affine_t BLS12_381_affine_t;
typedef struct BLS12_381_g2_affine_t BLS12_381_g2_affine_t;
typedef struct BLS12_381_scalar_t BLS12_381_scalar_t;
typedef cudaStream_t CudaStream_t;
int msm_cuda_bls12_381(
BLS12_381_projective_t* out, BLS12_381_affine_t* points, BLS12_381_scalar_t* scalars, size_t count, size_t device_id);
int msm_batch_cuda_bls12_381(
BLS12_381_projective_t* out,
BLS12_381_affine_t* points,
BLS12_381_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_scalar_t* d_scalars,
BLS12_381_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_scalar_t* d_scalars,
BLS12_381_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id);
int msm_g2_cuda_bls12_381(
BLS12_381_g2_projective_t* out,
BLS12_381_g2_affine_t* points,
BLS12_381_scalar_t* scalars,
size_t count,
size_t device_id);
int msm_batch_g2_cuda_bls12_381(
BLS12_381_g2_projective_t* out,
BLS12_381_g2_affine_t* points,
BLS12_381_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_g2_cuda_bls12_381(
BLS12_381_g2_projective_t* d_out,
BLS12_381_scalar_t* d_scalars,
BLS12_381_g2_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_g2_cuda_bls12_381(
BLS12_381_g2_projective_t* d_out,
BLS12_381_scalar_t* d_scalars,
BLS12_381_g2_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id,
cudaStream_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_381_MSM_H */

View File

@@ -0,0 +1,195 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// ntt.h
#ifndef _BLS12_381_NTT_H
#define _BLS12_381_NTT_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BLS12_381 projective and affine structs
typedef struct BLS12_381_projective_t BLS12_381_projective_t;
typedef struct BLS12_381_affine_t BLS12_381_affine_t;
typedef struct BLS12_381_scalar_t BLS12_381_scalar_t;
typedef struct BLS12_381_g2_projective_t BLS12_381_g2_projective_t;
typedef struct BLS12_381_g2_affine_t BLS12_381_g2_affine_t;
int ntt_cuda_bls12_381(BLS12_381_scalar_t* arr, uint32_t n, bool inverse, size_t device_id);
int ntt_batch_cuda_bls12_381(
BLS12_381_scalar_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
int ecntt_cuda_bls12_381(BLS12_381_projective_t* arr, uint32_t n, bool inverse, size_t device_id);
int ecntt_batch_cuda_bls12_381(
BLS12_381_projective_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
BLS12_381_scalar_t*
build_domain_cuda_bls12_381(uint32_t domain_size, uint32_t logn, bool inverse, size_t device_id, size_t stream);
int interpolate_scalars_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
unsigned device_id,
size_t stream);
int interpolate_scalars_batch_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_points_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
size_t device_id,
size_t stream);
int interpolate_points_batch_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_scalars_on_coset_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
BLS12_381_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int interpolate_scalars_batch_on_coset_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_evaluations,
BLS12_381_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
BLS12_381_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_scalars_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned device_id,
size_t stream);
int evaluate_scalars_batch_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int evaluate_points_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
size_t device_id,
size_t stream);
int evaluate_points_batch_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int evaluate_scalars_on_coset_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BLS12_381_scalar_t* coset_powers,
unsigned device_id,
size_t stream);
int evaluate_scalars_on_coset_batch_cuda_bls12_381(
BLS12_381_scalar_t* d_out,
BLS12_381_scalar_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BLS12_381_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BLS12_381_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_batch_cuda_bls12_381(
BLS12_381_projective_t* d_out,
BLS12_381_projective_t* d_coefficients,
BLS12_381_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BLS12_381_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int reverse_order_scalars_cuda_bls12_381(BLS12_381_scalar_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_scalars_batch_cuda_bls12_381(
BLS12_381_scalar_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int reverse_order_points_cuda_bls12_381(BLS12_381_projective_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_points_batch_cuda_bls12_381(
BLS12_381_projective_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int add_scalars_cuda_bls12_381(
BLS12_381_scalar_t* d_out, BLS12_381_scalar_t* d_in1, BLS12_381_scalar_t* d_in2, unsigned n, size_t stream);
int sub_scalars_cuda_bls12_381(
BLS12_381_scalar_t* d_out, BLS12_381_scalar_t* d_in1, BLS12_381_scalar_t* d_in2, unsigned n, size_t stream);
int to_montgomery_scalars_cuda_bls12_381(BLS12_381_scalar_t* d_inout, unsigned n, size_t stream);
int from_montgomery_scalars_cuda_bls12_381(BLS12_381_scalar_t* d_inout, unsigned n, size_t stream);
// points g1
int to_montgomery_proj_points_cuda_bls12_381(BLS12_381_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_cuda_bls12_381(BLS12_381_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_cuda_bls12_381(BLS12_381_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_cuda_bls12_381(BLS12_381_affine_t* d_inout, unsigned n, size_t stream);
// points g2
int to_montgomery_proj_points_g2_cuda_bls12_381(BLS12_381_g2_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_g2_cuda_bls12_381(BLS12_381_g2_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_g2_cuda_bls12_381(BLS12_381_g2_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_g2_cuda_bls12_381(BLS12_381_g2_affine_t* d_inout, unsigned n, size_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_381_NTT_H */

View File

@@ -0,0 +1,50 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// projective.h
#ifdef __cplusplus
extern "C" {
#endif
typedef struct BLS12_381_projective_t BLS12_381_projective_t;
typedef struct BLS12_381_g2_projective_t BLS12_381_g2_projective_t;
typedef struct BLS12_381_affine_t BLS12_381_affine_t;
typedef struct BLS12_381_g2_affine_t BLS12_381_g2_affine_t;
typedef struct BLS12_381_scalar_t BLS12_381_scalar_t;
bool projective_is_on_curve_bls12_381(BLS12_381_projective_t* point1);
int random_scalar_bls12_381(BLS12_381_scalar_t* out);
int random_projective_bls12_381(BLS12_381_projective_t* out);
BLS12_381_projective_t* projective_zero_bls12_381();
int projective_to_affine_bls12_381(BLS12_381_affine_t* out, BLS12_381_projective_t* point1);
int projective_from_affine_bls12_381(BLS12_381_projective_t* out, BLS12_381_affine_t* point1);
int random_g2_projective_bls12_381(BLS12_381_g2_projective_t* out);
int g2_projective_to_affine_bls12_381(BLS12_381_g2_affine_t* out, BLS12_381_g2_projective_t* point1);
int g2_projective_from_affine_bls12_381(BLS12_381_g2_projective_t* out, BLS12_381_g2_affine_t* point1);
bool g2_projective_is_on_curve_bls12_381(BLS12_381_g2_projective_t* point1);
bool eq_bls12_381(BLS12_381_projective_t* point1, BLS12_381_projective_t* point2);
bool eq_g2_bls12_381(BLS12_381_g2_projective_t* point1, BLS12_381_g2_projective_t* point2);
#ifdef __cplusplus
}
#endif

View File

@@ -0,0 +1,49 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
#include <cuda.h>
#include <stdbool.h>
// ve_mod_mult.h
#ifndef _BLS12_381_VEC_MULT_H
#define _BLS12_381_VEC_MULT_H
#ifdef __cplusplus
extern "C" {
#endif
typedef struct BLS12_381_projective_t BLS12_381_projective_t;
typedef struct BLS12_381_scalar_t BLS12_381_scalar_t;
int32_t vec_mod_mult_point_bls12_381(
BLS12_381_projective_t* inout, BLS12_381_scalar_t* scalar_vec, size_t n_elments, size_t device_id);
int32_t vec_mod_mult_scalar_bls12_381(
BLS12_381_scalar_t* inout, BLS12_381_scalar_t* scalar_vec, size_t n_elments, size_t device_id);
int32_t vec_mod_mult_device_scalar_bls12_381(
BLS12_381_scalar_t* inout, BLS12_381_scalar_t* scalar_vec, size_t n_elements, size_t device_id);
int32_t matrix_vec_mod_mult_bls12_381(
BLS12_381_scalar_t* matrix_flattened,
BLS12_381_scalar_t* input,
BLS12_381_scalar_t* output,
size_t n_elments,
size_t device_id);
#ifdef __cplusplus
}
#endif
#endif /* _BLS12_381_VEC_MULT_H */

View File

@@ -0,0 +1,209 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"errors"
"fmt"
"unsafe"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_381
// #include "msm.h"
import "C"
func Msm(out *G1ProjectivePoint, points []G1PointAffine, scalars []G1ScalarField, device_id int) (*G1ProjectivePoint, error) {
if len(points) != len(scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
pointsC := (*C.BLS12_381_affine_t)(unsafe.Pointer(&points[0]))
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&scalars[0]))
outC := (*C.BLS12_381_projective_t)(unsafe.Pointer(out))
ret := C.msm_cuda_bls12_381(outC, pointsC, scalarsC, C.size_t(len(points)), C.size_t(device_id))
if ret != 0 {
return nil, fmt.Errorf("msm_cuda_bls12_381 returned error code: %d", ret)
}
return out, nil
}
func MsmG2(out *G2Point, points []G2PointAffine, scalars []G1ScalarField, device_id int) (*G2Point, error) {
if len(points) != len(scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
pointsC := (*C.BLS12_381_g2_affine_t)(unsafe.Pointer(&points[0]))
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&scalars[0]))
outC := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(out))
ret := C.msm_g2_cuda_bls12_381(outC, pointsC, scalarsC, C.size_t(len(points)), C.size_t(device_id))
if ret != 0 {
return nil, fmt.Errorf("msm_g2_cuda_bls12_381 returned error code: %d", ret)
}
return out, nil
}
func MsmBatch(points *[]G1PointAffine, scalars *[]G1ScalarField, batchSize, deviceId int) ([]G1ProjectivePoint, error) {
// Check for nil pointers
if points == nil || scalars == nil {
return nil, errors.New("points or scalars is nil")
}
if len(*points) != len(*scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
// Check for empty slices
if len(*points) == 0 || len(*scalars) == 0 {
return nil, errors.New("points or scalars is empty")
}
// Check for zero batchSize
if batchSize <= 0 {
return nil, errors.New("error on: batchSize must be greater than zero")
}
out := make([]G1ProjectivePoint, batchSize)
for i := 0; i < len(out); i++ {
var p G1ProjectivePoint
p.SetZero()
out[i] = p
}
outC := (*C.BLS12_381_projective_t)(unsafe.Pointer(&out[0]))
pointsC := (*C.BLS12_381_affine_t)(unsafe.Pointer(&(*points)[0]))
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
msmSizeC := C.size_t(len(*points) / batchSize)
deviceIdC := C.size_t(deviceId)
batchSizeC := C.size_t(batchSize)
ret := C.msm_batch_cuda_bls12_381(outC, pointsC, scalarsC, batchSizeC, msmSizeC, deviceIdC)
if ret != 0 {
return nil, fmt.Errorf("msm_batch_cuda_bls12_381 returned error code: %d", ret)
}
return out, nil
}
func MsmG2Batch(points *[]G2PointAffine, scalars *[]G1ScalarField, batchSize, deviceId int) ([]G2Point, error) {
// Check for nil pointers
if points == nil || scalars == nil {
return nil, errors.New("points or scalars is nil")
}
if len(*points) != len(*scalars) {
return nil, errors.New("error on: len(points) != len(scalars)")
}
// Check for empty slices
if len(*points) == 0 || len(*scalars) == 0 {
return nil, errors.New("points or scalars is empty")
}
// Check for zero batchSize
if batchSize <= 0 {
return nil, errors.New("error on: batchSize must be greater than zero")
}
out := make([]G2Point, batchSize)
outC := (*C.BLS12_381_g2_projective_t)(unsafe.Pointer(&out[0]))
pointsC := (*C.BLS12_381_g2_affine_t)(unsafe.Pointer(&(*points)[0]))
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
msmSizeC := C.size_t(len(*points) / batchSize)
deviceIdC := C.size_t(deviceId)
batchSizeC := C.size_t(batchSize)
ret := C.msm_batch_g2_cuda_bls12_381(outC, pointsC, scalarsC, batchSizeC, msmSizeC, deviceIdC)
if ret != 0 {
return nil, fmt.Errorf("msm_batch_cuda_bls12_381 returned error code: %d", ret)
}
return out, nil
}
func Commit(d_out, d_scalars, d_points unsafe.Pointer, count, bucketFactor int) int {
d_outC := (*C.BLS12_381_projective_t)(d_out)
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
pointsC := (*C.BLS12_381_affine_t)(d_points)
countC := (C.size_t)(count)
largeBucketFactorC := C.uint(bucketFactor)
ret := C.commit_cuda_bls12_381(d_outC, scalarsC, pointsC, countC, largeBucketFactorC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitG2(d_out, d_scalars, d_points unsafe.Pointer, count, bucketFactor int) int {
d_outC := (*C.BLS12_381_g2_projective_t)(d_out)
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
pointsC := (*C.BLS12_381_g2_affine_t)(d_points)
countC := (C.size_t)(count)
largeBucketFactorC := C.uint(bucketFactor)
ret := C.commit_g2_cuda_bls12_381(d_outC, scalarsC, pointsC, countC, largeBucketFactorC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitBatch(d_out, d_scalars, d_points unsafe.Pointer, count, batch_size int) int {
d_outC := (*C.BLS12_381_projective_t)(d_out)
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
pointsC := (*C.BLS12_381_affine_t)(d_points)
countC := (C.size_t)(count)
batch_sizeC := (C.size_t)(batch_size)
ret := C.commit_batch_cuda_bls12_381(d_outC, scalarsC, pointsC, countC, batch_sizeC, 0)
if ret != 0 {
return -1
}
return 0
}
func CommitG2Batch(d_out, d_scalars, d_points unsafe.Pointer, count, batch_size int) int {
d_outC := (*C.BLS12_381_g2_projective_t)(d_out)
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
pointsC := (*C.BLS12_381_g2_affine_t)(d_points)
countC := (C.size_t)(count)
batch_sizeC := (C.size_t)(batch_size)
ret := C.msm_batch_g2_cuda_bls12_381(d_outC, pointsC, scalarsC, countC, batch_sizeC, 0)
if ret != 0 {
return -1
}
return 0
}

View File

@@ -0,0 +1,360 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"fmt"
"math"
"testing"
"time"
"unsafe"
"github.com/ingonyama-zk/icicle/goicicle"
"github.com/stretchr/testify/assert"
)
func GeneratePoints(count int) []G1PointAffine {
// Declare a slice of integers
var points []G1PointAffine
// populate the slice
for i := 0; i < 10; i++ {
var pointProjective G1ProjectivePoint
pointProjective.Random()
var pointAffine G1PointAffine
pointAffine.FromProjective(&pointProjective)
points = append(points, pointAffine)
}
log2_10 := math.Log2(10)
log2Count := math.Log2(float64(count))
log2Size := int(math.Ceil(log2Count - log2_10))
for i := 0; i < log2Size; i++ {
points = append(points, points...)
}
return points[:count]
}
func GeneratePointsProj(count int) []G1ProjectivePoint {
// Declare a slice of integers
var points []G1ProjectivePoint
// Use a loop to populate the slice
for i := 0; i < count; i++ {
var p G1ProjectivePoint
p.Random()
points = append(points, p)
}
return points
}
func GenerateScalars(count int, skewed bool) []G1ScalarField {
// Declare a slice of integers
var scalars []G1ScalarField
var rand G1ScalarField
var zero G1ScalarField
var one G1ScalarField
var randLarge G1ScalarField
zero.SetZero()
one.SetOne()
randLarge.Random()
if skewed && count > 1_200_000 {
for i := 0; i < count-1_200_000; i++ {
rand.Random()
scalars = append(scalars, rand)
}
for i := 0; i < 600_000; i++ {
scalars = append(scalars, randLarge)
}
for i := 0; i < 400_000; i++ {
scalars = append(scalars, zero)
}
for i := 0; i < 200_000; i++ {
scalars = append(scalars, one)
}
} else {
for i := 0; i < count; i++ {
rand.Random()
scalars = append(scalars, rand)
}
}
return scalars[:count]
}
func TestMSM(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GeneratePoints(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out := new(G1ProjectivePoint)
startTime := time.Now()
_, e := Msm(out, points, scalars, 0) // non mont
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
assert.Equal(t, e, nil, "error should be nil")
assert.True(t, out.IsOnCurve())
}
}
func TestCommitMSM(t *testing.T) {
for _, v := range []int{8} {
count := 1<<v - 1
points := GeneratePoints(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out_d, _ := goicicle.CudaMalloc(96)
pointsBytes := count * 64
points_d, _ := goicicle.CudaMalloc(pointsBytes)
goicicle.CudaMemCpyHtoD[G1PointAffine](points_d, points, pointsBytes)
scalarBytes := count * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
startTime := time.Now()
e := Commit(out_d, scalars_d, points_d, count, 10)
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
outHost := make([]G1ProjectivePoint, 1)
goicicle.CudaMemCpyDtoH[G1ProjectivePoint](outHost, out_d, 96)
assert.Equal(t, e, 0, "error should be 0")
assert.True(t, outHost[0].IsOnCurve())
}
}
func BenchmarkCommit(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GeneratePoints(msmSize)
scalars := GenerateScalars(msmSize, false)
out_d, _ := goicicle.CudaMalloc(96)
pointsBytes := msmSize * 64
points_d, _ := goicicle.CudaMalloc(pointsBytes)
goicicle.CudaMemCpyHtoD[G1PointAffine](points_d, points, pointsBytes)
scalarBytes := msmSize * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
b.Run(fmt.Sprintf("MSM %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
e := Commit(out_d, scalars_d, points_d, msmSize, 10)
if e != 0 {
panic("Error occurred")
}
}
})
}
}
func TestBatchMSM(t *testing.T) {
for _, batchPow2 := range []int{2, 4} {
for _, pow2 := range []int{4, 6} {
msmSize := 1 << pow2
batchSize := 1 << batchPow2
count := msmSize * batchSize
points := GeneratePoints(count)
scalars := GenerateScalars(count, false)
pointsResults, e := MsmBatch(&points, &scalars, batchSize, 0)
if e != nil {
t.Errorf("MsmBatchBLS12_381 returned an error: %v", e)
}
if len(pointsResults) != batchSize {
t.Errorf("Expected length %d, but got %d", batchSize, len(pointsResults))
}
for _, s := range pointsResults {
assert.True(t, s.IsOnCurve())
}
}
}
}
func BenchmarkMSM(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GeneratePoints(msmSize)
scalars := GenerateScalars(msmSize, false)
b.Run(fmt.Sprintf("MSM %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
out := new(G1ProjectivePoint)
_, e := Msm(out, points, scalars, 0)
if e != nil {
panic("Error occurred")
}
}
})
}
}
// G2
func GenerateG2Points(count int) []G2PointAffine {
// Declare a slice of integers
var points []G2PointAffine
// populate the slice
for i := 0; i < 10; i++ {
fmt.Print() // this prevents the test from hanging. TODO: figure out why
var p G2Point
p.Random()
var affine G2PointAffine
affine.FromProjective(&p)
points = append(points, affine)
}
log2_10 := math.Log2(10)
log2Count := math.Log2(float64(count))
log2Size := int(math.Ceil(log2Count - log2_10))
for i := 0; i < log2Size; i++ {
points = append(points, points...)
}
return points[:count]
}
func TestMsmG2BLS12_381(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GenerateG2Points(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
out := new(G2Point)
_, e := MsmG2(out, points, scalars, 0)
assert.Equal(t, e, nil, "error should be nil")
assert.True(t, out.IsOnCurve())
}
}
func BenchmarkMsmG2BLS12_381(b *testing.B) {
LOG_MSM_SIZES := []int{20, 21, 22, 23, 24, 25, 26}
for _, logMsmSize := range LOG_MSM_SIZES {
msmSize := 1 << logMsmSize
points := GenerateG2Points(msmSize)
scalars := GenerateScalars(msmSize, false)
b.Run(fmt.Sprintf("MSM G2 %d", logMsmSize), func(b *testing.B) {
for n := 0; n < b.N; n++ {
out := new(G2Point)
_, e := MsmG2(out, points, scalars, 0)
if e != nil {
panic("Error occurred")
}
}
})
}
}
func TestCommitG2MSM(t *testing.T) {
for _, v := range []int{8} {
count := 1 << v
points := GenerateG2Points(count)
fmt.Print("Finished generating points\n")
scalars := GenerateScalars(count, false)
fmt.Print("Finished generating scalars\n")
var sizeCheckG2PointAffine G2PointAffine
inputPointsBytes := count * int(unsafe.Sizeof(sizeCheckG2PointAffine))
var sizeCheckG2Point G2Point
out_d, _ := goicicle.CudaMalloc(int(unsafe.Sizeof(sizeCheckG2Point)))
points_d, _ := goicicle.CudaMalloc(inputPointsBytes)
goicicle.CudaMemCpyHtoD[G2PointAffine](points_d, points, inputPointsBytes)
scalarBytes := count * 32
scalars_d, _ := goicicle.CudaMalloc(scalarBytes)
goicicle.CudaMemCpyHtoD[G1ScalarField](scalars_d, scalars, scalarBytes)
startTime := time.Now()
e := CommitG2(out_d, scalars_d, points_d, count, 10)
fmt.Printf("icicle MSM took: %d ms\n", time.Since(startTime).Milliseconds())
outHost := make([]G2Point, 1)
goicicle.CudaMemCpyDtoH[G2Point](outHost, out_d, int(unsafe.Sizeof(sizeCheckG2Point)))
assert.Equal(t, e, 0, "error should be 0")
assert.Equal(t, len(outHost), 1)
result := outHost[0]
assert.True(t, result.IsOnCurve())
}
}
func TestBatchG2MSM(t *testing.T) {
for _, batchPow2 := range []int{2, 4} {
for _, pow2 := range []int{4, 6} {
msmSize := 1 << pow2
batchSize := 1 << batchPow2
count := msmSize * batchSize
points := GenerateG2Points(count)
scalars := GenerateScalars(count, false)
pointsResults, e := MsmG2Batch(&points, &scalars, batchSize, 0)
if e != nil {
t.Errorf("MsmBatchBLS12_381 returned an error: %v", e)
}
if len(pointsResults) != batchSize {
t.Errorf("Expected length %d, but got %d", batchSize, len(pointsResults))
}
for _, s := range pointsResults {
assert.True(t, s.IsOnCurve())
}
}
}
}

View File

@@ -0,0 +1,222 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"errors"
"fmt"
"unsafe"
"github.com/ingonyama-zk/icicle/goicicle"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_381
// #include "ntt.h"
import "C"
const (
NONE = 0
DIF = 1
DIT = 2
)
func Ntt(scalars *[]G1ScalarField, isInverse bool, deviceId int) uint64 {
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
ret := C.ntt_cuda_bls12_381(scalarsC, C.uint32_t(len(*scalars)), C.bool(isInverse), C.size_t(deviceId))
return uint64(ret)
}
func NttBatch(scalars *[]G1ScalarField, isInverse bool, batchSize, deviceId int) uint64 {
scalarsC := (*C.BLS12_381_scalar_t)(unsafe.Pointer(&(*scalars)[0]))
isInverseC := C.bool(isInverse)
batchSizeC := C.uint32_t(batchSize)
deviceIdC := C.size_t(deviceId)
ret := C.ntt_batch_cuda_bls12_381(scalarsC, C.uint32_t(len(*scalars)), batchSizeC, isInverseC, deviceIdC)
return uint64(ret)
}
func EcNtt(values *[]G1ProjectivePoint, isInverse bool, deviceId int) uint64 {
valuesC := (*C.BLS12_381_projective_t)(unsafe.Pointer(&(*values)[0]))
deviceIdC := C.size_t(deviceId)
isInverseC := C.bool(isInverse)
n := C.uint32_t(len(*values))
ret := C.ecntt_cuda_bls12_381(valuesC, n, isInverseC, deviceIdC)
return uint64(ret)
}
func EcNttBatch(values *[]G1ProjectivePoint, isInverse bool, batchSize, deviceId int) uint64 {
valuesC := (*C.BLS12_381_projective_t)(unsafe.Pointer(&(*values)[0]))
deviceIdC := C.size_t(deviceId)
isInverseC := C.bool(isInverse)
n := C.uint32_t(len(*values))
batchSizeC := C.uint32_t(batchSize)
ret := C.ecntt_batch_cuda_bls12_381(valuesC, n, batchSizeC, isInverseC, deviceIdC)
return uint64(ret)
}
func GenerateTwiddles(d_size int, log_d_size int, inverse bool) (up unsafe.Pointer, err error) {
domain_size := C.uint32_t(d_size)
logn := C.uint32_t(log_d_size)
is_inverse := C.bool(inverse)
dp := C.build_domain_cuda_bls12_381(domain_size, logn, is_inverse, 0, 0)
if dp == nil {
err = errors.New("nullptr returned from generating twiddles")
return unsafe.Pointer(nil), err
}
return unsafe.Pointer(dp), nil
}
// Reverses d_scalars in-place
func ReverseScalars(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
lenC := C.int(len)
if success := C.reverse_order_scalars_cuda_bls12_381(scalarsC, lenC, 0, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func Interpolate(scalars, twiddles, cosetPowers unsafe.Pointer, size int, isCoset bool) unsafe.Pointer {
size_d := size * 32
dp, err := goicicle.CudaMalloc(size_d)
if err != nil {
return nil
}
d_out := (*C.BLS12_381_scalar_t)(dp)
scalarsC := (*C.BLS12_381_scalar_t)(scalars)
twiddlesC := (*C.BLS12_381_scalar_t)(twiddles)
cosetPowersC := (*C.BLS12_381_scalar_t)(cosetPowers)
sizeC := C.uint(size)
var ret C.int
if isCoset {
ret = C.interpolate_scalars_on_coset_cuda_bls12_381(d_out, scalarsC, twiddlesC, sizeC, cosetPowersC, 0, 0)
} else {
ret = C.interpolate_scalars_cuda_bls12_381(d_out, scalarsC, twiddlesC, sizeC, 0, 0)
}
if ret != 0 {
fmt.Print("error interpolating")
}
return unsafe.Pointer(d_out)
}
func Evaluate(scalars_out, scalars, twiddles, coset_powers unsafe.Pointer, scalars_size, twiddles_size int, isCoset bool) int {
scalars_outC := (*C.BLS12_381_scalar_t)(scalars_out)
scalarsC := (*C.BLS12_381_scalar_t)(scalars)
twiddlesC := (*C.BLS12_381_scalar_t)(twiddles)
coset_powersC := (*C.BLS12_381_scalar_t)(coset_powers)
sizeC := C.uint(scalars_size)
twiddlesC_size := C.uint(twiddles_size)
var ret C.int
if isCoset {
ret = C.evaluate_scalars_on_coset_cuda_bls12_381(scalars_outC, scalarsC, twiddlesC, twiddlesC_size, sizeC, coset_powersC, 0, 0)
} else {
ret = C.evaluate_scalars_cuda_bls12_381(scalars_outC, scalarsC, twiddlesC, twiddlesC_size, sizeC, 0, 0)
}
if ret != 0 {
fmt.Print("error interpolating")
return -1
}
return 0
}
func VecScalarAdd(in1_d, in2_d unsafe.Pointer, size int) int {
in1_dC := (*C.BLS12_381_scalar_t)(in1_d)
in2_dC := (*C.BLS12_381_scalar_t)(in2_d)
sizeC := C.uint(size)
ret := C.add_scalars_cuda_bls12_381(in1_dC, in1_dC, in2_dC, sizeC, 0)
if ret != 0 {
fmt.Print("error adding scalar vectors")
return -1
}
return 0
}
func VecScalarSub(in1_d, in2_d unsafe.Pointer, size int) int {
in1_dC := (*C.BLS12_381_scalar_t)(in1_d)
in2_dC := (*C.BLS12_381_scalar_t)(in2_d)
sizeC := C.uint(size)
ret := C.sub_scalars_cuda_bls12_381(in1_dC, in1_dC, in2_dC, sizeC, 0)
if ret != 0 {
fmt.Print("error subtracting scalar vectors")
return -1
}
return 0
}
func ToMontgomery(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
lenC := C.uint(len)
if success := C.to_montgomery_scalars_cuda_bls12_381(scalarsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func FromMontgomery(d_scalars unsafe.Pointer, len int) (int, error) {
scalarsC := (*C.BLS12_381_scalar_t)(d_scalars)
lenC := C.uint(len)
if success := C.from_montgomery_scalars_cuda_bls12_381(scalarsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func AffinePointFromMontgomery(d_points unsafe.Pointer, len int) (int, error) {
pointsC := (*C.BLS12_381_affine_t)(d_points)
lenC := C.uint(len)
if success := C.from_montgomery_aff_points_cuda_bls12_381(pointsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}
func G2AffinePointFromMontgomery(d_points unsafe.Pointer, len int) (int, error) {
pointsC := (*C.BLS12_381_g2_affine_t)(d_points)
lenC := C.uint(len)
if success := C.from_montgomery_aff_points_g2_cuda_bls12_381(pointsC, lenC, 0); success != 0 {
return -1, errors.New("reversing failed")
}
return 0, nil
}

View File

@@ -0,0 +1,148 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
import (
"fmt"
"github.com/stretchr/testify/assert"
"reflect"
"testing"
)
func TestNttBLS12_381Batch(t *testing.T) {
count := 1 << 20
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
NttBatch(&nttResult, false, count, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestNttBLS12_381CompareToGnarkDIF(t *testing.T) {
count := 1 << 2
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, false, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestINttBLS12_381CompareToGnarkDIT(t *testing.T) {
count := 1 << 3
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, true, 0)
assert.NotEqual(t, nttResult, scalars)
assert.Equal(t, nttResult, nttResult)
}
func TestNttBLS12_381(t *testing.T) {
count := 1 << 3
scalars := GenerateScalars(count, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
assert.Equal(t, nttResult, scalars)
Ntt(&nttResult, false, 0)
assert.NotEqual(t, nttResult, scalars)
inttResult := make([]G1ScalarField, len(nttResult))
copy(inttResult, nttResult)
assert.Equal(t, inttResult, nttResult)
Ntt(&inttResult, true, 0)
assert.Equal(t, inttResult, scalars)
}
func TestNttBatchBLS12_381(t *testing.T) {
count := 1 << 5
batches := 4
scalars := GenerateScalars(count*batches, false)
var scalarVecOfVec [][]G1ScalarField = make([][]G1ScalarField, 0)
for i := 0; i < batches; i++ {
start := i * count
end := (i + 1) * count
batch := make([]G1ScalarField, len(scalars[start:end]))
copy(batch, scalars[start:end])
scalarVecOfVec = append(scalarVecOfVec, batch)
}
nttBatchResult := make([]G1ScalarField, len(scalars))
copy(nttBatchResult, scalars)
NttBatch(&nttBatchResult, false, count, 0)
var nttResultVecOfVec [][]G1ScalarField
for i := 0; i < batches; i++ {
// Clone the slice
clone := make([]G1ScalarField, len(scalarVecOfVec[i]))
copy(clone, scalarVecOfVec[i])
// Add it to the result vector of vectors
nttResultVecOfVec = append(nttResultVecOfVec, clone)
// Call the ntt_bls12_381 function
Ntt(&nttResultVecOfVec[i], false, 0)
}
assert.NotEqual(t, nttBatchResult, scalars)
// Check that the ntt of each vec of scalars is equal to the intt of the specific batch
for i := 0; i < batches; i++ {
if !reflect.DeepEqual(nttResultVecOfVec[i], nttBatchResult[i*count:((i+1)*count)]) {
t.Errorf("ntt of vec of scalars not equal to intt of specific batch")
}
}
}
func BenchmarkNTT(b *testing.B) {
LOG_NTT_SIZES := []int{12, 15, 20, 21, 22, 23, 24, 25, 26}
for _, logNTTSize := range LOG_NTT_SIZES {
nttSize := 1 << logNTTSize
b.Run(fmt.Sprintf("NTT %d", logNTTSize), func(b *testing.B) {
scalars := GenerateScalars(nttSize, false)
nttResult := make([]G1ScalarField, len(scalars)) // Make a new slice with the same length
copy(nttResult, scalars)
for n := 0; n < b.N; n++ {
Ntt(&nttResult, false, 0)
}
})
}
}

View File

@@ -0,0 +1,38 @@
package bls12381
import "encoding/binary"
// Function to convert [8]uint32 to [4]uint64
func ConvertUint32ArrToUint64Arr(arr32 [8]uint32) [4]uint64 {
var arr64 [4]uint64
for i := 0; i < len(arr32); i += 2 {
arr64[i/2] = (uint64(arr32[i]) << 32) | uint64(arr32[i+1])
}
return arr64
}
func ConvertUint64ArrToUint32Arr4(arr64 [4]uint64) [8]uint32 {
var arr32 [8]uint32
for i, v := range arr64 {
b := make([]byte, 8)
binary.LittleEndian.PutUint64(b, v)
arr32[i*2] = binary.LittleEndian.Uint32(b[0:4])
arr32[i*2+1] = binary.LittleEndian.Uint32(b[4:8])
}
return arr32
}
func ConvertUint64ArrToUint32Arr6(arr64 [6]uint64) [12]uint32 {
var arr32 [12]uint32
for i, v := range arr64 {
b := make([]byte, 8)
binary.LittleEndian.PutUint64(b, v)
arr32[i*2] = binary.LittleEndian.Uint32(b[0:4])
arr32[i*2+1] = binary.LittleEndian.Uint32(b[4:8])
}
return arr32
}

View File

@@ -0,0 +1,81 @@
package bls12381
import (
"testing"
)
func TestConvertUint32ArrToUint64Arr(t *testing.T) {
testCases := []struct {
name string
input [8]uint32
want [4]uint64
}{
{
name: "Test with incremental array",
input: [8]uint32{1, 2, 3, 4, 5, 6, 7, 8},
want: [4]uint64{4294967298, 12884901892, 21474836486, 30064771080},
},
{
name: "Test with all zeros",
input: [8]uint32{0, 0, 0, 0, 0, 0, 0, 0},
want: [4]uint64{0, 0, 0, 0},
},
{
name: "Test with maximum uint32 values",
input: [8]uint32{4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295},
want: [4]uint64{18446744073709551615, 18446744073709551615, 18446744073709551615, 18446744073709551615},
},
{
name: "Test with alternating min and max uint32 values",
input: [8]uint32{0, 4294967295, 0, 4294967295, 0, 4294967295, 0, 4294967295},
want: [4]uint64{4294967295, 4294967295, 4294967295, 4294967295},
},
{
name: "Test with alternating max and min uint32 values",
input: [8]uint32{4294967295, 0, 4294967295, 0, 4294967295, 0, 4294967295, 0},
want: [4]uint64{18446744069414584320, 18446744069414584320, 18446744069414584320, 18446744069414584320},
},
}
for _, tc := range testCases {
t.Run(tc.name, func(t *testing.T) {
got := ConvertUint32ArrToUint64Arr(tc.input)
if got != tc.want {
t.Errorf("got %v, want %v", got, tc.want)
}
})
}
}
func TestConvertUint64ArrToUint32Arr(t *testing.T) {
testCases := []struct {
name string
input [6]uint64
expected [12]uint32
}{
{
name: "test one",
input: [6]uint64{1, 2, 3, 4, 5, 6},
expected: [12]uint32{1, 0, 2, 0, 3, 0, 4, 0, 5, 0, 6, 0},
},
{
name: "test two",
input: [6]uint64{100, 200, 300, 400, 500, 600},
expected: [12]uint32{100, 0, 200, 0, 300, 0, 400, 0, 500, 0, 600, 0},
},
{
name: "test three",
input: [6]uint64{1000, 2000, 3000, 4000, 5000, 6000},
expected: [12]uint32{1000, 0, 2000, 0, 3000, 0, 4000, 0, 5000, 0, 6000, 0},
},
}
for _, tc := range testCases {
t.Run(tc.name, func(t *testing.T) {
got := ConvertUint64ArrToUint32Arr6(tc.input)
if got != tc.expected {
t.Errorf("got %v, want %v", got, tc.expected)
}
})
}
}

View File

@@ -0,0 +1,42 @@
// Copyright 2023 Ingonyama
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
// Code generated by Ingonyama DO NOT EDIT
package bls12381
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbls12_381
// #include "ve_mod_mult.h"
import "C"
import (
"fmt"
"unsafe"
)
func VecScalarMulMod(scalarVec1, scalarVec2 unsafe.Pointer, size int) int {
scalarVec1C := (*C.BLS12_381_scalar_t)(scalarVec1)
scalarVec2C := (*C.BLS12_381_scalar_t)(scalarVec2)
sizeC := C.size_t(size)
ret := C.vec_mod_mult_device_scalar_bls12_381(scalarVec1C, scalarVec2C, sizeC, 0)
if ret != 0 {
fmt.Print("error multiplying scalar vectors")
return -1
}
return 0
}

Some files were not shown because too many files have changed in this diff Show More