Compare commits

..

2 Commits

Author SHA1 Message Date
Vitalii
ae5c0681ad batch mult 2023-07-12 17:42:12 +02:00
Vitalii
58a2492d60 enforce c++17 standard 2023-06-26 21:10:07 +02:00
293 changed files with 7424 additions and 73555 deletions

View File

@@ -1,39 +0,0 @@
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

View File

@@ -1,3 +0,0 @@
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: type:bug
labels: 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: type:feature
labels: enhancement
---
## Description

View File

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

View File

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

46
.github/workflows/build.yml vendored Normal file
View File

@@ -0,0 +1,46 @@
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

View File

@@ -1,20 +0,0 @@
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

View File

@@ -1,50 +0,0 @@
# 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

View File

@@ -1,115 +0,0 @@
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

View File

@@ -1,47 +0,0 @@
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

View File

@@ -1,94 +0,0 @@
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

6
.gitignore vendored
View File

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

View File

@@ -1,10 +0,0 @@
# 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

View File

@@ -1,8 +0,0 @@
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"

28
Cargo.toml Normal file
View File

@@ -0,0 +1,28 @@
[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"
[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 = { version = "0.3.0", optional = true }
rand="*" #TODO: sort dependencies that are not required by the release
[build-dependencies]
cc = { version = "1.0", features = ["parallel"] }
[dev-dependencies]
rand="*"
[features]
default = ["bls12_381"]
bls12_381 = ["ark-bls12-381/curve"]

View File

@@ -1,28 +0,0 @@
# 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"]

187
README.md
View File

@@ -1,124 +1,88 @@
# ICICLE
<div align="center">Icicle is a library for ZK acceleration using CUDA-enabled GPUs.</div>
**<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>
![image (4)](https://user-images.githubusercontent.com/2446179/223707486-ed8eb5ab-0616-4601-8557-12050df8ccf7.png)
## 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.
We believe GPUs are as important for ZK as for AI.
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.
- 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.
## Zero Knowledge on GPU
## Getting Started
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.
ICICLE is a CUDA implementation of general functions widely used in ZKP.
### Supported primitives
> [!NOTE]
> Developers: We highly recommend reading our [documentation]
- Fields
- Scalars
- Points
- Projective: {x, y, z}
- Affine: {x, y}
- Curves
- [BLS12-381]
> [!TIP]
> Try out ICICLE by running some [examples] using ICICLE in C++ and our Rust bindings
> NOTE: _Support for BN254 and BLS12-377 are planned_
### Prerequisites
## Build and usage
- [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).
> NOTE: [NVCC] is a prerequisite for building.
> [!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.
### Accessing Hardware
If you don't have access to a Nvidia GPU we have some options for you.
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].
If you require more compute and have an interesting research project, we have [bounty and grant programs][GRANT_PROGRAM].
### Build systems
ICICLE has three build systems.
- [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
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].
### Compiling ICICLE
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
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.
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
```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
```
docker build -t <name_of_your_choice> .
docker run --gpus all -it <name_of_your_choice> /bin/bash
### Testing the CUDA code
We are using [googletest] library for testing. To build and run the test suite for finite field and elliptic curve arithmetic, run from the `icicle` folder:
```sh
mkdir -p build
cmake -S . -B build
cmake --build build
cd build && ctest
```
### Rust Bindings
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
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.
> NOTE: A GPU must be detectable and therefore installed if the `ARCH_TYPE` is not set.
Once you have your parameters set, run:
```sh
cargo build --release
```
You'll find a release ready library at `target/release/libicicle_utils.rlib`.
### Example Usage
An example of using the Rust bindings library can be found in our [fast-danksharding implementation][FDI]
## Contributions
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 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)
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.
## License
@@ -127,26 +91,13 @@ ICICLE is distributed under the terms of the MIT License.
See [LICENSE-MIT][LMIT] for details.
<!-- Begin Links -->
[BLS12-381]: ./icicle/curves/
[BLS12-377]: ./icicle/curves/
[BN254]: ./icicle/curves/
[BW6-671]: ./icicle/curves/
[BLS12-381]: ./icicle/curves/bls12_381.cuh
[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 -->

31
build.rs Normal file
View File

@@ -0,0 +1,31 @@
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);
nvcc.cuda(true);
nvcc.debug(false);
nvcc.flag(&arch);
nvcc.flag("--std=c++17");
nvcc.files([
"./icicle/appUtils/ntt/ntt.cu",
"./icicle/appUtils/msm/msm.cu",
"./icicle/appUtils/vector_manipulation/ve_mod_mult.cu",
"./icicle/primitives/projective.cu",
]);
nvcc.compile("ingo_icicle"); //TODO: extension??
}

View File

@@ -1,23 +0,0 @@
# 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

@@ -1,25 +0,0 @@
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

@@ -1,72 +0,0 @@
# 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

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

View File

@@ -1,152 +0,0 @@
#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

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

View File

@@ -1,25 +0,0 @@
# 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

@@ -1,21 +0,0 @@
{
"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

@@ -1,25 +0,0 @@
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

@@ -1,52 +0,0 @@
# 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

View File

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

View File

@@ -1,180 +0,0 @@
#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;
}

View File

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

View File

@@ -1,23 +0,0 @@
# 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

@@ -1,24 +0,0 @@
{
"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

@@ -1,25 +0,0 @@
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

@@ -1,41 +0,0 @@
# 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

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

View File

@@ -1,163 +0,0 @@
#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;
}

View File

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

View File

@@ -1,25 +0,0 @@
# 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

@@ -1,22 +0,0 @@
{
"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

@@ -1,26 +0,0 @@
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

@@ -1,32 +0,0 @@
# 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`

View File

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

View File

@@ -1,99 +0,0 @@
#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;
}

View File

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

View File

@@ -1,27 +0,0 @@
# 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

@@ -1,23 +0,0 @@
{
"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

@@ -1,19 +0,0 @@
[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

@@ -1,56 +0,0 @@
# 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

@@ -1,192 +0,0 @@
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

@@ -1,27 +0,0 @@
# 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

@@ -1,23 +0,0 @@
{
"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

@@ -1,20 +0,0 @@
[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

@@ -1,65 +0,0 @@
# 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

@@ -1,157 +0,0 @@
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

@@ -1,27 +0,0 @@
# 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

@@ -1,23 +0,0 @@
{
"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

@@ -1,14 +0,0 @@
[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

@@ -1,50 +0,0 @@
# 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

@@ -1,50 +0,0 @@
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
View File

@@ -1,17 +0,0 @@
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
View File

@@ -1,20 +0,0 @@
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=

View File

@@ -1,34 +0,0 @@
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)

View File

@@ -1,82 +0,0 @@
# 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

@@ -1,328 +0,0 @@
// 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

@@ -1,198 +0,0 @@
// 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

@@ -1,102 +0,0 @@
// 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

@@ -1,79 +0,0 @@
// 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

@@ -1,98 +0,0 @@
// 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

@@ -1,195 +0,0 @@
// 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

@@ -1,50 +0,0 @@
// 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

@@ -1,49 +0,0 @@
// 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

@@ -1,209 +0,0 @@
// 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

@@ -1,360 +0,0 @@
// 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

@@ -1,222 +0,0 @@
// 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

@@ -1,148 +0,0 @@
// 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

@@ -1,38 +0,0 @@
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

@@ -1,42 +0,0 @@
// 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

@@ -1,328 +0,0 @@
// 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

@@ -1,198 +0,0 @@
// 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

@@ -1,102 +0,0 @@
// 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

@@ -1,79 +0,0 @@
// 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

@@ -1,98 +0,0 @@
// 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

@@ -1,195 +0,0 @@
// 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

@@ -1,50 +0,0 @@
// 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

@@ -1,49 +0,0 @@
// 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

@@ -1,209 +0,0 @@
// 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

@@ -1,360 +0,0 @@
// 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

@@ -1,222 +0,0 @@
// 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

@@ -1,148 +0,0 @@
// 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

@@ -1,38 +0,0 @@
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

@@ -1,81 +0,0 @@
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

@@ -1,42 +0,0 @@
// 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
}

View File

@@ -1,328 +0,0 @@
// 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 bn254
import (
"unsafe"
"encoding/binary"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbn254
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
const SCALAR_SIZE = 8
const BASE_SIZE = 8
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.BN254_projective_t)(unsafe.Pointer(p))
in := (*C.BN254_affine_t)(unsafe.Pointer(affine))
C.projective_from_affine_bn254(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.BN254_scalar_t)(unsafe.Pointer(p))
C.random_scalar_bn254(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
}
/*
* PointBN254
*/
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 *PointBN254 to *C.BN254_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.BN254_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BN254_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_bn254(pC, pCompareC))
}
func (p *G1ProjectivePoint) IsOnCurve() bool {
point := (*C.BN254_projective_t)(unsafe.Pointer(p))
res := C.projective_is_on_curve_bn254(point)
return bool(res)
}
func (p *G1ProjectivePoint) Random() *G1ProjectivePoint {
outC := (*C.BN254_projective_t)(unsafe.Pointer(p))
C.random_projective_bn254(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
}
/*
* PointAffineNoInfinityBN254
*/
type G1PointAffine struct {
X, Y G1BaseField
}
func (p *G1PointAffine) FromProjective(projective *G1ProjectivePoint) *G1PointAffine {
in := (*C.BN254_projective_t)(unsafe.Pointer(projective))
out := (*C.BN254_affine_t)(unsafe.Pointer(p))
C.projective_to_affine_bn254(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.BN254_projective_t)(unsafe.Pointer(&a[0]))
scalarsC := (*C.BN254_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_point_bn254(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.BN254_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BN254_scalar_t)(unsafe.Pointer(&b[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.vec_mod_mult_scalar_bn254(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.BN254_scalar_t)(unsafe.Pointer(&a[0]))
bC := (*C.BN254_scalar_t)(unsafe.Pointer(&b[0]))
cC := (*C.BN254_scalar_t)(unsafe.Pointer(&c[0]))
deviceIdC := C.size_t(deviceID)
nElementsC := C.size_t(len(a))
C.matrix_vec_mod_mult_bn254(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

@@ -1,198 +0,0 @@
// 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 bn254
import (
"encoding/binary"
"testing"
"github.com/stretchr/testify/assert"
)
func TestNewFieldBN254One(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 TestNewFieldBN254Zero(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 TestFieldBN254ToBytesLe(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 TestNewPointBN254Zero(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 TestBN254Eq(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 TestBN254StripZ(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 TestPointBN254fromLimbs(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 TestNewPointAffineNoInfinityBN254Zero(t *testing.T) {
var zeroP G1PointAffine
var zeroSanity G1BaseField
zeroSanity.SetZero()
assert.Equal(t, zeroP.X, zeroSanity)
assert.Equal(t, zeroP.Y, zeroSanity)
}
func TestPointAffineNoInfinityBN254FromLimbs(t *testing.T) {
// Initialize your test values
x := [8]uint32{1, 2, 3, 4, 5, 6, 7, 8}
y := [8]uint32{9, 10, 11, 12, 13, 14, 15, 16}
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

@@ -1,102 +0,0 @@
// 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 bn254
import (
"encoding/binary"
"unsafe"
)
// #cgo CFLAGS: -I./include/
// #cgo CFLAGS: -I/usr/local/cuda/include
// #cgo LDFLAGS: -L${SRCDIR}/../../ -lbn254
// #include "projective.h"
// #include "ve_mod_mult.h"
import "C"
// G2 extension field
type G2Element [4]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.BN254_g2_projective_t)(unsafe.Pointer(p))
C.random_g2_projective_bn254(outC)
return p
}
func (p *G2Point) FromAffine(affine *G2PointAffine) *G2Point {
out := (*C.BN254_g2_projective_t)(unsafe.Pointer(p))
in := (*C.BN254_g2_affine_t)(unsafe.Pointer(affine))
C.g2_projective_from_affine_bn254(out, in)
return p
}
func (p *G2Point) Eq(pCompare *G2Point) bool {
// Cast *PointBN254 to *C.BN254_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.BN254_g2_projective_t)(unsafe.Pointer(p))
pCompareC := (*C.BN254_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_bn254(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.BN254_g2_affine_t)(unsafe.Pointer(p))
in := (*C.BN254_g2_projective_t)(unsafe.Pointer(projective))
C.g2_projective_to_affine_bn254(out, in)
return p
}
func (p *G2Point) IsOnCurve() bool {
// Directly copy memory from the C struct to the Go struct
point := (*C.BN254_g2_projective_t)(unsafe.Pointer(p))
res := C.g2_projective_is_on_curve_bn254(point)
return bool(res)
}

View File

@@ -1,79 +0,0 @@
// 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 bn254
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

@@ -1,94 +0,0 @@
// 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 _BN254_MSM_H
#define _BN254_MSM_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BN254 projective and affine structs
typedef struct BN254_projective_t BN254_projective_t;
typedef struct BN254_g2_projective_t BN254_g2_projective_t;
typedef struct BN254_affine_t BN254_affine_t;
typedef struct BN254_g2_affine_t BN254_g2_affine_t;
typedef struct BN254_scalar_t BN254_scalar_t;
typedef cudaStream_t CudaStream_t;
int msm_cuda_bn254(
BN254_projective_t* out, BN254_affine_t* points, BN254_scalar_t* scalars, size_t count, size_t device_id);
int msm_batch_cuda_bn254(
BN254_projective_t* out,
BN254_affine_t* points,
BN254_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_cuda_bn254(
BN254_projective_t* d_out,
BN254_scalar_t* d_scalars,
BN254_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_cuda_bn254(
BN254_projective_t* d_out,
BN254_scalar_t* d_scalars,
BN254_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id);
int msm_g2_cuda_bn254(
BN254_g2_projective_t* out, BN254_g2_affine_t* points, BN254_scalar_t* scalars, size_t count, size_t device_id);
int msm_batch_g2_cuda_bn254(
BN254_g2_projective_t* out,
BN254_g2_affine_t* points,
BN254_scalar_t* scalars,
size_t batch_size,
size_t msm_size,
size_t device_id);
int commit_g2_cuda_bn254(
BN254_g2_projective_t* d_out,
BN254_scalar_t* d_scalars,
BN254_g2_affine_t* d_points,
size_t count,
unsigned large_bucket_factor,
size_t device_id);
int commit_batch_g2_cuda_bn254(
BN254_g2_projective_t* d_out,
BN254_scalar_t* d_scalars,
BN254_g2_affine_t* d_points,
size_t count,
size_t batch_size,
size_t device_id,
cudaStream_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BN254_MSM_H */

View File

@@ -1,193 +0,0 @@
// 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 _BN254_NTT_H
#define _BN254_NTT_H
#ifdef __cplusplus
extern "C" {
#endif
// Incomplete declaration of BN254 projective and affine structs
typedef struct BN254_projective_t BN254_projective_t;
typedef struct BN254_affine_t BN254_affine_t;
typedef struct BN254_scalar_t BN254_scalar_t;
typedef struct BN254_g2_projective_t BN254_g2_projective_t;
typedef struct BN254_g2_affine_t BN254_g2_affine_t;
int ntt_cuda_bn254(BN254_scalar_t* arr, uint32_t n, bool inverse, size_t device_id);
int ntt_batch_cuda_bn254(BN254_scalar_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
int ecntt_cuda_bn254(BN254_projective_t* arr, uint32_t n, bool inverse, size_t device_id);
int ecntt_batch_cuda_bn254(
BN254_projective_t* arr, uint32_t arr_size, uint32_t batch_size, bool inverse, size_t device_id);
BN254_scalar_t*
build_domain_cuda_bn254(uint32_t domain_size, uint32_t logn, bool inverse, size_t device_id, size_t stream);
int interpolate_scalars_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
unsigned device_id,
size_t stream);
int interpolate_scalars_batch_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_points_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
size_t device_id,
size_t stream);
int interpolate_points_batch_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int interpolate_scalars_on_coset_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
BN254_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int interpolate_scalars_batch_on_coset_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_evaluations,
BN254_scalar_t* d_domain,
unsigned n,
unsigned batch_size,
BN254_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_scalars_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned device_id,
size_t stream);
int evaluate_scalars_batch_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
size_t device_id,
size_t stream);
int evaluate_points_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
size_t device_id,
size_t stream);
int evaluate_points_batch_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_coefficients,
BN254_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_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BN254_scalar_t* coset_powers,
unsigned device_id,
size_t stream);
int evaluate_scalars_on_coset_batch_cuda_bn254(
BN254_scalar_t* d_out,
BN254_scalar_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BN254_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
BN254_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int evaluate_points_on_coset_batch_cuda_bn254(
BN254_projective_t* d_out,
BN254_projective_t* d_coefficients,
BN254_scalar_t* d_domain,
unsigned domain_size,
unsigned n,
unsigned batch_size,
BN254_scalar_t* coset_powers,
size_t device_id,
size_t stream);
int reverse_order_scalars_cuda_bn254(BN254_scalar_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_scalars_batch_cuda_bn254(BN254_scalar_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int reverse_order_points_cuda_bn254(BN254_projective_t* arr, int n, size_t device_id, size_t stream);
int reverse_order_points_batch_cuda_bn254(
BN254_projective_t* arr, int n, int batch_size, size_t device_id, size_t stream);
int add_scalars_cuda_bn254(
BN254_scalar_t* d_out, BN254_scalar_t* d_in1, BN254_scalar_t* d_in2, unsigned n, size_t stream);
int sub_scalars_cuda_bn254(
BN254_scalar_t* d_out, BN254_scalar_t* d_in1, BN254_scalar_t* d_in2, unsigned n, size_t stream);
int to_montgomery_scalars_cuda_bn254(BN254_scalar_t* d_inout, unsigned n, size_t stream);
int from_montgomery_scalars_cuda_bn254(BN254_scalar_t* d_inout, unsigned n, size_t stream);
// points g1
int to_montgomery_proj_points_cuda_bn254(BN254_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_cuda_bn254(BN254_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_cuda_bn254(BN254_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_cuda_bn254(BN254_affine_t* d_inout, unsigned n, size_t stream);
// points g2
int to_montgomery_proj_points_g2_cuda_bn254(BN254_g2_projective_t* d_inout, unsigned n, size_t stream);
int from_montgomery_proj_points_g2_cuda_bn254(BN254_g2_projective_t* d_inout, unsigned n, size_t stream);
int to_montgomery_aff_points_g2_cuda_bn254(BN254_g2_affine_t* d_inout, unsigned n, size_t stream);
int from_montgomery_aff_points_g2_cuda_bn254(BN254_g2_affine_t* d_inout, unsigned n, size_t stream);
#ifdef __cplusplus
}
#endif
#endif /* _BN254_NTT_H */

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