mirror of
https://github.com/pseXperiments/icicle.git
synced 2026-05-01 03:00:14 -04:00
Compare commits
10 Commits
v2.0.1
...
hadar/msm-
| Author | SHA1 | Date | |
|---|---|---|---|
|
fad3f9ae61 | ||
|
|
90bf45a2f0 | ||
|
|
641380e9b0 | ||
|
|
eeb6f6e452 | ||
|
|
34f0212c0d | ||
|
|
f6758f3447 | ||
|
|
e2ad621f97 | ||
|
|
bdc3da98d6 | ||
|
|
36e288c1fa | ||
|
|
f8d15e2613 |
31
.github/workflows/golang.yml
vendored
31
.github/workflows/golang.yml
vendored
@@ -99,11 +99,40 @@ jobs:
|
||||
path: |
|
||||
icicle/build/lib/libingo_field_${{ matrix.field.name }}.a
|
||||
retention-days: 1
|
||||
|
||||
build-hashes-linux:
|
||||
name: Build hashes on Linux
|
||||
runs-on: [self-hosted, Linux, X64, icicle]
|
||||
needs: [check-changed-files, check-format]
|
||||
strategy:
|
||||
matrix:
|
||||
hash:
|
||||
- name: keccak
|
||||
build_args:
|
||||
steps:
|
||||
- name: Checkout Repo
|
||||
uses: actions/checkout@v4
|
||||
- name: Setup go
|
||||
uses: actions/setup-go@v5
|
||||
with:
|
||||
go-version: '1.20.0'
|
||||
- name: Build
|
||||
working-directory: ./wrappers/golang
|
||||
if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
|
||||
run: ./build.sh -hash=${{ matrix.hash.name }} ${{ matrix.hash.build_args }} # builds a single hash algorithm
|
||||
- name: Upload ICICLE lib artifacts
|
||||
uses: actions/upload-artifact@v4
|
||||
if: needs.check-changed-files.outputs.golang == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
|
||||
with:
|
||||
name: icicle-builds-${{ matrix.hash.name }}-${{ github.workflow }}-${{ github.sha }}
|
||||
path: |
|
||||
icicle/build/lib/libingo_hash.a
|
||||
retention-days: 1
|
||||
|
||||
test-linux:
|
||||
name: Test on Linux
|
||||
runs-on: [self-hosted, Linux, X64, icicle]
|
||||
needs: [check-changed-files, build-curves-linux, build-fields-linux]
|
||||
needs: [check-changed-files, build-curves-linux, build-fields-linux, build-hashes-linux]
|
||||
steps:
|
||||
- name: Checkout Repo
|
||||
uses: actions/checkout@v4
|
||||
|
||||
13
.github/workflows/rust.yml
vendored
13
.github/workflows/rust.yml
vendored
@@ -60,10 +60,10 @@ jobs:
|
||||
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
|
||||
# ignored tests are polynomial tests. Since they conflict with NTT tests, they are executed sperately
|
||||
# ignored tests are polynomial tests. Since they conflict with NTT tests, they are executed separately
|
||||
run: |
|
||||
cargo test --workspace --exclude icicle-babybear --release --verbose --features=g2 -- --test-threads=2 --ignored
|
||||
cargo test --workspace --exclude icicle-babybear --release --verbose --features=g2 -- --test-threads=2
|
||||
cargo test --workspace --exclude icicle-babybear --exclude icicle-stark252 --release --verbose --features=g2 -- --test-threads=2 --ignored
|
||||
cargo test --workspace --exclude icicle-babybear --exclude icicle-stark252 --release --verbose --features=g2 -- --test-threads=2
|
||||
|
||||
- name: Run baby bear tests
|
||||
working-directory: ./wrappers/rust/icicle-fields/icicle-babybear
|
||||
@@ -72,6 +72,13 @@ jobs:
|
||||
cargo test --release --verbose -- --ignored
|
||||
cargo test --release --verbose
|
||||
|
||||
- name: Run stark252 tests
|
||||
working-directory: ./wrappers/rust/icicle-fields/icicle-stark252
|
||||
if: needs.check-changed-files.outputs.rust == 'true' || needs.check-changed-files.outputs.cpp_cuda == 'true'
|
||||
run: |
|
||||
cargo test --release --verbose -- --ignored
|
||||
cargo test --release --verbose
|
||||
|
||||
build-windows:
|
||||
name: Build on Windows
|
||||
runs-on: windows-2022
|
||||
|
||||
@@ -62,8 +62,8 @@ import (
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
...
|
||||
```
|
||||
|
||||
@@ -68,8 +68,8 @@ func GetDefaultNTTConfig[T any](cosetGen T) NTTConfig[T]
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func Main() {
|
||||
|
||||
@@ -42,9 +42,9 @@ package main
|
||||
import (
|
||||
"log"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
@@ -85,9 +85,9 @@ package main
|
||||
import (
|
||||
"log"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
g2 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
g2 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/g2"
|
||||
)
|
||||
|
||||
func main() {
|
||||
|
||||
@@ -11,9 +11,9 @@
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
@@ -167,7 +167,7 @@ Now you may import `g2` package of the specified curve.
|
||||
|
||||
```go
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls254/g2"
|
||||
)
|
||||
```
|
||||
|
||||
@@ -177,9 +177,9 @@ This package include `G2Projective` and `G2Affine` points as well as a `G2Msm` m
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
g2 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
g2 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/g2"
|
||||
)
|
||||
|
||||
func main() {
|
||||
|
||||
@@ -21,9 +21,9 @@ import (
|
||||
"fmt"
|
||||
"sync"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
|
||||
@@ -10,9 +10,9 @@
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr/fft"
|
||||
)
|
||||
|
||||
@@ -15,9 +15,9 @@ Icicle is exposing a number of vector operations which a user can control:
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
@@ -41,9 +41,9 @@ func main() {
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
@@ -67,9 +67,9 @@ func main() {
|
||||
package main
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func main() {
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
function(check_field)
|
||||
set(SUPPORTED_FIELDS babybear)
|
||||
set(SUPPORTED_FIELDS babybear;stark252)
|
||||
|
||||
set(IS_FIELD_SUPPORTED FALSE)
|
||||
set(I 1000)
|
||||
@@ -14,4 +14,4 @@ function(check_field)
|
||||
if (NOT IS_FIELD_SUPPORTED)
|
||||
message( FATAL_ERROR "The value of FIELD variable: ${FIELD} is not one of the supported fields: ${SUPPORTED_FIELDS}" )
|
||||
endif ()
|
||||
endfunction()
|
||||
endfunction()
|
||||
@@ -8,9 +8,9 @@
|
||||
#include "hash/keccak/keccak.cuh"
|
||||
|
||||
extern "C" cudaError_t
|
||||
keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config);
|
||||
keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, keccak::KeccakConfig& config);
|
||||
|
||||
extern "C" cudaError_t
|
||||
keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config);
|
||||
keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, keccak::KeccakConfig& config);
|
||||
|
||||
#endif
|
||||
47
icicle/include/api/stark252.h
Normal file
47
icicle/include/api/stark252.h
Normal file
@@ -0,0 +1,47 @@
|
||||
// WARNING: This file is auto-generated by a script.
|
||||
// Any changes made to this file may be overwritten.
|
||||
// Please modify the code generation script instead.
|
||||
// Path to the code generation script: scripts/gen_c_api.py
|
||||
|
||||
#pragma once
|
||||
#ifndef STARK252_API_H
|
||||
#define STARK252_API_H
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
#include "gpu-utils/device_context.cuh"
|
||||
#include "fields/stark_fields/stark252.cuh"
|
||||
#include "ntt/ntt.cuh"
|
||||
#include "vec_ops/vec_ops.cuh"
|
||||
|
||||
extern "C" cudaError_t stark252_mul_cuda(
|
||||
stark252::scalar_t* vec_a, stark252::scalar_t* vec_b, int n, vec_ops::VecOpsConfig& config, stark252::scalar_t* result);
|
||||
|
||||
extern "C" cudaError_t stark252_add_cuda(
|
||||
stark252::scalar_t* vec_a, stark252::scalar_t* vec_b, int n, vec_ops::VecOpsConfig& config, stark252::scalar_t* result);
|
||||
|
||||
extern "C" cudaError_t stark252_sub_cuda(
|
||||
stark252::scalar_t* vec_a, stark252::scalar_t* vec_b, int n, vec_ops::VecOpsConfig& config, stark252::scalar_t* result);
|
||||
|
||||
extern "C" cudaError_t stark252_transpose_matrix_cuda(
|
||||
const stark252::scalar_t* input,
|
||||
uint32_t row_size,
|
||||
uint32_t column_size,
|
||||
stark252::scalar_t* output,
|
||||
device_context::DeviceContext& ctx,
|
||||
bool on_device,
|
||||
bool is_async);
|
||||
|
||||
extern "C" void stark252_generate_scalars(stark252::scalar_t* scalars, int size);
|
||||
|
||||
extern "C" cudaError_t stark252_scalar_convert_montgomery(
|
||||
stark252::scalar_t* d_inout, size_t n, bool is_into, device_context::DeviceContext& ctx);
|
||||
|
||||
extern "C" cudaError_t stark252_initialize_domain(
|
||||
stark252::scalar_t* primitive_root, device_context::DeviceContext& ctx, bool fast_twiddles_mode);
|
||||
|
||||
extern "C" cudaError_t stark252_ntt_cuda(
|
||||
const stark252::scalar_t* input, int size, ntt::NTTDir dir, ntt::NTTConfig<stark252::scalar_t>& config, stark252::scalar_t* output);
|
||||
|
||||
extern "C" cudaError_t stark252_release_domain(device_context::DeviceContext& ctx);
|
||||
|
||||
#endif
|
||||
@@ -74,8 +74,13 @@ public:
|
||||
return {X3, Y3, Z3};
|
||||
}
|
||||
|
||||
friend HOST_DEVICE_INLINE Projective operator+(Projective p1, const Projective& p2)
|
||||
friend HOST_DEVICE_INLINE Projective operator+(Projective p1, const Projective& p2)
|
||||
{
|
||||
// FF X3 = p1.x * p2.x;
|
||||
// FF Y3 = p1.y * p2.y;
|
||||
// FF Z3 = p1.y * p2.z;
|
||||
|
||||
|
||||
const FF X1 = p1.x; // < 2
|
||||
const FF Y1 = p1.y; // < 2
|
||||
const FF Z1 = p1.z; // < 2
|
||||
@@ -109,15 +114,174 @@ public:
|
||||
const auto t24 = FF::mul_wide(t12, t23); // t24 ← t12 · t23 < 2
|
||||
const auto t25 = FF::mul_wide(t07, t22); // t25 ← t07 · t22 < 2
|
||||
const FF X3 = FF::reduce(t25 - t24); // X3 ← t25 − t24 < 2
|
||||
// const FF X3 = FF::Wide::get_higher(t25 - t24); // X3 ← t25 − t24 < 2
|
||||
const auto t27 = FF::mul_wide(t23, t19); // t27 ← t23 · t19 < 2
|
||||
const auto t28 = FF::mul_wide(t22, t21); // t28 ← t22 · t21 < 2
|
||||
const FF Y3 = FF::reduce(t28 + t27); // Y3 ← t28 + t27 < 2
|
||||
// const FF Y3 = FF::Wide::get_higher(t28 + t27); // Y3 ← t28 + t27 < 2
|
||||
const auto t30 = FF::mul_wide(t19, t07); // t30 ← t19 · t07 < 2
|
||||
const auto t31 = FF::mul_wide(t21, t12); // t31 ← t21 · t12 < 2
|
||||
const FF Z3 = FF::reduce(t31 + t30); // Z3 ← t31 + t30 < 2
|
||||
// const FF Z3 = FF::Wide::get_higher(t31 + t30); // Z3 ← t31 + t30 < 2
|
||||
return {X3, Y3, Z3};
|
||||
}
|
||||
|
||||
friend DEVICE_INLINE Projective ec_add(Projective p1, const Projective& p2)
|
||||
{
|
||||
// FF X3 = modmul(p1.x,p2.x);
|
||||
// FF Y3 = modmul(p1.y,p2.y);
|
||||
// FF Z3 = modmul(p1.y,p2.z);
|
||||
|
||||
FF t00 = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
FF t01 = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
FF t02 = modmul(p1.y,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
FF t03 = p1.x + p1.y; // t03 ← X1 + Y1 < 4
|
||||
FF t04 = p2.x + p2.y; // t04 ← X2 + Y2 < 4
|
||||
t03 = modmul(t03, t04); // t03 ← t03 · t04 < 3
|
||||
t04 = t00 + t01; // t06 ← t00 + t01 < 4
|
||||
FF t07 = t03 - t04; // t05 ← t05 − t06 < 2
|
||||
FF t08 = p1.y + p1.z; // t08 ← Y1 + Z1 < 4
|
||||
FF t09 = p2.y + p2.z; // t09 ← Y2 + Z2 < 4
|
||||
t08 = modmul(t08, t09); // t10 ← t08 · t09 < 3
|
||||
t09 = t01 + t02; // t11 ← t01 + t02 < 4
|
||||
FF t12 = t08 - t09; // t12 ← t10 − t11 < 2
|
||||
FF t13 = p1.x + p1.z; // t13 ← X1 + Z1 < 4
|
||||
FF t14 = p2.x + p2.z; // t14 ← X2 + Z2 < 4
|
||||
t13 = modmul(t13, t14); // t15 ← t13 · t14 < 3
|
||||
t14 = t00 + t02; // t16 ← t00 + t02 < 4
|
||||
FF t17 = t13 - t14; // t17 ← t15 − t16 < 2
|
||||
FF t18 = t00 + t00; // t18 ← t00 + t00 < 2
|
||||
FF t19 = t18 + t00; // t19 ← t18 + t00 < 2
|
||||
t00 = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(t02)); // t20 ← b3 · t02 < 2
|
||||
FF t21 = t01 + t00; // t21 ← t01 + t20 < 2
|
||||
FF t22 = t01 - t00; // t22 ← t01 − t20 < 2
|
||||
FF t23 = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(t17)); // t23 ← b3 · t17 < 2
|
||||
|
||||
// auto t24 = FF::mul_wide(t12, t23); // t24 ← t12 · t23 < 2
|
||||
// auto t25 = FF::mul_wide(t07, t22); // t25 ← t07 · t22 < 2
|
||||
// FF X3 = FF::reduce(t25 - t24); // X3 ← t25 − t24 < 2
|
||||
// auto t27 = FF::mul_wide(t23, t19); // t27 ← t23 · t19 < 2
|
||||
// auto t28 = FF::mul_wide(t22, t21); // t28 ← t22 · t21 < 2
|
||||
// FF Y3 = FF::reduce(t28 + t27); // Y3 ← t28 + t27 < 2
|
||||
// auto t30 = FF::mul_wide(t19, t07); // t30 ← t19 · t07 < 2
|
||||
// auto t31 = FF::mul_wide(t21, t12); // t31 ← t21 · t12 < 2
|
||||
// FF Z3 = FF::reduce(t31 + t30); // Z3 ← t31 + t30 < 2
|
||||
|
||||
FF X3 = modmul(t12,t23)-modmul(t07,t22); // X3 ← t25 − t24 < 2
|
||||
FF Y3 = modmul(t19,t23)+modmul(t22,t21); // Y3 ← t28 + t27 < 2
|
||||
FF Z3 = modmul(t07,t19)+modmul(t21,t12); // Z3 ← t31 + t30 < 2
|
||||
|
||||
// FF X3 = t21;
|
||||
// FF Y3 = t22;
|
||||
// FF Z3 = t23;
|
||||
return {X3, Y3, Z3};
|
||||
}
|
||||
|
||||
static HOST_DEVICE_INLINE void accumulate_proj(Projective &p1, const Projective& p2)
|
||||
{
|
||||
FF T1,T2,T3,T4,T5,T6;
|
||||
typename FF::Wide W1,W2;
|
||||
T1 = p1.x + p1.y; // t03 ← X1 + Y1 < 4
|
||||
T2 = p2.x + p2.y; // t04 ← X2 + Y2 < 4
|
||||
T1 = T1 * T2; // t03 ← t03 · t04 < 3
|
||||
T2 = p1.y + p1.z; // t08 ← Y1 + Z1 < 4
|
||||
T3 = p2.y + p2.z; // t09 ← Y2 + Z2 < 4
|
||||
T2 = T2 * T3; // t10 ← t08 · t09 < 3
|
||||
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
|
||||
T3 = p1.x + p1.z; // t13 ← X1 + Z1 < 4
|
||||
T4 = p2.x + p2.z; // t14 ← X2 + Z2 < 4
|
||||
T3 = T3 * T4; // t15 ← t13 · t14 < 3
|
||||
|
||||
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
|
||||
T4 = p1.x + p1.y; // t06 ← t00 + t01 < 4
|
||||
T5 = p1.x + p1.z; // t16 ← t00 + t02 < 4
|
||||
T6 = p1.y + p1.z; // t11 ← t01 + t02 < 4
|
||||
|
||||
|
||||
T1 = T1 - T4; // t05 ← t05 − t06 < 2
|
||||
T2 = T2 - T5; // t12 ← t10 − t11 < 2
|
||||
T3 = T3 - T6; // t17 ← t15 − t16 < 2
|
||||
T3 = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(T3)); // t23 ← b3 · t17 < 2
|
||||
|
||||
T4 = p1.y + p1.z; // t21 ← t01 + t20 < 2
|
||||
T5 = p1.y - p1.z; // t22 ← t01 − t20 < 2
|
||||
p1.x = p1.x + p1.x; // t18 ← t00 + t00 < 2
|
||||
T6 = p1.x + p1.x; // t19 ← t18 + t00 < 2
|
||||
p1.z = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(p1.z)); // t20 ← b3 · t02 < 2
|
||||
|
||||
W1 = FF::mul_wide(T1, T5); // t24 ← t12 · t23 < 2
|
||||
W2 = FF::mul_wide(T2, T3); // t25 ← t07 · t22 < 2
|
||||
p1.x = FF::reduce(W1 - W2); // X3 ← t25 − t24 < 2
|
||||
W1 = FF::mul_wide(T4, T5); // t27 ← t23 · t19 < 2
|
||||
W2 = FF::mul_wide(T2, T6); // t28 ← t22 · t21 < 2
|
||||
p1.y = FF::reduce(W1 + W2); // Y3 ← t28 + t27 < 2
|
||||
W1 = FF::mul_wide(T3, T4); // t30 ← t19 · t07 < 2
|
||||
W2 = FF::mul_wide(T1, T6); // t31 ← t21 · t12 < 2
|
||||
p1.z = FF::reduce(W1 + W2); // Z3 ← t31 + t30 < 2
|
||||
// return {X3, Y3, Z3};
|
||||
}
|
||||
|
||||
static HOST_DEVICE_INLINE void accumulate_simple(Projective &p1, const Projective& p2)
|
||||
{
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p1.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p1.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p1.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p1.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p1.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p1.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
p1.x = p1.x * p2.x; // t00 ← X1 · X2 < 2
|
||||
p1.y = p1.y * p2.y; // t01 ← Y1 · Y2 < 2
|
||||
p1.z = p1.z * p2.z; // t02 ← Z1 · Z2 < 2
|
||||
}
|
||||
static HOST_DEVICE_INLINE void accumulate_simple_zpmul(Projective &p1, const Projective& p2)
|
||||
{
|
||||
p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
p1.x = modmul(p1.x,p1.x); // t00 ← X1 · X2 < 2
|
||||
p1.y = modmul(p1.y,p1.y); // t01 ← Y1 · Y2 < 2
|
||||
p1.z = modmul(p1.z,p1.z); // t02 ← Z1 · Z2 < 2
|
||||
p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
// p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
// p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
// p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
// p1.x = modmul(p1.x,p1.x); // t00 ← X1 · X2 < 2
|
||||
// p1.y = modmul(p1.y,p1.y); // t01 ← Y1 · Y2 < 2
|
||||
// p1.z = modmul(p1.z,p1.z); // t02 ← Z1 · Z2 < 2
|
||||
// p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
// p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
// p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
// p1.x = modmul(p1.x,p2.x); // t00 ← X1 · X2 < 2
|
||||
// p1.y = modmul(p1.y,p2.y); // t01 ← Y1 · Y2 < 2
|
||||
// p1.z = modmul(p1.z,p2.z); // t02 ← Z1 · Z2 < 2
|
||||
}
|
||||
|
||||
|
||||
friend HOST_DEVICE_INLINE Projective operator-(Projective p1, const Projective& p2) { return p1 + neg(p2); }
|
||||
|
||||
friend HOST_DEVICE_INLINE Projective operator+(Projective p1, const Affine<FF>& p2)
|
||||
@@ -227,6 +391,13 @@ public:
|
||||
for (int i = 0; i < size; i++)
|
||||
out[i] = (i % size < 100) ? to_affine(rand_host()) : out[i - 100];
|
||||
}
|
||||
|
||||
|
||||
static __device__ __forceinline__ void warpShuffleProjective(Projective& r, const Projective& pt, const uint32_t lane) {
|
||||
FF::warpShuffleField(r.x, pt.x, lane);
|
||||
FF::warpShuffleField(r.y, pt.y, lane);
|
||||
FF::warpShuffleField(r.z, pt.z, lane);
|
||||
}
|
||||
};
|
||||
|
||||
template <typename FF, class SCALAR_FF, const FF& B_VALUE, const FF& GENERATOR_X, const FF& GENERATOR_Y>
|
||||
|
||||
@@ -24,6 +24,9 @@
|
||||
#include "host_math.cuh"
|
||||
#include "ptx.cuh"
|
||||
#include "storage.cuh"
|
||||
#include "asm.cu"
|
||||
#include "Chain.cu"
|
||||
#include "MP.cu"
|
||||
|
||||
#include <iomanip>
|
||||
#include <iostream>
|
||||
@@ -828,7 +831,7 @@ public:
|
||||
// As mentioned, either 2 or 1 reduction can be performed depending on the field in question.
|
||||
if (num_of_reductions() == 2) {
|
||||
carry = sub_limbs<true>(r.limbs_storage, get_modulus<2>(), r_reduced);
|
||||
if (carry == 0) r = Field{r_reduced};
|
||||
if (carry == 0) r = Field{r_reduced}; //r=c*r+(1-c)*r_red
|
||||
}
|
||||
carry = sub_limbs<true>(r.limbs_storage, get_modulus<1>(), r_reduced);
|
||||
if (carry == 0) r = Field{r_reduced};
|
||||
@@ -839,9 +842,23 @@ public:
|
||||
friend HOST_DEVICE_INLINE Field operator*(const Field& xs, const Field& ys)
|
||||
{
|
||||
Wide xy = mul_wide(xs, ys); // full mult
|
||||
// return Wide::get_higher(xy); // reduce mod p
|
||||
return reduce(xy); // reduce mod p
|
||||
}
|
||||
|
||||
friend DEVICE_INLINE Field modmul(const Field& xs, const Field& ys)
|
||||
{
|
||||
Field r;
|
||||
Field localN;
|
||||
uint64_t evenOdd[TLC];
|
||||
bool carry;
|
||||
|
||||
localN = Field{get_modulus()};
|
||||
carry=mp_mul_red_cl<TLC>(evenOdd, xs.limbs_storage.limbs, ys.limbs_storage.limbs, localN.limbs_storage.limbs); //generalize qterm
|
||||
mp_merge_cl<TLC>(r.limbs_storage.limbs, evenOdd, carry);
|
||||
return r;
|
||||
}
|
||||
|
||||
friend HOST_DEVICE_INLINE bool operator==(const Field& xs, const Field& ys)
|
||||
{
|
||||
#ifdef __CUDA_ARCH__
|
||||
@@ -872,6 +889,7 @@ public:
|
||||
for (unsigned i = 1; i < TLC; i++)
|
||||
is_u32 &= (mul.limbs_storage.limbs[i] == 0);
|
||||
|
||||
// return mul_unsigned<multiplier.limbs_storage.limbs[0], Field>(xs);
|
||||
if (is_u32) return mul_unsigned<multiplier.limbs_storage.limbs[0], Field>(xs);
|
||||
return mul * xs;
|
||||
}
|
||||
@@ -990,6 +1008,13 @@ public:
|
||||
}
|
||||
return (u == one) ? b : c;
|
||||
}
|
||||
|
||||
|
||||
static __device__ __forceinline__ void warpShuffleField(Field& r, const Field& f, const uint32_t lane) {
|
||||
#pragma unroll
|
||||
for(int32_t i=0;i<TLC;i++)
|
||||
r.limbs_storage.limbs[i]=__shfl_sync(0xFFFFFFFF, f.limbs_storage.limbs[i], lane);
|
||||
}
|
||||
};
|
||||
|
||||
template <class CONFIG>
|
||||
|
||||
@@ -30,6 +30,9 @@ namespace field_config = grumpkin;
|
||||
#elif FIELD_ID == BABY_BEAR
|
||||
#include "fields/stark_fields/babybear.cuh"
|
||||
namespace field_config = babybear;
|
||||
#elif FIELD_ID == STARK_252
|
||||
#include "fields/stark_fields/stark252.cuh"
|
||||
namespace field_config = stark252;
|
||||
#endif
|
||||
|
||||
#endif
|
||||
@@ -9,5 +9,6 @@
|
||||
#define GRUMPKIN 5
|
||||
|
||||
#define BABY_BEAR 1001
|
||||
#define STARK_252 1002
|
||||
|
||||
#endif
|
||||
631
icicle/include/fields/stark_fields/stark252.cuh
Normal file
631
icicle/include/fields/stark_fields/stark252.cuh
Normal file
@@ -0,0 +1,631 @@
|
||||
#pragma once
|
||||
|
||||
#include "fields/storage.cuh"
|
||||
#include "fields/field.cuh"
|
||||
|
||||
// modulus = 3618502788666131213697322783095070105623107215331596699973092056135872020481 (2^251+17*2^192+1)
|
||||
namespace stark252 {
|
||||
struct fp_config {
|
||||
static constexpr unsigned limbs_count = 8;
|
||||
static constexpr unsigned modulus_bit_count = 252;
|
||||
static constexpr unsigned num_of_reductions = 1;
|
||||
static constexpr unsigned omegas_count = 192;
|
||||
|
||||
static constexpr storage<limbs_count> modulus = {0x00000001, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000011, 0x08000000};
|
||||
static constexpr storage<limbs_count> modulus_2 = {0x00000002, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000022, 0x10000000};
|
||||
static constexpr storage<limbs_count> modulus_4 = {0x00000004, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000044, 0x20000000};
|
||||
static constexpr storage<limbs_count> neg_modulus = {0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff,
|
||||
0xffffffff, 0xffffffff, 0xffffffee, 0xf7ffffff};
|
||||
static constexpr storage<2 * limbs_count> modulus_wide = {
|
||||
0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000011, 0x08000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000};
|
||||
static constexpr storage<2 * limbs_count> modulus_squared = {
|
||||
0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000022, 0x10000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000121, 0x10000000, 0x00000001, 0x00400000};
|
||||
static constexpr storage<2 * limbs_count> modulus_squared_2 = {
|
||||
0x00000002, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000044, 0x20000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000242, 0x20000000, 0x00000002, 0x00800000};
|
||||
static constexpr storage<2 * limbs_count> modulus_squared_4 = {
|
||||
0x00000004, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000088, 0x40000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000484, 0x40000000, 0x00000004, 0x01000000};
|
||||
static constexpr storage<limbs_count> m = {0x8c81fffb, 0x00000002, 0xfeccf000, 0xffffffff,
|
||||
0x0000907f, 0x00000000, 0xffffffbc, 0x1fffffff};
|
||||
static constexpr storage<limbs_count> one = {0x00000001, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000};
|
||||
static constexpr storage<limbs_count> zero = {0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000000, 0x00000000, 0x00000000, 0x00000000};
|
||||
static constexpr storage<limbs_count> montgomery_r = {0xffffffe1, 0xffffffff, 0xffffffff, 0xffffffff,
|
||||
0xffffffff, 0xffffffff, 0xfffffdf0, 0x07ffffff};
|
||||
static constexpr storage<limbs_count> montgomery_r_inv = {0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
0x00000121, 0x10000000, 0x00000001, 0x00400000};
|
||||
|
||||
static constexpr storage_array<omegas_count, limbs_count> omega = {
|
||||
{{0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000011, 0x08000000},
|
||||
{0xf41337e3, 0x2a616626, 0xac8320da, 0xc5268e56, 0x4329f8c7, 0x53312066, 0x29a2995b, 0x06250239},
|
||||
{0xee6feebb, 0x3ada5e1d, 0xe4412e87, 0x98c62155, 0x2f9c676e, 0xc90adb1e, 0x0de874d9, 0x063365fe},
|
||||
{0x6021e539, 0x8337c45f, 0xbbf30245, 0xb0bdf467, 0x514425f3, 0x4537602d, 0x88826aba, 0x05ec467b},
|
||||
{0x9b48a8ab, 0x2225638f, 0x1a8e7981, 0x26da375d, 0xce6246af, 0xfcdca219, 0x9ecd5c85, 0x0789ad45},
|
||||
{0xb2703765, 0xd6871506, 0xf9e225ec, 0xd09bd064, 0x10826800, 0x5e869a07, 0xe82b2bb5, 0x0128f0fe},
|
||||
{0xdd4af20f, 0xfdab65db, 0x56f9ddbc, 0xefa66822, 0x1b03a097, 0x587781ce, 0x9556f9b8, 0x000fcad1},
|
||||
{0xff0cb347, 0x9f1bc8d7, 0xd0e87cd5, 0xc4d78992, 0xdd51a717, 0xbc7924d5, 0xfd121b58, 0x00c92ecb},
|
||||
{0xc13a1d0b, 0xcc4074a0, 0xe3bc8e32, 0xa1f811a9, 0x6d4b9bd4, 0x0234b46e, 0x7880b4dc, 0x011d07d9},
|
||||
{0xec89c4f1, 0xa206c054, 0xdc125289, 0x653d9e35, 0x711825f5, 0x72406af6, 0x46a03edd, 0x0659d839},
|
||||
{0x0fa30710, 0x45391692, 0x11b54c6c, 0xd439f572, 0xa3492c1e, 0xed5ebbf4, 0xb5d9a6de, 0x010f4d91},
|
||||
{0x7afd187f, 0x9273dbbc, 0x91ee171f, 0xdb5375bc, 0x6749ae3d, 0xc061f425, 0x6ec477cf, 0x003d14df},
|
||||
{0x3112b02d, 0x8171e1da, 0xadf9bf78, 0x5c4564eb, 0x5689b232, 0x68c34184, 0x6538624f, 0x0363d70a},
|
||||
{0x606082e1, 0x3e5a42f0, 0x76fc314a, 0x5edd09f0, 0x0f673d7c, 0xd650df25, 0x34832dba, 0x0393a32b},
|
||||
{0x13a77460, 0xe3efc75d, 0x62ef8a01, 0x93898bc8, 0x8bdbd9b3, 0x1c3a6e5c, 0x611b7206, 0x034b5d5d},
|
||||
{0x309d9da9, 0x80ee9837, 0xf51eddbc, 0x1646d633, 0x4901fab8, 0xb9d2cd85, 0x9978ee09, 0x01eb6d84},
|
||||
{0x2755bfac, 0xa7b1f98c, 0xeb7aa1c1, 0x9ec8116c, 0x3109e611, 0x0eeadedd, 0xc9761a8a, 0x06a6f98d},
|
||||
{0x9745a046, 0xce7b0a8b, 0xe411ee63, 0x7ff61841, 0x635f8799, 0x34f67453, 0xef852560, 0x04768803},
|
||||
{0xbffaa9db, 0x1727fce0, 0xf973dc22, 0x858f5918, 0x223f6558, 0x3e277fa0, 0xf71614e3, 0x02d25658},
|
||||
{0x8574e81f, 0xe3d47b99, 0x7fc4c648, 0xc727c9af, 0xee93dc85, 0x581d81ca, 0xca8a00d9, 0x0594beaf},
|
||||
{0x0e5ffcb8, 0x00654744, 0xe7c1b2fd, 0x030530a6, 0xecbf157b, 0x27e46d76, 0xbeea04f1, 0x01f4c2bf},
|
||||
{0x3e3a2f4b, 0xead33145, 0xd6482f17, 0xd841544d, 0x8d24a344, 0x9822fb10, 0x31eeac7c, 0x03e43835},
|
||||
{0xb40bdbe8, 0x01af11c3, 0xb32a3b23, 0xd7c9c0a1, 0xcd0be360, 0x81cb2e43, 0xafb3df1a, 0x01054544},
|
||||
{0x77156db2, 0xf6b13488, 0xddc0f211, 0x1ad6f3be, 0xd664f4da, 0xe643d3ea, 0x174a8e80, 0x071a47b8},
|
||||
{0x4ca88ffc, 0xb86b03a4, 0x8ef9a25a, 0x6e3398e6, 0xf5fa4665, 0xce9a0d37, 0x5c437763, 0x06e8e769},
|
||||
{0x4586dbc3, 0x32609f1d, 0xaa2da684, 0x03148f22, 0x4795d346, 0xa679e36b, 0x9e51225c, 0x03d8d2c7},
|
||||
{0xea5f81cf, 0xeac5be9e, 0x64c12e72, 0x102e16b2, 0xfee282e4, 0xce0bc0d9, 0xa93b28f3, 0x01f05206},
|
||||
{0xbb6422f9, 0x258e96d2, 0x617c5468, 0x751615d8, 0x6056f032, 0x27145cb6, 0x81c06d84, 0x057a7971},
|
||||
{0xb030713c, 0xf42231bb, 0x3a96c59e, 0xae9c3f9a, 0xf1ee840c, 0x5397e8e2, 0xf2b87657, 0x05e7deca},
|
||||
{0xf81f58b4, 0x209745aa, 0x91af248d, 0x74a64310, 0xc04b00b7, 0xe566a8e1, 0x80fb4cea, 0x022bde40},
|
||||
{0x5de74517, 0x8265b62b, 0xb9b9f2c9, 0x6a788149, 0xa9565d98, 0x6fec2239, 0x573f0c28, 0x060ac0c4},
|
||||
{0xd3ce8992, 0xc129d0f1, 0x81c43de5, 0x719252eb, 0x48221e1a, 0xfea566de, 0x0be8ced2, 0x050732ed},
|
||||
{0x2216f1c8, 0x9aae0db3, 0xd7220015, 0x95e231ac, 0x6340df6f, 0xbd6ae160, 0x16a6e39c, 0x0166c8e2},
|
||||
{0x76b0a92e, 0x3ccd9d2b, 0x7d671a9d, 0x1feb39d7, 0x2109fd56, 0x3c49a630, 0x5d4ec292, 0x07badc4b},
|
||||
{0x5dd8c4c3, 0x081c3166, 0xec14ba21, 0x9dca12d8, 0xcf93b2e5, 0xf58069e2, 0x571ddc34, 0x02399005},
|
||||
{0x08a616fc, 0x65a19cf4, 0x8aea6ff7, 0x860d442c, 0x6896a559, 0x4f24ab19, 0x3d7f5ae6, 0x0685db92},
|
||||
{0x622478c4, 0x051093f0, 0x3fab8962, 0x5c200627, 0x21254c39, 0x2aa7ae1b, 0x7b116fb9, 0x0100fff9},
|
||||
{0x00637050, 0x2693b834, 0x22440235, 0x3fef7c1b, 0x3481c4fe, 0x31150ac1, 0xf261b6de, 0x0772cb7a},
|
||||
{0xd990d491, 0x6966804c, 0xc7505f35, 0x46aba1bc, 0xaceeb7f7, 0x4f696cba, 0x6474b8f0, 0x02b73cad},
|
||||
{0xf39cd3e8, 0x7d13e948, 0x62a1db76, 0xd5c33593, 0x4d1be159, 0x7fd3b59b, 0x3676644e, 0x066d3f61},
|
||||
{0xb3bd8b7e, 0x5a896ef3, 0xba5762ab, 0x2319450a, 0x1a545f8b, 0x226f0a07, 0x55446d35, 0x02760973},
|
||||
{0x140e5623, 0x38eaa186, 0x94be15ba, 0x5a48d469, 0xad75d32a, 0xe4f1f15b, 0x2f14e2f1, 0x039ccdaa},
|
||||
{0xe6fcfdb2, 0xad7108d3, 0x9c9f7f04, 0xfadfc050, 0x9df95366, 0xdbb20071, 0xe555c739, 0x02c4d3fa},
|
||||
{0xc3111bcb, 0xb640956f, 0xbb11fb86, 0xcd942bbd, 0xa3db81cd, 0xa4b4eb09, 0x684fdb65, 0x041ed5ed},
|
||||
{0xdd5ca525, 0x462b41fa, 0x153c3d28, 0xbcc17ccd, 0x6b06db5c, 0x8a81d137, 0x4a050358, 0x05f5cf39},
|
||||
{0xcc60fb85, 0x374012a6, 0x34d1905d, 0x978f9785, 0x4e17ff38, 0x713383d4, 0x1055c25d, 0x07f3796f},
|
||||
{0x0643771f, 0x852ba56e, 0x86781a31, 0xadfa956c, 0xb26a3811, 0x2ee2fccf, 0xdbd56ba7, 0x009214ce},
|
||||
{0x68bc148c, 0xe2bf6c4b, 0x01c203ce, 0xd38dbf38, 0x97923b55, 0x27f73df4, 0x5081f7d9, 0x030a2e81},
|
||||
{0xf11422a0, 0xbe23b78f, 0x99cdc2e0, 0xd4f3510d, 0xaa13ffe5, 0xcb05b3da, 0xc724e0c5, 0x028d98a5},
|
||||
{0x96934000, 0x15277271, 0x588c8a51, 0x8013dd5e, 0x9ed55af8, 0x77772f7c, 0x03549e60, 0x020895f8},
|
||||
{0x34db29f8, 0xc0cc8556, 0x67455b5d, 0x5582a9ff, 0x8a9a38b5, 0x12862a43, 0xa59fd242, 0x059655bc},
|
||||
{0x94ceaf98, 0x39bc5131, 0xc71ccc0d, 0x99f4d1a0, 0x54acb87c, 0xc565794d, 0xc33590ef, 0x0593fcef},
|
||||
{0xe97bf51c, 0xa2922d09, 0x3200d367, 0xdbb866a2, 0x4ad9302d, 0x05849ed8, 0xdf93f2b5, 0x000c447e},
|
||||
{0x850fb317, 0x2755d6c2, 0xd45eb3f5, 0x36feeeea, 0xdfbc1d97, 0x4f4471d7, 0x4e3003f8, 0x07ec8926},
|
||||
{0xb6a791f1, 0x38b8dc2a, 0x27a1bbb1, 0x79d6de48, 0xcad54cf2, 0x78c40b06, 0xa43bc898, 0x036dd150},
|
||||
{0x1cc4133c, 0xefa72477, 0x477d39be, 0x5327d617, 0x2c5db3a4, 0xfd1de1f9, 0xc9a18a1c, 0x0147819b},
|
||||
{0xf8133966, 0x275e6b02, 0x87969b48, 0x82bc79b9, 0x5d1e2f0e, 0x85b1f9bd, 0xc819531b, 0x00f9ea29},
|
||||
{0x120edfab, 0x9e0392a5, 0xe3681a15, 0x07403ad4, 0x8a1c3817, 0xa8d469d8, 0x89f15c6f, 0x0395e7fc},
|
||||
{0x641826ac, 0x7f405a9f, 0x6861e2ce, 0xa566e755, 0xba82a050, 0x8a3a08ba, 0xea63598d, 0x071dd923},
|
||||
{0x5f65c188, 0x1d2b7538, 0xd6fc9625, 0xcb704d0f, 0xf59deccc, 0x18729111, 0x52fe1979, 0x07595020},
|
||||
{0x8a08756f, 0x0175aa1c, 0x7fa7c6c4, 0x9a76a312, 0x6e93f6f3, 0x0bfa523a, 0x258c2f23, 0x03d70de4},
|
||||
{0x8229376d, 0x8a0b9d02, 0x2c65c94e, 0x08421430, 0xd34b0aa6, 0x1160b441, 0xbbfb9491, 0x03b9eb75},
|
||||
{0x827caf53, 0x91874856, 0x37e8a006, 0xdfdcae7a, 0x04e3af6b, 0x6dcfc3f2, 0xba66ff37, 0x0592823d},
|
||||
{0x72fb8b0d, 0xb0a6628d, 0xa72b1f03, 0x7d3eef8b, 0x8dd54dbe, 0x5be965ba, 0x96d1fe4c, 0x0114a278},
|
||||
{0x06051d55, 0x0256d8e6, 0xb9fa9dcc, 0xbf152353, 0x44140d6e, 0x6ef2c68c, 0xc9c0fea6, 0x015f291a},
|
||||
{0xed992efc, 0xa1826724, 0x771da991, 0x9a58fd99, 0xd0b370a1, 0xce51a153, 0x826df846, 0x03c53bf5},
|
||||
{0xcc7bf8c3, 0x3909aad7, 0xb08ddfa2, 0xd408ae7d, 0xff94d9fc, 0x2e9ab5d6, 0xf11cbcf6, 0x0020a1b2},
|
||||
{0x3e257b43, 0x448fff07, 0x5fd9edca, 0x00f4a128, 0x7b429f71, 0x6f8987e3, 0x0fc8b522, 0x013336c1},
|
||||
{0x062bd860, 0xef78ac4c, 0xf5d787d2, 0x6539ee52, 0xbb65576e, 0x113b6071, 0x9f3d7f85, 0x0160e952},
|
||||
{0xf966d24e, 0x0c4e7c07, 0x318277e8, 0x011853d8, 0x7c287f58, 0x93bae650, 0xf64289f7, 0x00b974a1},
|
||||
{0x30408cb9, 0x66d19420, 0x0430b017, 0x709ca6c6, 0x23d95951, 0xb174ad46, 0x111f4192, 0x030762f8},
|
||||
{0xf246c901, 0xb9d70015, 0x57a1cdec, 0xd3616cb1, 0x0d732fdb, 0x61aab25e, 0x12d620d8, 0x0712858b},
|
||||
{0x16334e1a, 0x8ec7e113, 0xa96aeeab, 0x0021a55b, 0xfd639175, 0x8f4c1366, 0x69bc866a, 0x07acdde9},
|
||||
{0x23088fc7, 0x1fb24e5e, 0x92a88089, 0xcacd65df, 0x17343c48, 0x103ec3c8, 0xc387a3b5, 0x03d296b9},
|
||||
{0xcd9fedee, 0xae703c5b, 0x7853b30d, 0xd0c3e0c6, 0x12abaef5, 0xc1e326b3, 0x5d57bb23, 0x04f42d7f},
|
||||
{0x1824b92c, 0x19cd1b4e, 0x81ebc117, 0xc5daaff4, 0xb8183a1d, 0xeeedaa59, 0xe28baf8a, 0x069d8f0c},
|
||||
{0x9dc50729, 0x9733e8df, 0xf1b9f411, 0xd7e0dbb9, 0x50edf7ea, 0x59e4dbd2, 0x4059cb5f, 0x002259fe},
|
||||
{0xb79a92b1, 0x5e3197fc, 0x59086db1, 0xbfddf5c5, 0xdbea4a69, 0x234d8639, 0x4d0a367d, 0x05dd79b0},
|
||||
{0xa86eec0c, 0x8cc1d845, 0x573b44d7, 0x3cac8839, 0x7b0de880, 0x8b8d8735, 0x68c99722, 0x01c5ef12},
|
||||
{0xc2ba0f23, 0x12680395, 0x471f947e, 0xd43bcf85, 0xcc9d9b24, 0x19935b68, 0x108eec6a, 0x06263e1e},
|
||||
{0x5b7be972, 0x29617bad, 0xc55b1b68, 0x0ab73eef, 0x2544381e, 0x07f12359, 0x63a080a0, 0x0161444d},
|
||||
{0x312f9080, 0x07a4b921, 0x2f530413, 0x64c25a07, 0x7d71ca2f, 0x3f6903d7, 0x04838ba1, 0x06917cab},
|
||||
{0x10bdb6cc, 0xec7cfc1f, 0x3bcf85c7, 0x7046910d, 0x7bc3ff5f, 0x7ef09e22, 0x385306d4, 0x004b0b60},
|
||||
{0x3a41158a, 0x82d06d78, 0xaa690d1f, 0x37c4a361, 0x7117c44a, 0x700766e1, 0xab40d7e4, 0x031261d0},
|
||||
{0x91b88258, 0x384c5e8b, 0x009b84dc, 0xd777abd5, 0xe7eed224, 0x02102b55, 0xdbefe5e9, 0x03b22830},
|
||||
{0x8770a4be, 0xec982f60, 0x961f56ad, 0x4b92533d, 0xf428c4b9, 0x7df85fbb, 0x2d9291a4, 0x057e4876},
|
||||
{0xf4910a60, 0x6ace9477, 0x9fc63b7f, 0xdb5a705f, 0x72328369, 0x4cc157b4, 0xc282db6f, 0x05b8acbc},
|
||||
{0x57269216, 0x4c69edd9, 0xbfee24ac, 0xd04f1eeb, 0x2a069b18, 0xacda8418, 0x5990b523, 0x03761a4f},
|
||||
{0xc608d246, 0x7f2e2048, 0x4664959b, 0xd4f52ed2, 0x11c1d565, 0x354e3bf7, 0x457eabd3, 0x0156d837},
|
||||
{0xd455f483, 0xea8cbefd, 0x5d940684, 0x33cd5725, 0x8091a287, 0x2d89a777, 0x939b3ef3, 0x06159e4a},
|
||||
{0x4fa405aa, 0xe43439f1, 0xdbe5763d, 0xa258cfc7, 0x78d7b607, 0x9491173a, 0x9ad23eac, 0x01775d66},
|
||||
{0xd772d637, 0x2413e92c, 0x5eac4588, 0x22c99c9f, 0x71a0cdd2, 0xa2bd1d06, 0xfdd73a36, 0x05e88acb},
|
||||
{0xb2bfa1ad, 0x68886b35, 0x35d2dfb6, 0x7a969b62, 0x9767a44a, 0x359ddb45, 0x52e5da6d, 0x00f1a46e},
|
||||
{0x1c5a4861, 0x4ef9fe94, 0x1c841a89, 0x1540cf67, 0xa9bed4f5, 0x8b51336f, 0xf63c32ab, 0x0240fc41},
|
||||
{0x87086e50, 0x7f5c626d, 0x049c46e2, 0x38ec0386, 0x0c597ea7, 0x30b003fd, 0x6660a912, 0x07a8faa1},
|
||||
{0x7dac5d19, 0x2810d2b4, 0x80339f39, 0x040470c4, 0xc946ab30, 0x30d97769, 0x52667151, 0x019fa1f9},
|
||||
{0x5e7c57a2, 0x00e13c8e, 0x2a0fb7bd, 0x95490ca0, 0x08451e35, 0x6af2b76d, 0xcf78c579, 0x04c3a3a1},
|
||||
{0x55e39071, 0xa848b2f2, 0xf132ce21, 0x6831da1d, 0xe080e2ec, 0x439bdda4, 0xadd19a7d, 0x06680f09},
|
||||
{0x6be27786, 0xfebd2a8b, 0x093a5a7f, 0x2cdd8f78, 0xdcb004b3, 0xbc0746a1, 0xd12450ed, 0x005f950a},
|
||||
{0x39759f39, 0xe1462ca6, 0x7bbe087d, 0x0c37dca2, 0x0c8661cb, 0x198de347, 0x7e531b52, 0x03602655},
|
||||
{0x66d7eb25, 0xaf24ead2, 0x5ee6eb03, 0x27cea560, 0x4f6267c7, 0xe9aa6d50, 0xe5dd28e0, 0x00c962b1},
|
||||
{0xb11706c9, 0x3c3407a5, 0xcf0e1b88, 0x44370686, 0x9fbda5e3, 0x5d0e7af0, 0x41cf0a6b, 0x010d235f},
|
||||
{0x358cfcc2, 0x1fbc42a3, 0xc78f7dac, 0x5a2e6ea2, 0xa12773f2, 0x33e089ca, 0xed7788c1, 0x04bef156},
|
||||
{0xbea42f88, 0xdb150649, 0x5f3fb72a, 0x71329f69, 0x86b82de7, 0x7aa46ad0, 0xc6093912, 0x07913b17},
|
||||
{0xb3b67067, 0xb2b074ae, 0xc55f4455, 0x4f17674d, 0xdeb0740d, 0x9a112816, 0x316cc0d3, 0x06bd0cde},
|
||||
{0x1a264ab3, 0x962ceb6b, 0xd99f7159, 0xd5930255, 0x24a4096e, 0x7db961b0, 0x3e50dfed, 0x050c8e5c},
|
||||
{0x443af109, 0xc3eebe54, 0x86946633, 0x2ca03fcb, 0x04badff6, 0x6e6eef04, 0x82210754, 0x05d92ab7},
|
||||
{0xa5c0dca4, 0xcbadd8ad, 0x5ac103a0, 0x4cf688cf, 0x26e5d435, 0x571dbdb9, 0x220fc7db, 0x074ffc4d},
|
||||
{0x88740c3e, 0x70b80432, 0x03821aa8, 0x4a959d50, 0xe4df06d8, 0x3eb8c3a0, 0xcac57496, 0x025a425b},
|
||||
{0x55205413, 0xdcadfd29, 0x90b17b01, 0xda7456d2, 0x73696a28, 0x437c2fda, 0x329f6855, 0x00a8a188},
|
||||
{0xa828431e, 0x3cde2cdd, 0x9ed29340, 0x60e6c362, 0x7c13e145, 0xef00dfa9, 0xba288c0b, 0x04159bec},
|
||||
{0x9065f8ee, 0x41d351cd, 0xa4845868, 0x4e2e298f, 0xbdb3834a, 0xbcba6ac1, 0xea85f2ec, 0x042c8871},
|
||||
{0x1fda880f, 0xc4dc0d20, 0x26fc2d5c, 0x4f0f9dc4, 0x86839de7, 0x2c555343, 0xf698dd8f, 0x04d12da8},
|
||||
{0x21bd655a, 0x3a6299bd, 0x8cfd772f, 0x2e4aea22, 0xd2c2590d, 0x09716ad9, 0xb298587d, 0x053b143c},
|
||||
{0xa95e3cbf, 0xd35f3e32, 0x04eac3cf, 0xe380dee7, 0x0f7e3e6b, 0x27e6570a, 0xbed46774, 0x008cd288},
|
||||
{0x9583f023, 0xe42676b0, 0x75cfaa7e, 0x39d57dd6, 0x4f0bb727, 0x10d4a8d0, 0x27c81bdd, 0x016b03c9},
|
||||
{0x4decc603, 0x89b394f7, 0xd24690f4, 0xd7322ee9, 0x947a00fd, 0xbbc12961, 0x82e8fa75, 0x00886d23},
|
||||
{0xeb0faad4, 0x7b48a33b, 0x60e0b0c8, 0x4c11ef26, 0x36f0f791, 0x4163a401, 0xa4074faf, 0x07986fea},
|
||||
{0x31d9587e, 0x96044919, 0x9049fd2d, 0xb1cab341, 0x9c0eea09, 0xf28c83c9, 0x5c6620aa, 0x033b74dd},
|
||||
{0x13ee028c, 0xde558d16, 0x5d4233b0, 0x4dcf3932, 0x2e422803, 0x7bd46887, 0xe1261bff, 0x04b4757d},
|
||||
{0xd48e9b00, 0x6c80848f, 0x10b6a121, 0x937c1e6e, 0xe9f2008c, 0x7782f8b8, 0x2bc7171c, 0x00217358},
|
||||
{0x324228d8, 0xba523265, 0x682ee17c, 0x4ebe5506, 0x3be009f9, 0x6c646fe8, 0x8594b924, 0x046de7bc},
|
||||
{0x3b50645a, 0x270aa33a, 0x2a9c6282, 0x28fd23fd, 0xcfe96515, 0x5b2fa771, 0x3f812377, 0x063039de},
|
||||
{0xaba4060a, 0xa1da52b0, 0x0374be67, 0x7f191fd6, 0x0d7d2126, 0x14c64d05, 0xf7f77381, 0x00419cb7},
|
||||
{0xe4b19319, 0x07eda692, 0x0fef654e, 0x6190d3f6, 0x0b21ca7e, 0x893b0916, 0x073c48b4, 0x0367a3c7},
|
||||
{0xc520e3ea, 0x8fd405b2, 0x487e93c9, 0x73b4f714, 0xd5142cff, 0x70b7ee88, 0xa320eca2, 0x058fb800},
|
||||
{0x72ef3623, 0x3b5a8740, 0xaff370fd, 0xbff4af42, 0xe338258e, 0x64c137b0, 0xc7afafca, 0x05ac9917},
|
||||
{0x82ccc89a, 0x99c46a0d, 0x9ff87868, 0x05ae3209, 0xa489481f, 0x6249b2a4, 0xbaead348, 0x0056c235},
|
||||
{0xba0ea95e, 0x5a0640f3, 0xc03af976, 0x518db5cd, 0x5a250a06, 0x1c3223aa, 0xbc3442eb, 0x0397b942},
|
||||
{0xacf14a4f, 0x164f0705, 0x33eb6c0e, 0x386c2325, 0xd7264573, 0xdfaceff6, 0xd1e22f80, 0x00e94509},
|
||||
{0x9ff51bc7, 0x8964ee48, 0x57bbca04, 0x3e0f5037, 0x6510630c, 0xe78d6c8d, 0xdf0a61c1, 0x041d6351},
|
||||
{0x45aa1b58, 0x47892f3b, 0x915c1c70, 0x5a1787ba, 0x67f20d25, 0xbaa23359, 0x0c4bc4be, 0x00e1919f},
|
||||
{0xb9975332, 0x2a87c37a, 0xcdecebc9, 0x95db523f, 0x1d0db226, 0x703949ee, 0x4c3842dd, 0x03152c1d},
|
||||
{0xecfb6f72, 0x0eff7e6a, 0x9493a628, 0xb3a83455, 0xd596cd51, 0xced58dd1, 0x25ee51ff, 0x033dee78},
|
||||
{0x72a30547, 0x1f4047ca, 0xd40b6d0f, 0x9feefa06, 0x94db1b38, 0x836ffd80, 0xa0992ed5, 0x037c79f6},
|
||||
{0xceb3dffd, 0x7ffa095d, 0x768e2cb3, 0x23097a65, 0x373f6222, 0xd228b1f9, 0xc57feea2, 0x06309a6b},
|
||||
{0xecd4c6f7, 0x7a5bead4, 0x7e70f7de, 0xab92043c, 0x220db8d8, 0xf78f890e, 0x2865a07e, 0x052eeb98},
|
||||
{0xdf253531, 0x8e9a6336, 0xbafa937b, 0xb24b664a, 0x303b1f5a, 0xc89f660e, 0x876bd8c7, 0x07ea9749},
|
||||
{0x1d4c3fec, 0xd958e726, 0x06fbef31, 0xa5eb368f, 0xba6a027d, 0x0c911679, 0x5f80f992, 0x06321b51},
|
||||
{0x046b49b2, 0x3ca61d9e, 0x6aa9c29a, 0x616a47d6, 0x9e9462dc, 0x27a7ffeb, 0x8971b70e, 0x0794ed38},
|
||||
{0x9f47496f, 0xdb259a57, 0xa6b0481c, 0x7f3e3f90, 0x4afab47a, 0x76f42726, 0xc5a79505, 0x07b9da96},
|
||||
{0x57e7aeed, 0x908e6450, 0x81648127, 0xe86db2fb, 0x8dd76882, 0x53f3c573, 0x72327da6, 0x02b37324},
|
||||
{0x73a220ec, 0x82a941c9, 0x7f25beea, 0xb4cbecb7, 0xbfb061d6, 0x746ded71, 0x641b3f3d, 0x00f7af27},
|
||||
{0xcbd4ba67, 0x69b8f4df, 0x3d526981, 0x5ee3ac6f, 0x145cef8c, 0x9372af4e, 0x72a31ef1, 0x05cc1cc6},
|
||||
{0x62d1ba57, 0xce898b0d, 0xee3fa47e, 0x86ba0504, 0x4395b70d, 0xc68233b1, 0x80eb8d60, 0x024cfa58},
|
||||
{0x74d51c41, 0x8fa83850, 0x60f8f9da, 0x5824a285, 0xaf1bea48, 0xa7a2067e, 0x5455acc3, 0x04ba49f2},
|
||||
{0x324c6039, 0x0a1e223e, 0x7b18a9d0, 0x28312228, 0x88b6ecda, 0xb60c1f93, 0x687ba365, 0x053097d8},
|
||||
{0xa7dae551, 0x5604b398, 0xe2e11609, 0x51f02e33, 0xe58e2094, 0x0b51a085, 0x3a3ecc28, 0x078679d6},
|
||||
{0x92d52444, 0xe24b5528, 0x33d0fa70, 0xf77e35ad, 0x9bcbfb57, 0x8af5a7b7, 0x022748d2, 0x015c5f15},
|
||||
{0xc993b168, 0xc002185c, 0x293ad856, 0x5586addb, 0x8ec50726, 0x69c1bfcf, 0x5fd97ea1, 0x00d514fc},
|
||||
{0x8866c747, 0x52d7a9a2, 0x01d6ee05, 0x9bd77465, 0xc3a87a88, 0x576adf96, 0xfa69f0ec, 0x0693e89a},
|
||||
{0x05903be3, 0xcfe50d90, 0xcf739179, 0xbe651dd1, 0x2ae70678, 0xba80ffda, 0xb55b06cc, 0x051dbe40},
|
||||
{0x5585a6f0, 0x4adb5947, 0x9fa37e68, 0x14634b99, 0xa2a910a8, 0x27da5fbf, 0xa99c704d, 0x022a91ce},
|
||||
{0xe2ddaacd, 0xfabab7b8, 0x60cf9603, 0x1edf6a83, 0xbfadddd3, 0x20b04218, 0xa81dbffa, 0x03e0ddb6},
|
||||
{0xda25c9fd, 0xf9c1e3a3, 0xac57ece3, 0x41ff4e1e, 0xdd684055, 0x9ba50868, 0x46d8156a, 0x01b30314},
|
||||
{0xab76a462, 0x30e067cc, 0x08f1b99b, 0x2d84c4c2, 0x73edc56f, 0x6b399ae0, 0x62cfacb2, 0x02f187e1},
|
||||
{0x34fc5356, 0xb085758e, 0xf805fedf, 0xbafe9a1c, 0x95272d01, 0x0bcf423c, 0x1feca651, 0x01df4a81},
|
||||
{0x4c264e97, 0xd3bd9833, 0xc08b1798, 0xc0b192be, 0xdc3ed49e, 0x42724e80, 0xbaee9a58, 0x04100303},
|
||||
{0xe49749c9, 0xb653c919, 0x09f8e2fc, 0x07dbe557, 0xca71e551, 0xbb172d28, 0x7989c8fd, 0x07f5f801},
|
||||
{0xdf1d9004, 0x9412a9f3, 0xbe90d67e, 0xddcf6d66, 0x4692f803, 0x1dbfd679, 0x524c2944, 0x04f4fae1},
|
||||
{0x5707d134, 0xd413afdf, 0x887fd7e9, 0xf8a339cf, 0x84883580, 0xf74544f4, 0x851739e0, 0x0554f72a},
|
||||
{0x59824907, 0xe3827564, 0x421182c9, 0x352eab2a, 0x8f8530f2, 0x19138257, 0x20275950, 0x04e3bf44},
|
||||
{0x33f928b7, 0xef7660f9, 0xf5952362, 0xb7cb0619, 0xf17eb8d7, 0x5b24913b, 0x8e8b8082, 0x00f4804c},
|
||||
{0x5bd84f3e, 0xe7020613, 0x736a1659, 0x7ee777e1, 0x0795844b, 0x34ca7cb6, 0x7503ddc3, 0x07ce12e4},
|
||||
{0x6d8408a5, 0xbbbafb3f, 0x519dadca, 0xe0f02915, 0x0670f5d4, 0x5acba199, 0x4a93340f, 0x0056db45},
|
||||
{0xe404f6c5, 0x73f8a435, 0x01731858, 0x68cd3f7a, 0xd01f3de9, 0x214d3134, 0xd5d75a88, 0x05fb76be},
|
||||
{0xf976eb41, 0x3a66ad86, 0xcd08787a, 0x6401b6d3, 0x7d1e82a8, 0x575950f3, 0x55ee9d49, 0x00e34b33},
|
||||
{0x0cc5cbf4, 0xbff2f4e6, 0xec205dcd, 0x5a6b430d, 0xc94862af, 0xa8114ab3, 0x2fe8be1f, 0x0247ecf5},
|
||||
{0x8b98bf40, 0xded3bc57, 0xe26b66b3, 0xb658c8c4, 0x8d4220db, 0x8bd91c55, 0x94d2adea, 0x00d109f2},
|
||||
{0xedeaec42, 0x0fbfd336, 0x5d407ae8, 0xd94f928d, 0x727e74b5, 0xe5e4a16b, 0xc8c22dd8, 0x06a550df},
|
||||
{0x135e0ee9, 0xe378a012, 0x856a1aef, 0x5be86512, 0xd8febe77, 0x7de04ce2, 0xea43d59b, 0x03ddeed6},
|
||||
{0x005a1d86, 0xc04dc48c, 0x6f29053d, 0x64f4bbd2, 0x9be0aef5, 0x10b1b3db, 0xcc625a0b, 0x03745ca5},
|
||||
{0x1f4f0e85, 0x6c72bd40, 0xc2069cba, 0x4234afd0, 0xb99395f4, 0xc25b262f, 0xae0874e2, 0x0605f6a2},
|
||||
{0xdd756b6d, 0x9513e0d4, 0xf0c137cd, 0x5127a167, 0x7f01c538, 0x1a12a425, 0x00a4483b, 0x068b3aaf},
|
||||
{0x79bc6c86, 0x7a5b3e70, 0x375dc240, 0x5a337909, 0xe111d6ce, 0x46d6fe3c, 0x2ff2ca50, 0x02708b05},
|
||||
{0x1524ad8c, 0x1181eb95, 0x52294490, 0xd0744ddc, 0x848605cf, 0x88ed5b7b, 0xb478c12a, 0x04b9cb49},
|
||||
{0x27105dae, 0x98cb2411, 0xed5c1361, 0x3efa8fae, 0xd498e337, 0x6fa736a5, 0x1e369b4f, 0x038e3b07},
|
||||
{0x98c8db7f, 0xbc5915ae, 0x50425ae8, 0x1f3c8f96, 0xfa86658a, 0x77d60416, 0x28ec2dda, 0x02bc8b30},
|
||||
{0xb94bc10e, 0xad6794f2, 0x7e80093a, 0x7463b3f3, 0x90db4c79, 0x7bf5af53, 0x965c0cc4, 0x031531c6},
|
||||
{0x7cc1083d, 0x66425289, 0xa45d785f, 0x778ba471, 0xbbc94c16, 0xe3f5c599, 0x9b92e036, 0x02606413},
|
||||
{0xcf287faf, 0x191a2ea9, 0x823ddf07, 0xe6406a78, 0xaabe912b, 0xabcf2825, 0x7c48649a, 0x021dab44},
|
||||
{0x65375f6c, 0x9465d77c, 0x65370520, 0x924e189c, 0x918f0105, 0x8be0ca5f, 0xb1925509, 0x07586d27},
|
||||
{0x9302ac44, 0xe4fa93cb, 0xbf87d840, 0xf381ebbd, 0x44793049, 0x5027e7d9, 0xd3f09392, 0x0230b5c3},
|
||||
{0x31d48a82, 0x123e992e, 0x729d40e2, 0xef2990c6, 0x0f331903, 0x946813e3, 0x112a2c4d, 0x022f575e},
|
||||
{0xd4ee8cf7, 0x4b44764e, 0xdb576ebc, 0x4d44cff8, 0x0ab93ba1, 0xc6185d3a, 0x7e3f1e78, 0x0520c2d3},
|
||||
{0xbc46b8b4, 0xd9446736, 0x91e2ede1, 0xc7776293, 0x87689930, 0x0323845f, 0x379293ae, 0x061e359f},
|
||||
{0xb49b3a0a, 0x767a1747, 0x2b58f45e, 0x17e69346, 0x1425ad98, 0x10820519, 0x1b487ae5, 0x0367f384},
|
||||
{0x92f8ac25, 0xe0407696, 0x2beb71a6, 0x9ca9d269, 0x2f0c2471, 0x914017ea, 0xf421a10d, 0x07709cc3},
|
||||
{0xc3bb6a8f, 0x2c8ed622, 0xa2a1a8f2, 0x31c57cb6, 0x4bf6c316, 0x053924d5, 0x09563089, 0x0727b76a},
|
||||
{0x09dc6b5c, 0x567be37f, 0x9476eb5d, 0x57e36f45, 0xee5be5b6, 0xf68488dd, 0x2884c2d7, 0x05ac1ff1},
|
||||
{0x04173760, 0x0fc5b934, 0xda828f00, 0xe43272df, 0x2fad6e9c, 0x7e2ab5fe, 0x0a4995b3, 0x00e0a5eb},
|
||||
{0x42f8ef94, 0x6070024f, 0xe11a6161, 0xad187148, 0x9c8b0fa5, 0x3f046451, 0x87529cfa, 0x005282db}}};
|
||||
|
||||
static constexpr storage_array<omegas_count, limbs_count> omega_inv = {
|
||||
{{0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x00000011, 0x08000000},
|
||||
{0x0becc81e, 0xd59e99d9, 0x537cdf25, 0x3ad971a9, 0xbcd60738, 0xaccedf99, 0xd65d66b5, 0x01dafdc6},
|
||||
{0x4bc9ca34, 0xc8e6df6f, 0x5397aaca, 0xab8bfbc5, 0x94813e6e, 0xb5ea6773, 0xe295dda2, 0x0446ed3c},
|
||||
{0x8145aa75, 0xd7981c5b, 0x3d174c52, 0xb14011ea, 0xe4721c1e, 0x647c9ba3, 0x6f6ac6dd, 0x05c3ed0c},
|
||||
{0x6e0bef41, 0x9de8c5cf, 0xcee1b9b0, 0xec349cbb, 0x2121589c, 0xfe72ab05, 0x24c7669c, 0x03b1c96a},
|
||||
{0x246766d8, 0xb878549e, 0xb5a03ab4, 0x8c5d8531, 0x7f1ec75e, 0x334a83ab, 0x46b146d7, 0x01342b29},
|
||||
{0x31055652, 0x8c71bd50, 0x6081f8c3, 0x2eedac49, 0xab013740, 0x25164a76, 0xbca84bf7, 0x05c0a717},
|
||||
{0xd0a6b4f5, 0x1ad37af3, 0x8ca50294, 0x6dc49fe3, 0x5d9529c3, 0x8357a7ff, 0xcefe8efe, 0x02c161bc},
|
||||
{0x296fbf1c, 0x90a5fa7f, 0xc977b113, 0x18226a39, 0xc178262e, 0x9362d5c9, 0x40d28de5, 0x03a362d3},
|
||||
{0x125ca33a, 0x04eeb1c0, 0x8437c604, 0xaa47a4c0, 0xa4d6bafe, 0x064426a2, 0xb8cc76db, 0x00ffbb44},
|
||||
{0x179e2ebe, 0xecf0daf8, 0x2574403b, 0x942e643e, 0x6bf06f7c, 0x684d31aa, 0x244c675c, 0x003b2bde},
|
||||
{0xfeccfccc, 0x96bc19dc, 0x269130b4, 0xbb26f74e, 0xd511649f, 0x15d57a9f, 0x7dcde3c3, 0x02d852a4},
|
||||
{0x44ad0610, 0xb4a47f4c, 0x06fa1b55, 0xdc2f028f, 0xd25979ac, 0xd73ddcd4, 0x076e7f5d, 0x06ba7cbe},
|
||||
{0x349eea63, 0xb0f43dd2, 0x3e64660d, 0x5e64466c, 0xc3bb94ce, 0x7206f426, 0xed4327aa, 0x036cb7c6},
|
||||
{0xf248b36c, 0x6503e80b, 0xe36060ec, 0xb93dd56f, 0x95c2c067, 0x6d3b2763, 0x155023a7, 0x038e7d59},
|
||||
{0xcdf92351, 0x140437ad, 0x2a5ab630, 0xb7a6e1b4, 0xd48175a5, 0xaa80b742, 0xd4afae89, 0x06a50046},
|
||||
{0xaea51997, 0xe8cde2cd, 0x417e3754, 0x612806f6, 0xb940adf4, 0xe40a4a07, 0xa33929b2, 0x063f5efa},
|
||||
{0x0c07573f, 0x0c0926df, 0xd8d4bee3, 0xa84e9027, 0x6bcd79ea, 0xf3776dfa, 0x523f55a8, 0x043a8517},
|
||||
{0x66984d05, 0x5b7e4e45, 0xdb8c30c4, 0xb9381de7, 0xae86e4f6, 0xd7c15128, 0x809daae7, 0x0718f1ad},
|
||||
{0xc1eae1a6, 0xe4fb0a7d, 0xa90a0813, 0xe5484134, 0x895df525, 0x24cca8f9, 0x1cedd2ee, 0x035fd390},
|
||||
{0x82e87775, 0x0a87a942, 0x971f450b, 0x9f2b4b62, 0x8eae6f09, 0x1dc5aecd, 0x1c5686a6, 0x07547fa3},
|
||||
{0x2e35511a, 0x785975cc, 0xa085c456, 0x4266bc82, 0x3abd5bfd, 0x45cf52e1, 0x7bd95ece, 0x019e8e43},
|
||||
{0xae580194, 0xfad72a75, 0x2989ac16, 0xf2bb5a00, 0x55f2b4d0, 0x53fee728, 0x9c7a91e5, 0x02b9f95d},
|
||||
{0x71200963, 0xb0062d2c, 0x1ac57a23, 0xe16e9f91, 0xc4bd9d3e, 0xaae7b169, 0x7f505f35, 0x07462151},
|
||||
{0x57e31913, 0xcf7bd10e, 0x6a4d0ee4, 0x1a360a91, 0x31869e35, 0xb2ba4914, 0x18005db4, 0x07a62d5c},
|
||||
{0xb4344711, 0x431f11e2, 0x6192c47e, 0x0cc3049c, 0xeb9c1bc3, 0x375dff93, 0x42071ee8, 0x03a75790},
|
||||
{0x9ed81498, 0x4eb14251, 0x98b804ef, 0x5852dbc5, 0x56d7f20c, 0xe0c1be13, 0x20d69181, 0x023e7f68},
|
||||
{0xe34f2d55, 0xf2eeb9b5, 0x2aad6f84, 0x63459f16, 0xbe37dbea, 0xf12099e7, 0x11b1a0fd, 0x06e45493},
|
||||
{0x0d6c93ed, 0x63032f6a, 0x5a04829f, 0xd99cbcc8, 0x89608b5e, 0x80f20416, 0x9df329f4, 0x00bf4231},
|
||||
{0x2710f927, 0xc7fc3d1b, 0x90d8503e, 0xc72d19af, 0x9940e689, 0xa9dcd3b8, 0x2da77ac9, 0x06fd386e},
|
||||
{0x08b27bc2, 0xc800035f, 0x4dfacc03, 0xd98987cf, 0x1256e525, 0x24f8fdbf, 0x1f104273, 0x04c575f1},
|
||||
{0x256c604a, 0x68b16e90, 0x6eba097d, 0x7f51023a, 0x1aeba9c8, 0x52c7629c, 0x4809d8da, 0x0575e850},
|
||||
{0x4ac81249, 0x7439d2f9, 0x4fc31ff2, 0x351e4a62, 0xb3906ded, 0x68fb8313, 0x08507a35, 0x007d43d8},
|
||||
{0x98859a12, 0xa87902b8, 0x73af55b3, 0x2f0d13e0, 0x1b9783c2, 0x5a46c66a, 0x2f5f71d4, 0x01045b06},
|
||||
{0x604fce1e, 0x0c379595, 0x7fccc2b4, 0x20ab6eb8, 0xf1820ae7, 0xac0bc709, 0x93fb2b07, 0x07e7654f},
|
||||
{0x246c4bf0, 0xa0e40811, 0x816b15e0, 0xe12accf5, 0x17938138, 0xee417239, 0x2c9a34fb, 0x004e092e},
|
||||
{0xad2cd984, 0x6304351b, 0x4bf1aafc, 0x38546ca6, 0xf310e99f, 0x1fb81192, 0xb5376275, 0x07e89896},
|
||||
{0x7b2d141d, 0xe4376a0b, 0x6dac220c, 0xea1795e5, 0xb19e1901, 0xd778ab50, 0xa94c274f, 0x077df905},
|
||||
{0x16fcd6c7, 0x7039bab1, 0xa6ea1c94, 0x8eececb7, 0x0f122046, 0x84d26ab5, 0x22fd55a1, 0x053c5d48},
|
||||
{0x72f11f65, 0xd43eb7bb, 0xb2a566d6, 0xfb538785, 0x3f35cbf5, 0xccc2cdc6, 0x7112504a, 0x06df5a9e},
|
||||
{0x60ce9c30, 0x75efb55c, 0x3c541437, 0x991873ed, 0xdf0cbb3b, 0x37eaedcb, 0xb04c2858, 0x0278d7f0},
|
||||
{0x1a06866b, 0x5757dd4e, 0x6570fa7f, 0x15c176b1, 0xafe89a1d, 0x9981b57f, 0xee0cb14c, 0x03c57f4d},
|
||||
{0x503c31cd, 0x3438cd66, 0xc0736d4b, 0x34437e52, 0x2a9d1b28, 0xe825b769, 0x73c06ee7, 0x06955a3a},
|
||||
{0x5c5e530e, 0xbbf0995a, 0x6569a2f9, 0xdee304b3, 0x5bd1a886, 0x3b9c993c, 0xc9cd050a, 0x00f66017},
|
||||
{0xee755737, 0x3666e752, 0x74d0e317, 0xa13bfafc, 0x01d2f1bf, 0x17ab672a, 0x0778f525, 0x079dde3a},
|
||||
{0xed8a25e9, 0x96a003c2, 0x8f347cec, 0x45d258fe, 0x96ea14ac, 0x68ff148d, 0xe148eda9, 0x058f4ec7},
|
||||
{0xe2a700ab, 0x23baf732, 0x5202a945, 0x6434725a, 0x2e693363, 0xa19a338d, 0xbf2f39c6, 0x01d0ea7a},
|
||||
{0x3ab52589, 0x5e571cad, 0x92240361, 0xe2916bb2, 0xdff5e354, 0xe6f8897b, 0x2ffa4707, 0x02a62880},
|
||||
{0xef649a85, 0xaf446c62, 0xed4e461f, 0x14d8072f, 0x59993efa, 0x5a07f4e5, 0x72a3a652, 0x00dc28b6},
|
||||
{0xf21511df, 0x139299d7, 0x4854ebc3, 0x8914e707, 0xbfd102a9, 0x9f3b5913, 0x3a5af894, 0x009dc24f},
|
||||
{0x1f4ba4fa, 0x650e1d91, 0x1977bff0, 0x6ba67806, 0xaa9bbc1b, 0xffbdc531, 0x997408aa, 0x057b69b2},
|
||||
{0x65fb1a91, 0x25c03e81, 0x7fd22618, 0x8682f98b, 0xf46cb453, 0xcad67f13, 0x5a80e5c6, 0x060ca599},
|
||||
{0x94188f2a, 0xa7978a90, 0xdbb9338e, 0xd5fc8f0b, 0xcbdd84f0, 0xf8387e6d, 0xbbc743a3, 0x073ae131},
|
||||
{0x0415bbcc, 0xafd00c46, 0x0df4a52a, 0x1a00eb6c, 0x0b96b594, 0x1ec67c64, 0x8e26b699, 0x01cb82a5},
|
||||
{0x7f740f93, 0xf56319fb, 0x2e2f6ed7, 0xb40d559b, 0x75e19784, 0x63f96f04, 0xc31ba061, 0x06406929},
|
||||
{0xfa5a3239, 0x22349e8b, 0xb9ca6bf9, 0xe1236395, 0x9b0017a4, 0x76ae5a8b, 0x17b7af03, 0x06cfb4ce},
|
||||
{0xb51abfe6, 0x34938785, 0x1249edb6, 0x21f54c80, 0xab038972, 0x3bd1cc16, 0xa4a57a81, 0x0636b37f},
|
||||
{0xf88717cf, 0xfda4a9a1, 0xee19d402, 0xf8fcba35, 0x47c9ba1b, 0x1ac940f6, 0xdd991440, 0x013c0ab3},
|
||||
{0x3743adf4, 0x5082318a, 0x22440f94, 0x3293bae1, 0x8dd2d761, 0x4c2e6d7f, 0xcdc38c82, 0x07124118},
|
||||
{0x76198779, 0xb031f8b7, 0x1b6c1944, 0x6742f602, 0x894a6134, 0xa18290db, 0xaba037dc, 0x035289d8},
|
||||
{0x9f8a9b07, 0x4579e855, 0x4dca3764, 0x1e580662, 0xb8c8ef49, 0xda92152e, 0x8b54508a, 0x0444085a},
|
||||
{0x34696648, 0x7f670ce1, 0xc05768d9, 0x2f00108f, 0x390fb519, 0x2d00a444, 0x1cd6f914, 0x015c468b},
|
||||
{0xfe46c5f2, 0x00666cbf, 0x9f7174d6, 0xca4051c5, 0x8e4277f4, 0x1629882a, 0x6ee002a3, 0x00b3f261},
|
||||
{0xc1dbb4f6, 0x418a2b86, 0x9a6ca270, 0x9f453ccc, 0x1d457b20, 0x1966471f, 0x80fd1319, 0x00b4d831},
|
||||
{0x1c76c8b1, 0xa12f86a8, 0xc0125e48, 0x2772e424, 0x1459dfb8, 0x8d650644, 0xad06d01c, 0x02128e5c},
|
||||
{0x3472799c, 0xcc8cc7f6, 0x2f511cae, 0xfbd97f95, 0x5ebbff71, 0xadd8818b, 0x09af0983, 0x00520540},
|
||||
{0x8ec654cc, 0xcaab5dd4, 0x17ba15a9, 0xc05ad0a7, 0x36300a00, 0x4bda7469, 0x41bb0610, 0x02e486cd},
|
||||
{0x2d6be8b5, 0x077ba983, 0xfe89eb7d, 0xdd5e728f, 0x63f9c51f, 0xe3c872fb, 0xce639995, 0x01f2f7a8},
|
||||
{0xaa2ea7eb, 0xd82b1599, 0xa16489e0, 0x1be5d254, 0x173d3219, 0x19cb236a, 0x1fe63b23, 0x007dd45f},
|
||||
{0x19dba628, 0xa27cc4d3, 0x5fd2e061, 0xf04ac441, 0x9307a758, 0xc7405333, 0x28c40fe4, 0x0103c707},
|
||||
{0x54662aab, 0xb5129fd1, 0x59158f32, 0x2ec5b69b, 0x12c44eec, 0x6c7e6492, 0xe527abb2, 0x046e7c11},
|
||||
{0xe32d46fe, 0xb9bf4936, 0xb08ef006, 0xf23ae18c, 0xe6a5179e, 0x5352cc59, 0x5bf7c0b8, 0x0753a621},
|
||||
{0x9318db3a, 0x19f65bc2, 0x7e3d0014, 0x93ff3f79, 0x6beb580d, 0xf7f93c7f, 0xddd72603, 0x04fdb898},
|
||||
{0xe184a935, 0xf7e1f88f, 0x1ad510f0, 0x82a0f047, 0x4c9ab6ca, 0xce0f7c44, 0x5104a95a, 0x0552304e},
|
||||
{0x985bba5c, 0x06615580, 0xf487a1fb, 0x8ccd29a8, 0xeecf758d, 0xb3e15ed0, 0x857ce648, 0x05328783},
|
||||
{0x6cb042b0, 0x5d1d5a22, 0x0277083c, 0x64375cf4, 0x5fa82215, 0xe8947dab, 0x86932495, 0x05e72829},
|
||||
{0x8c3e2849, 0x5bf6f46a, 0x4924c8f4, 0x7e40314c, 0xdffd6118, 0x3c74a4ba, 0x2f8de20a, 0x05247cdd},
|
||||
{0xd0042d11, 0x25a418c5, 0x2f7da60c, 0x1b60ee9f, 0x02c0b69f, 0x61c041ad, 0x15670214, 0x0632d33a},
|
||||
{0x90e05a92, 0x32b03a5e, 0x78d1e8d6, 0xfb12a1b1, 0x5bc2f5d5, 0xb8af534e, 0xa032918a, 0x05ab4772},
|
||||
{0x0a711a9d, 0x096878a8, 0x6b083c8c, 0x87d070da, 0x87d06afb, 0x77931578, 0xf3104057, 0x03705277},
|
||||
{0xdf993e46, 0x502d2374, 0x35baf646, 0xc1cd2868, 0xe30aa213, 0xa61b54b6, 0xbce34b74, 0x02511017},
|
||||
{0x90a6b9b9, 0xcfb6c51a, 0x8be6ade8, 0x4e0b29ef, 0xd3832d74, 0xa8292467, 0x41ca1e45, 0x02ce7977},
|
||||
{0x3e672d5b, 0x25ee10aa, 0x28597504, 0xb0e60c63, 0xe263c827, 0x4a8d0567, 0xfadefeba, 0x01f4ec42},
|
||||
{0xa5a26158, 0x8b4b15e0, 0x88a71cf2, 0xa59b2df9, 0x5d734341, 0xde44f2e7, 0x4db8d2e8, 0x007a18a0},
|
||||
{0xb4d18100, 0x30fcf001, 0xf8ae0b4f, 0xcdaa5334, 0xe325615a, 0x67017b2b, 0xf0ccbf57, 0x016c6d47},
|
||||
{0xba937732, 0x66afc115, 0xc20be386, 0x917d4890, 0xa017c59d, 0x5dadccff, 0x986c39c1, 0x043fa44e},
|
||||
{0x08baa72a, 0xc57ec886, 0x052364ed, 0xe65a4680, 0x85f9a523, 0x0536b505, 0xfe744ee2, 0x03580609},
|
||||
{0x1bab1ab8, 0x88109415, 0x62f0fa74, 0x02244b19, 0x915618e0, 0x837fcd10, 0x942f12d2, 0x061b83d0},
|
||||
{0x687b7798, 0x823d0bba, 0x84a49784, 0x5f93174a, 0x2574af37, 0xcfd64159, 0xe108057c, 0x0290722e},
|
||||
{0x58a66036, 0x900a7031, 0x6153c2ae, 0xcb443378, 0xa6ccdffe, 0x4c48b8dd, 0xa06e955a, 0x049a9211},
|
||||
{0xea0b9dd9, 0x1b034532, 0x638c79ec, 0x11cba08f, 0x7c5b2d15, 0x16d00728, 0xbb9a759c, 0x05abcbcd},
|
||||
{0x1552d6af, 0x21b4f60e, 0xbed54865, 0x2f7ea9d2, 0x738befdb, 0x39378802, 0x97845360, 0x02adf76c},
|
||||
{0x4026bb92, 0x6e5eb2ca, 0xcbed5570, 0x18f3d8bf, 0xb655ac26, 0x2a5fc8cd, 0x3809a1c5, 0x0031cd25},
|
||||
{0x0ef5e011, 0x2d698950, 0xc018b82d, 0xc0668c45, 0xf520d325, 0xd180ff47, 0xa38122b1, 0x046714c7},
|
||||
{0x12df2cc7, 0x8dec8a4b, 0x963031f8, 0x5eb84a1b, 0x88525708, 0xb75ad701, 0x07df57bd, 0x02054a99},
|
||||
{0x82b2f616, 0xe0013d43, 0x7b385914, 0x2ad34c97, 0x11108f4b, 0xc9969223, 0x9c9fad59, 0x0183f639},
|
||||
{0x06b4dc38, 0xaca9dfbc, 0x962d5774, 0x85596bbc, 0x22f1cd7d, 0xd7023923, 0x2067b180, 0x04d3c939},
|
||||
{0xe4004173, 0x6d13e6ab, 0xaafe8726, 0x3495d095, 0x33dc3303, 0xa22d3e4a, 0x776d2e14, 0x0276dbb2},
|
||||
{0x68c539b6, 0xa03f83cb, 0x7b42a06e, 0xfd3fa839, 0xe8d45ac3, 0xea0f1f15, 0xa414b012, 0x061adb94},
|
||||
{0xb33fb188, 0xd22fc6e3, 0xf723dc18, 0xbebc7978, 0xf6c99f34, 0xa874b584, 0xf67ff454, 0x049beb53},
|
||||
{0x754bed16, 0x7c247948, 0xe50eac10, 0x4a84bcfb, 0xade97580, 0xc00d65df, 0xca79c5ae, 0x0763d73c},
|
||||
{0x7aadbe1a, 0x696e27af, 0x9d8e2a1f, 0x113535e0, 0x4c011766, 0x6953003f, 0xbb52558c, 0x0498a75f},
|
||||
{0x6e09cee7, 0xcf26e897, 0x299b63c7, 0x813a76f2, 0x0939904c, 0x67c02fa7, 0x7e0b9483, 0x045c41a9},
|
||||
{0x4af5adcc, 0xad979914, 0xc2c7c068, 0x7d9267f9, 0x21b4a0a7, 0xda4fa3f8, 0x3386c423, 0x03f4bcc9},
|
||||
{0xd1228595, 0xe5fcd634, 0x12fc8b7c, 0x5571b994, 0x244857f8, 0xd50dcd33, 0x263b93f0, 0x060dc1d6},
|
||||
{0xfee59c89, 0x7040a236, 0x78ceb168, 0x91a4301b, 0x19cdb36a, 0x973b55bd, 0x71008400, 0x06a1c58e},
|
||||
{0x6af1f351, 0x1d3c7ad7, 0xe8ad24dc, 0x8493c0c1, 0x48d5ffd9, 0x076f9dea, 0x5931555f, 0x00b9b2bf},
|
||||
{0xeaa5731c, 0xa3d54d89, 0xba84ee02, 0xfcc41a45, 0xcc1cdac8, 0x7c828f73, 0x5bfe9d23, 0x009c426b},
|
||||
{0x3f1f352c, 0x36fb314c, 0x9feb1120, 0x750a2a5f, 0xd7b06171, 0x3a2f19e8, 0x3b550cd9, 0x06de1885},
|
||||
{0xb69183f6, 0xefc03237, 0x979ee075, 0xb5a14fc3, 0x2dcb1d51, 0xbf114125, 0xb8eca2d3, 0x062364f7},
|
||||
{0x95375861, 0x575f1ea7, 0x80cc8dba, 0x30608586, 0xcf7a8f9f, 0x2beca9f5, 0x5fe60da4, 0x00dfc078},
|
||||
{0x0f86ded5, 0x312928eb, 0xb9c4f0cc, 0x646f5d3e, 0x2fbf14dd, 0x23c69382, 0xc44caa0e, 0x023aae90},
|
||||
{0x13e16243, 0xa7c92faf, 0x92efd5fc, 0x035a3e75, 0x86a744ea, 0x32f44d08, 0x1ea28333, 0x05b45217},
|
||||
{0xc41fdf22, 0xb557d203, 0x4bbc8f76, 0x9697570c, 0x81eaf742, 0x3a6a2cb5, 0xb0d03a0f, 0x07f2c08a},
|
||||
{0x2a18b73a, 0xca806385, 0xdb6a953d, 0xf2015d6d, 0xba5f67b9, 0x51d21a8e, 0x14807dd6, 0x051439d5},
|
||||
{0xf75051de, 0x7b6e0c13, 0x14dd1aa0, 0x114681fb, 0x0fd95a37, 0x72a1cccc, 0xa39e5bb8, 0x02f29d4c},
|
||||
{0x116529cd, 0x4808a0de, 0x5b941d1c, 0x1cf38580, 0xd70796f7, 0xc96a451e, 0x3f24e64f, 0x016d083f},
|
||||
{0x3cf155ee, 0xc71b78d0, 0x0c361b67, 0x0c04a134, 0x7756e4a9, 0xdb546edc, 0x2988eb2c, 0x03474404},
|
||||
{0xf30cef17, 0x1a0b3585, 0x864abd80, 0x63c1de29, 0xc0687c8e, 0x0c171d6e, 0xc9763a97, 0x0353aec8},
|
||||
{0x94192fb8, 0x0a2c9cff, 0x1a7f5bbf, 0x27320b93, 0xe5ceeb75, 0x465d2f9f, 0xd78f1cc3, 0x07ce6f99},
|
||||
{0xe8d1b26d, 0x0f899233, 0xb87a2984, 0xed4b44d2, 0x0bd6354a, 0x0c0712c6, 0xc7032f5c, 0x01eb2a31},
|
||||
{0x46b03b57, 0xc4c03fbd, 0x785ebbe8, 0x989b0ff3, 0x7f0bcb19, 0x5cada62a, 0xa97557c9, 0x01426410},
|
||||
{0x96fb0a26, 0xf1d2e82b, 0x1edb9ce3, 0xe270bc10, 0xfc7aaed8, 0x9549cfd0, 0xd90d7c9c, 0x03e8256c},
|
||||
{0x43ac9984, 0x14eef0ee, 0xa16d6770, 0x2903ff22, 0xa38fbfc0, 0xc66c2690, 0x8755440e, 0x0032a202},
|
||||
{0xf3601782, 0x46a07cf2, 0xaa71d137, 0x79f410f9, 0x8bcabc59, 0xc320c6f1, 0xf8ab64d8, 0x00a706cf},
|
||||
{0x8dbd8d4f, 0x8848a9f0, 0x0085061d, 0xeff89e69, 0xfee62fbe, 0x90e634a7, 0x2ffb456b, 0x03983046},
|
||||
{0xb272ed5c, 0x91ec28a8, 0xdc0cbb77, 0xf8529918, 0x3648d2c5, 0x8f896ddb, 0x74edaf19, 0x0668a86c},
|
||||
{0x128c9bd9, 0x341d5fc8, 0x6b3241c5, 0x592f87d8, 0xb2cc3c97, 0xf8cba6f2, 0x03f396ed, 0x03463bf1},
|
||||
{0xafd9d239, 0xcf3ae525, 0xea20b753, 0x06b8b7b9, 0x3408a993, 0xb2be1e49, 0x9f47063f, 0x02bcb200},
|
||||
{0xa0bd0bc8, 0x7ca02722, 0xb862774d, 0xce8b32ee, 0x5f8da059, 0x424ba5f0, 0x3bb422a0, 0x05c81961},
|
||||
{0x32fd8907, 0x137dad8c, 0xc95a3a5d, 0x301d5119, 0x8937ac08, 0x144b38c3, 0x39338de7, 0x00e66f0e},
|
||||
{0xcfc10885, 0xe68b8875, 0x96147e68, 0x4f24d49a, 0x43032c15, 0x5da9e6fd, 0x9bf25e12, 0x061ab0e6},
|
||||
{0x455c65ad, 0xeab29bbd, 0x2448be64, 0x1c7da0e7, 0x8eedfa1f, 0x8c2c1bcd, 0x698c1197, 0x0400e2d2},
|
||||
{0x04549c13, 0x335d3e9e, 0xd31585cc, 0x546f0d82, 0xe16dbbac, 0x350d5ed5, 0x113c53fd, 0x05f77544},
|
||||
{0x7d8f3b7e, 0x6aa75c04, 0x10a641ae, 0xc70851dd, 0x9a0750fe, 0x4d33edd4, 0xcd1b230f, 0x022802cf},
|
||||
{0xef8170e3, 0x59fa1903, 0x62995788, 0x464a73ef, 0x13369717, 0x338be7fd, 0x52d21278, 0x02e97589},
|
||||
{0x4856ddd5, 0x3f2deca8, 0xfced10e2, 0x969b10e2, 0x52860ee7, 0x09620dde, 0xb620fa3f, 0x04a169bf},
|
||||
{0xa03b49f1, 0xd9beb712, 0xe9af606e, 0x0798af09, 0x63e70b9a, 0xe37f9aea, 0xb35abd7c, 0x02542a44},
|
||||
{0xf6e78973, 0x335d4000, 0x76f1bb23, 0x7bc28fde, 0x1b30e9ca, 0x6cfdc907, 0x0400b651, 0x03ff88aa},
|
||||
{0x36433eaf, 0xfb862981, 0x4111cfa3, 0x15fdc659, 0xeab2909d, 0x569574b9, 0x3cd80f84, 0x01442360},
|
||||
{0xe85c4af3, 0xa8ed8f31, 0xe6aaf3da, 0xf7680fee, 0xc5c1772c, 0x2240e931, 0xaebeeb70, 0x04f44f6f},
|
||||
{0x8846e0af, 0x29de323f, 0x42c25319, 0x33f91593, 0x6cbadd58, 0x863099c1, 0xfd83e5b3, 0x06a603cf},
|
||||
{0x86c77703, 0x1bdd17f3, 0xe02db671, 0x8cee8e78, 0x0b6dffce, 0xed1627af, 0xa0d9b3cc, 0x04491984},
|
||||
{0xcb583661, 0x177f8f9c, 0x73d05bfc, 0x54122d0c, 0xebe37b4a, 0xa9231660, 0xd4826038, 0x06e885db},
|
||||
{0x13c253b9, 0x64cde875, 0x2fbc98a9, 0x8484bccb, 0x4885a9af, 0xbad877c5, 0x0cbc33b6, 0x03007c90},
|
||||
{0x47cfa357, 0x41eb9173, 0x325309ad, 0xb3f06289, 0xaa85421b, 0x029da7c1, 0x84de4bd4, 0x07b7eb0d},
|
||||
{0x56b831e2, 0x2c459a80, 0x321aba19, 0x2b99d098, 0xea73c0e1, 0x96237364, 0xe25ed0ed, 0x02f2c638},
|
||||
{0x9b388bf4, 0xfc8c3228, 0x82cd081d, 0xa4c371e4, 0xc85f75df, 0x11239026, 0x8892896e, 0x01f01c5e},
|
||||
{0x73457917, 0xce1dde59, 0x16dd8b49, 0xdfdaeb19, 0xbfd17b1e, 0x4289a976, 0xc842870a, 0x05e2cf7e},
|
||||
{0xc7705532, 0x72faa825, 0x8f7fe8c2, 0xd24bf942, 0xb695e31b, 0xb7403e13, 0xfc85a0c6, 0x02eac9e7},
|
||||
{0x1ddb2dff, 0xc47638e3, 0x799bb649, 0x78b91a13, 0x552588ed, 0x001800de, 0x9cd9425c, 0x01d0640c},
|
||||
{0xfb431e10, 0x159891e7, 0xa012b461, 0x2f2fb29a, 0xb3333e5d, 0xc1dca804, 0x9a47200d, 0x05b918ec},
|
||||
{0x2d5ce760, 0x379119b5, 0xda2ccdab, 0xf9911f75, 0x47b5c054, 0x92b09490, 0x7298d065, 0x0742a31e},
|
||||
{0x4a73d1f1, 0xe2a1046b, 0xc6ab4d9c, 0xbc85a747, 0xba0701f8, 0x79b0e699, 0xeebc6762, 0x05e5c2cb},
|
||||
{0xe0c0db50, 0xdc644b37, 0x2b8444d2, 0x26f7f083, 0x63479a84, 0x90acf2e7, 0x90ffe372, 0x0590d880},
|
||||
{0x83c0fc9c, 0x3dd1aba4, 0xcfb43020, 0x30a1051f, 0xaf5be716, 0x7d1ca380, 0x1ed8aed9, 0x01d56947},
|
||||
{0x0fa23690, 0x657df8c4, 0x32111be3, 0x61a12fe4, 0xe78236c9, 0xd6cc9942, 0x85e66191, 0x01709635},
|
||||
{0xc6a054f0, 0x96bf35ed, 0x004113cc, 0x9d1e411a, 0x1ac7a3ec, 0xccdb9bc3, 0xd08016b8, 0x07362425},
|
||||
{0x9721b035, 0x72744cce, 0x0beb72e3, 0xb87eb606, 0x60870c2e, 0x00c5e70c, 0x685d7c14, 0x029fa4d3},
|
||||
{0x86e52af4, 0x06d3a7a3, 0x70020878, 0x7b1c814a, 0x52e68007, 0x44373cb7, 0xe403540f, 0x041cf8c0},
|
||||
{0x76a27949, 0xd5dbc8bf, 0x27d9cd12, 0xb41449bc, 0xa7a667a1, 0x93740020, 0x0fbb4e77, 0x000bf807},
|
||||
{0x9969cfe9, 0x274ce281, 0x259ec27c, 0x3234d283, 0xe0b44f04, 0x9ff85b71, 0xffcc1006, 0x0298d060},
|
||||
{0x68ab54f8, 0x5cd8b289, 0x437eaab8, 0x42e3877f, 0x9318bd3e, 0x6490dc61, 0x4e54d968, 0x075b01f3},
|
||||
{0x7b64243c, 0x73100d65, 0x5c802f82, 0x692378be, 0x88184c0c, 0x00283dbb, 0xab6f4f0e, 0x0442efad},
|
||||
{0x72015722, 0xbe83b708, 0xe1cdcf0e, 0x2035319f, 0x398347da, 0x2b1b3351, 0x1a14b8dc, 0x061823d8},
|
||||
{0x378d9803, 0x1090948c, 0x4725c64b, 0x61a558cc, 0x7d7fcd91, 0x9e5bd3b5, 0x57ebda25, 0x061e02a0},
|
||||
{0xf8324dc8, 0x166b4a3c, 0x38133fda, 0xa25b9d11, 0x917171a5, 0x9d602950, 0x417d104e, 0x0632e48b},
|
||||
{0x6a61d5e0, 0x03b9f1b9, 0xe59cfbb7, 0xd906b740, 0x7892fbe4, 0x99a93267, 0xad1b8171, 0x06ddc2a6},
|
||||
{0x67fc3874, 0x6ae4355d, 0xb1ada695, 0x4fa456d8, 0x9f91ac43, 0x4e234065, 0x829d173e, 0x028da309},
|
||||
{0xfc695c2c, 0x1e08dd18, 0xfa687112, 0x1c0a2fad, 0xffd6302a, 0xeb5ebf01, 0xfd1d10f5, 0x012fd387},
|
||||
{0x236e65c9, 0x0b907f2e, 0xb1281d54, 0x92ba7a15, 0xc13f1d75, 0x07f0a6ad, 0xcd6d1e9c, 0x05dfe4e3},
|
||||
{0xc45f33f8, 0xd99cc41a, 0xd373165c, 0xc1c10a71, 0x2ce2936a, 0x6c809230, 0xa0498cf5, 0x018dc832},
|
||||
{0x7b222ad8, 0x8e881eab, 0xb6194efb, 0xc8b48774, 0x963c6b6b, 0x38452dfd, 0xe4c4e0f8, 0x02847f5a},
|
||||
{0x2bf4ad95, 0x2950bb4a, 0xdc39ffb0, 0x37f42c9b, 0x101253a8, 0x3814fa42, 0xb67f2ca5, 0x04d4a34c},
|
||||
{0xa9684ba0, 0x6c40fece, 0x3b13bca4, 0xc7108aad, 0xe7bff9be, 0x98ccc7ea, 0xe9b3b316, 0x048b3a6a},
|
||||
{0x08390a2b, 0x4d908260, 0x74b070bc, 0xd5a641d0, 0x910015c5, 0xc3b19274, 0xd5a998a7, 0x02ac8e74},
|
||||
{0x9698d605, 0x8de03acc, 0xa4c9137f, 0x3b8b720c, 0x354faf46, 0x5bbad6e4, 0xfd9e842f, 0x0054c120},
|
||||
{0xd65aead5, 0x305fa33f, 0x0fe296f9, 0xba02b164, 0x708efc94, 0x64cba43c, 0x8ad7f0ef, 0x034b9ffe},
|
||||
{0x13c2e8f4, 0x59e1179e, 0xc572f8a8, 0x5d823d59, 0x74003bce, 0x0cfdb6ee, 0x011c179e, 0x00763941},
|
||||
{0xa47999a8, 0x29b692ee, 0xbfcd80d8, 0x6436c3f1, 0x959768d7, 0x553444f3, 0x583896d4, 0x01d45a26},
|
||||
{0xc150b3f8, 0x0ce0791d, 0xf493c135, 0x7d3a0c1f, 0x5ede0712, 0x4d37cc23, 0x34fbae9c, 0x036a6a38},
|
||||
{0x2ca1eb78, 0xa8ee8204, 0x66d8b759, 0xc713a1dc, 0xac061800, 0x1813508d, 0x3b1f0da2, 0x05725ca0},
|
||||
{0xf2f391c1, 0xbe6826df, 0x232878f0, 0xeb85b046, 0xf7e1d662, 0xf5a96510, 0xe38c2b64, 0x0419a43b},
|
||||
{0xe69e791b, 0x4b54889b, 0xb5c95ea5, 0xb371eeb0, 0x0b2f26a3, 0x9f53ccca, 0x66f45f71, 0x0040592d},
|
||||
{0xad2e5d5b, 0x4ced12db, 0x0987b849, 0x5f57b16d, 0xd9ec045b, 0xcab0e2e9, 0x6cfbf4df, 0x03e4e405},
|
||||
{0x3ecb72a4, 0xd71a1eee, 0x03a13fb7, 0x6bd9f7ec, 0x5877c6c7, 0xb74a54c8, 0xa28236a5, 0x0377689b},
|
||||
{0x74b3354c, 0x6f558a20, 0x3f776b18, 0xb67f6d10, 0x01165ed8, 0x8c447df2, 0xf3889308, 0x056b8991},
|
||||
{0x0d306b7a, 0x9482eb10, 0xd441cd03, 0xdd738e0f, 0x2de5dfd7, 0x6d186de5, 0x75fd1833, 0x00781b3e},
|
||||
{0x77ec28e5, 0xdbc14748, 0xd26e050c, 0x02ceee41, 0x18457c96, 0x8e5aef74, 0x1823c60f, 0x0461a6e2},
|
||||
{0x2be17c8b, 0x172e551d, 0x49c6a7b8, 0x90e25fa2, 0xa1b3478f, 0x6219e63e, 0xd063a517, 0x00c412f8},
|
||||
{0x65a9b68e, 0xb136b848, 0x673c6cbc, 0x9a9b7169, 0xf8ec7473, 0x15fa1875, 0x3033a5d6, 0x022d72f6}}};
|
||||
|
||||
static constexpr storage_array<omegas_count, limbs_count> inv = {
|
||||
{{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x80000000, 0x00000008, 0x04000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xc0000000, 0x0000000c, 0x06000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xe0000000, 0x0000000e, 0x07000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xf0000000, 0x0000000f, 0x07800000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0x78000000, 0x00000010, 0x07c00000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xbc000000, 0x00000010, 0x07e00000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xde000000, 0x00000010, 0x07f00000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xef000000, 0x00000010, 0x07f80000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xf7800000, 0x00000010, 0x07fc0000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfbc00000, 0x00000010, 0x07fe0000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfde00000, 0x00000010, 0x07ff0000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfef00000, 0x00000010, 0x07ff8000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xff780000, 0x00000010, 0x07ffc000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffbc0000, 0x00000010, 0x07ffe000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffde0000, 0x00000010, 0x07fff000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffef0000, 0x00000010, 0x07fff800},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfff78000, 0x00000010, 0x07fffc00},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffbc000, 0x00000010, 0x07fffe00},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffde000, 0x00000010, 0x07ffff00},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffef000, 0x00000010, 0x07ffff80},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffff7800, 0x00000010, 0x07ffffc0},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffbc00, 0x00000010, 0x07ffffe0},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffde00, 0x00000010, 0x07fffff0},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffef00, 0x00000010, 0x07fffff8},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffff780, 0x00000010, 0x07fffffc},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffffbc0, 0x00000010, 0x07fffffe},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffffde0, 0x00000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xfffffef0, 0x80000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffff78, 0xc0000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffffbc, 0xe0000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffffde, 0xf0000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x00000000, 0xffffffef, 0xf8000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x80000000, 0xfffffff7, 0xfc000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xc0000000, 0xfffffffb, 0xfe000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xe0000000, 0xfffffffd, 0xff000010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xf0000000, 0xfffffffe, 0xff800010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0x78000000, 0xffffffff, 0xffc00010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xbc000000, 0xffffffff, 0xffe00010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xde000000, 0xffffffff, 0xfff00010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xef000000, 0xffffffff, 0xfff80010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xf7800000, 0xffffffff, 0xfffc0010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfbc00000, 0xffffffff, 0xfffe0010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfde00000, 0xffffffff, 0xffff0010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfef00000, 0xffffffff, 0xffff8010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xff780000, 0xffffffff, 0xffffc010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffbc0000, 0xffffffff, 0xffffe010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffde0000, 0xffffffff, 0xfffff010, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffef0000, 0xffffffff, 0xfffff810, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfff78000, 0xffffffff, 0xfffffc10, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffbc000, 0xffffffff, 0xfffffe10, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffde000, 0xffffffff, 0xffffff10, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffef000, 0xffffffff, 0xffffff90, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffff7800, 0xffffffff, 0xffffffd0, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffbc00, 0xffffffff, 0xfffffff0, 0x07ffffff},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffde00, 0xffffffff, 0x00000000, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffef00, 0xffffffff, 0x00000008, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffff780, 0xffffffff, 0x0000000c, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffffbc0, 0xffffffff, 0x0000000e, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffffde0, 0xffffffff, 0x0000000f, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xfffffef0, 0x7fffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffff78, 0xbfffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffffbc, 0xdfffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffffde, 0xefffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x00000000, 0xffffffef, 0xf7ffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x80000000, 0xfffffff7, 0xfbffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xc0000000, 0xfffffffb, 0xfdffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xe0000000, 0xfffffffd, 0xfeffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xf0000000, 0xfffffffe, 0xff7fffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0x78000000, 0xffffffff, 0xffbfffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xbc000000, 0xffffffff, 0xffdfffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xde000000, 0xffffffff, 0xffefffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xef000000, 0xffffffff, 0xfff7ffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xf7800000, 0xffffffff, 0xfffbffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfbc00000, 0xffffffff, 0xfffdffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfde00000, 0xffffffff, 0xfffeffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfef00000, 0xffffffff, 0xffff7fff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xff780000, 0xffffffff, 0xffffbfff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffbc0000, 0xffffffff, 0xffffdfff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffde0000, 0xffffffff, 0xffffefff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffef0000, 0xffffffff, 0xfffff7ff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfff78000, 0xffffffff, 0xfffffbff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffbc000, 0xffffffff, 0xfffffdff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffde000, 0xffffffff, 0xfffffeff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffef000, 0xffffffff, 0xffffff7f, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffff7800, 0xffffffff, 0xffffffbf, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffbc00, 0xffffffff, 0xffffffdf, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffde00, 0xffffffff, 0xffffffef, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffef00, 0xffffffff, 0xfffffff7, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffff780, 0xffffffff, 0xfffffffb, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffffbc0, 0xffffffff, 0xfffffffd, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffffde0, 0xffffffff, 0xfffffffe, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xfffffef0, 0x7fffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffff78, 0xbfffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffffbc, 0xdfffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffffde, 0xefffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x00000000, 0xffffffef, 0xf7ffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x80000000, 0xfffffff7, 0xfbffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xc0000000, 0xfffffffb, 0xfdffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xe0000000, 0xfffffffd, 0xfeffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xf0000000, 0xfffffffe, 0xff7fffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0x78000000, 0xffffffff, 0xffbfffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xbc000000, 0xffffffff, 0xffdfffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xde000000, 0xffffffff, 0xffefffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xef000000, 0xffffffff, 0xfff7ffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xf7800000, 0xffffffff, 0xfffbffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfbc00000, 0xffffffff, 0xfffdffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfde00000, 0xffffffff, 0xfffeffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfef00000, 0xffffffff, 0xffff7fff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xff780000, 0xffffffff, 0xffffbfff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffbc0000, 0xffffffff, 0xffffdfff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffde0000, 0xffffffff, 0xffffefff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffef0000, 0xffffffff, 0xfffff7ff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfff78000, 0xffffffff, 0xfffffbff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffbc000, 0xffffffff, 0xfffffdff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffde000, 0xffffffff, 0xfffffeff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffef000, 0xffffffff, 0xffffff7f, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffff7800, 0xffffffff, 0xffffffbf, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffbc00, 0xffffffff, 0xffffffdf, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffde00, 0xffffffff, 0xffffffef, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffef00, 0xffffffff, 0xfffffff7, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffff780, 0xffffffff, 0xfffffffb, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffffbc0, 0xffffffff, 0xfffffffd, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffffde0, 0xffffffff, 0xfffffffe, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xfffffef0, 0x7fffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffff78, 0xbfffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffffbc, 0xdfffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffffde, 0xefffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x00000000, 0xffffffef, 0xf7ffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x80000000, 0xfffffff7, 0xfbffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xc0000000, 0xfffffffb, 0xfdffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xe0000000, 0xfffffffd, 0xfeffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xf0000000, 0xfffffffe, 0xff7fffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0x78000000, 0xffffffff, 0xffbfffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xbc000000, 0xffffffff, 0xffdfffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xde000000, 0xffffffff, 0xffefffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xef000000, 0xffffffff, 0xfff7ffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xf7800000, 0xffffffff, 0xfffbffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfbc00000, 0xffffffff, 0xfffdffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfde00000, 0xffffffff, 0xfffeffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfef00000, 0xffffffff, 0xffff7fff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xff780000, 0xffffffff, 0xffffbfff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffbc0000, 0xffffffff, 0xffffdfff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffde0000, 0xffffffff, 0xffffefff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffef0000, 0xffffffff, 0xfffff7ff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfff78000, 0xffffffff, 0xfffffbff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffbc000, 0xffffffff, 0xfffffdff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffde000, 0xffffffff, 0xfffffeff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffef000, 0xffffffff, 0xffffff7f, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffff7800, 0xffffffff, 0xffffffbf, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffbc00, 0xffffffff, 0xffffffdf, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffde00, 0xffffffff, 0xffffffef, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffef00, 0xffffffff, 0xfffffff7, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffff780, 0xffffffff, 0xfffffffb, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffffbc0, 0xffffffff, 0xfffffffd, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffffde0, 0xffffffff, 0xfffffffe, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xfffffef0, 0x7fffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffff78, 0xbfffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffffbc, 0xdfffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffffde, 0xefffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x00000001, 0xffffffef, 0xf7ffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x80000001, 0xfffffff7, 0xfbffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xc0000001, 0xfffffffb, 0xfdffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xe0000001, 0xfffffffd, 0xfeffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xf0000001, 0xfffffffe, 0xff7fffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0x78000001, 0xffffffff, 0xffbfffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xbc000001, 0xffffffff, 0xffdfffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xde000001, 0xffffffff, 0xffefffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xef000001, 0xffffffff, 0xfff7ffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xf7800001, 0xffffffff, 0xfffbffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfbc00001, 0xffffffff, 0xfffdffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfde00001, 0xffffffff, 0xfffeffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfef00001, 0xffffffff, 0xffff7fff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xff780001, 0xffffffff, 0xffffbfff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffbc0001, 0xffffffff, 0xffffdfff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffde0001, 0xffffffff, 0xffffefff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffef0001, 0xffffffff, 0xfffff7ff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfff78001, 0xffffffff, 0xfffffbff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffbc001, 0xffffffff, 0xfffffdff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffde001, 0xffffffff, 0xfffffeff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffef001, 0xffffffff, 0xffffff7f, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffff7801, 0xffffffff, 0xffffffbf, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffbc01, 0xffffffff, 0xffffffdf, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffde01, 0xffffffff, 0xffffffef, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffef01, 0xffffffff, 0xfffffff7, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffff781, 0xffffffff, 0xfffffffb, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffffbc1, 0xffffffff, 0xfffffffd, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffffde1, 0xffffffff, 0xfffffffe, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffffef1, 0x7fffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffff79, 0xbfffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffffbd, 0xdfffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xffffffdf, 0xefffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000},
|
||||
{0xfffffff0, 0xf7ffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0xffffffff, 0x00000010, 0x08000000}}};
|
||||
};
|
||||
|
||||
/**
|
||||
* Scalar field. Is always a prime field.
|
||||
*/
|
||||
typedef Field<fp_config> scalar_t;
|
||||
} // namespace stark252
|
||||
@@ -50,7 +50,7 @@ namespace keccak {
|
||||
*/
|
||||
template <int C, int D>
|
||||
cudaError_t
|
||||
keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config);
|
||||
keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config);
|
||||
} // namespace keccak
|
||||
|
||||
#endif
|
||||
@@ -79,6 +79,7 @@ namespace msm {
|
||||
* non-blocking and you'd need to synchronize it explicitly by running
|
||||
* `cudaStreamSynchronize` or `cudaDeviceSynchronize`. If set to false, the MSM
|
||||
* function will block the current CPU thread. */
|
||||
bool segments_reduction;
|
||||
};
|
||||
|
||||
/**
|
||||
@@ -103,6 +104,7 @@ namespace msm {
|
||||
false, // are_results_on_device
|
||||
false, // is_big_triangle
|
||||
false, // is_async
|
||||
false, // segments_reduction
|
||||
};
|
||||
return config;
|
||||
}
|
||||
|
||||
@@ -32,6 +32,7 @@ namespace mxntt {
|
||||
S* external_twiddles,
|
||||
S* internal_twiddles,
|
||||
S* basic_twiddles,
|
||||
S* linear_twiddle, // twiddles organized as [1,w,w^2,...] for coset-eval in fast-tw mode
|
||||
int ntt_size,
|
||||
int max_logn,
|
||||
int batch_size,
|
||||
|
||||
@@ -224,7 +224,7 @@ namespace keccak {
|
||||
|
||||
template <int C, int D>
|
||||
cudaError_t
|
||||
keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config)
|
||||
keccak_hash(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config)
|
||||
{
|
||||
CHK_INIT_IF_RETURN();
|
||||
cudaStream_t& stream = config.ctx.stream;
|
||||
@@ -245,7 +245,7 @@ namespace keccak {
|
||||
CHK_IF_RETURN(cudaMallocAsync(&output_device, number_of_blocks * (D / 8), stream));
|
||||
}
|
||||
|
||||
int number_of_threads = 1024;
|
||||
int number_of_threads = 512;
|
||||
int number_of_gpu_blocks = (number_of_blocks - 1) / number_of_threads + 1;
|
||||
keccak_hash_blocks<C, D><<<number_of_gpu_blocks, number_of_threads, 0, stream>>>(
|
||||
input_device, input_block_size, number_of_blocks, output_device);
|
||||
@@ -262,13 +262,13 @@ namespace keccak {
|
||||
}
|
||||
|
||||
extern "C" cudaError_t
|
||||
keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config)
|
||||
keccak256_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config)
|
||||
{
|
||||
return keccak_hash<512, 256>(input, input_block_size, number_of_blocks, output, config);
|
||||
}
|
||||
|
||||
extern "C" cudaError_t
|
||||
keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig config)
|
||||
keccak512_cuda(uint8_t* input, int input_block_size, int number_of_blocks, uint8_t* output, KeccakConfig& config)
|
||||
{
|
||||
return keccak_hash<1024, 512>(input, input_block_size, number_of_blocks, output, config);
|
||||
}
|
||||
|
||||
@@ -1,4 +1,4 @@
|
||||
#include "utils/device_context.cuh"
|
||||
#include "gpu-utils/device_context.cuh"
|
||||
#include "keccak.cu"
|
||||
|
||||
// #define DEBUG
|
||||
@@ -51,7 +51,7 @@ int main(int argc, char* argv[])
|
||||
|
||||
START_TIMER(keccak_timer);
|
||||
KeccakConfig config = default_keccak_config();
|
||||
keccak256(in_ptr, input_block_size, number_of_blocks, out_ptr, config);
|
||||
keccak256_cuda(in_ptr, input_block_size, number_of_blocks, out_ptr, config);
|
||||
END_TIMER(keccak_timer, "Keccak")
|
||||
|
||||
for (int i = 0; i < number_of_blocks; i++) {
|
||||
|
||||
@@ -1,4 +1,8 @@
|
||||
build_msm:
|
||||
mkdir -p work
|
||||
nvcc -o work/test_msm -std=c++17 -arch=sm_80 -I. -I../../include tests/msm_test.cu
|
||||
|
||||
test_msm:
|
||||
mkdir -p work
|
||||
nvcc -o work/test_msm -std=c++17 -I. -I../../include tests/msm_test.cu
|
||||
work/test_msm
|
||||
nvcc -o work/test_msm -std=c++17 -arch=sm_80 -I. -I../../include tests/msm_test.cu
|
||||
work/test_msm
|
||||
@@ -16,6 +16,132 @@
|
||||
#include "gpu-utils/error_handler.cuh"
|
||||
#include "utils/mont.cuh"
|
||||
|
||||
typedef uint32_t Field1[12];
|
||||
template<class Field>
|
||||
class PointXYZ {
|
||||
public:
|
||||
Field x, y, z;
|
||||
};
|
||||
|
||||
__device__ __forceinline__ void setZero(uint32_t* r) {
|
||||
#pragma unroll
|
||||
for(int32_t i=0;i<12;i++)
|
||||
r[i]=0;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void initializeXYZ(PointXYZ<Field1>& acc) {
|
||||
setZero(acc.x);
|
||||
setZero(acc.y);
|
||||
setZero(acc.z);
|
||||
}
|
||||
|
||||
|
||||
template <typename FF>
|
||||
__device__ __forceinline__ void accumulate_xyz(FF &x1,FF &y1,FF &z1, const FF& x2, const FF& y2, const FF& z2)
|
||||
{
|
||||
// FF B_VALUE{0x00000001, 0x00000000, 0x00000000, 0x00000000,
|
||||
// 0x00000000, 0x00000000, 0x00000000, 0x00000000,
|
||||
// 0x00000000, 0x00000000, 0x00000000, 0x00000000};
|
||||
FF T1,T2,T3,T4,T5,T6;
|
||||
typename FF::Wide W1,W2;
|
||||
T1 = x1 + y1; // t03 ← X1 + Y1 < 4
|
||||
T2 = x2 + y2; // t04 ← X2 + Y2 < 4
|
||||
T1 = T1 * T2; // t03 ← t03 · t04 < 3
|
||||
T2 = y1 + z1; // t08 ← Y1 + Z1 < 4
|
||||
T3 = y2 + z2; // t09 ← Y2 + Z2 < 4
|
||||
T2 = T2 * T3; // t10 ← t08 · t09 < 3
|
||||
|
||||
x1 = x1 * x2; // t00 ← X1 · X2 < 2
|
||||
|
||||
T3 = x1 + z1; // t13 ← X1 + Z1 < 4
|
||||
T4 = x2 + z2; // t14 ← X2 + Z2 < 4
|
||||
T3 = T3 * T4; // t15 ← t13 · t14 < 3
|
||||
|
||||
|
||||
y1 = y1 * y2; // t01 ← Y1 · Y2 < 2
|
||||
z1 = z1 * z2; // t02 ← Z1 · Z2 < 2
|
||||
|
||||
T4 = x1 + y1; // t06 ← t00 + t01 < 4
|
||||
T5 = x1 + z1; // t16 ← t00 + t02 < 4
|
||||
T6 = y1 + z1; // t11 ← t01 + t02 < 4
|
||||
|
||||
|
||||
T1 = T1 - T4; // t05 ← t05 − t06 < 2
|
||||
T2 = T2 - T5; // t12 ← t10 − t11 < 2
|
||||
T3 = T3 - T6; // t17 ← t15 − t16 < 2
|
||||
// T3 = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(T3)); // t23 ← b3 · t17 < 2
|
||||
|
||||
T4 = y1 + z1; // t21 ← t01 + t20 < 2
|
||||
T5 = y1 - z1; // t22 ← t01 − t20 < 2
|
||||
x1 = x1 + x1; // t18 ← t00 + t00 < 2
|
||||
T6 = x1 + x1; // t19 ← t18 + t00 < 2
|
||||
// z1 = FF::template mul_unsigned<3>(FF::template mul_const<B_VALUE>(z1)); // t20 ← b3 · t02 < 2
|
||||
|
||||
W1 = FF::mul_wide(T1, T5); // t24 ← t12 · t23 < 2
|
||||
W2 = FF::mul_wide(T2, T3); // t25 ← t07 · t22 < 2
|
||||
x1 = FF::reduce(W1 - W2); // X3 ← t25 − t24 < 2
|
||||
W1 = FF::mul_wide(T4, T5); // t27 ← t23 · t19 < 2
|
||||
W2 = FF::mul_wide(T2, T6); // t28 ← t22 · t21 < 2
|
||||
y1 = FF::reduce(W1 + W2); // Y3 ← t28 + t27 < 2
|
||||
W1 = FF::mul_wide(T3, T4); // t30 ← t19 · t07 < 2
|
||||
W2 = FF::mul_wide(T1, T6); // t31 ← t21 · t12 < 2
|
||||
z1 = FF::reduce(W1 + W2); // Z3 ← t31 + t30 < 2
|
||||
// return {X3, Y3, Z3};
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void setN(Field1& r) {
|
||||
r[0]=0x00000001; r[1]=0x8508C000; r[2]=0x30000000; r[3]=0x170B5D44;
|
||||
r[4]=0xBA094800; r[5]=0x1EF3622F; r[6]=0x00F5138F; r[7]=0x1A22D9F3;
|
||||
r[8]=0x6CA1493B; r[9]=0xC63B05C0; r[10]=0x17C510EA; r[11]=0x01AE3A46;
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void modmul2(Field1& r, const Field1& a, const Field1& b)
|
||||
{
|
||||
Field1 localN; //={0x00000001, 0x8508c000, 0x30000000, 0x170b5d44,
|
||||
// 0xba094800, 0x1ef3622f, 0x00f5138f, 0x1a22d9f3,
|
||||
// 0x6ca1493b, 0xc63b05c0, 0x17c510ea, 0x01ae3a46};
|
||||
uint64_t evenOdd[12];
|
||||
bool carry;
|
||||
|
||||
setN(localN);
|
||||
carry=mp_mul_red_cl<12>(evenOdd, a, b, localN);
|
||||
mp_merge_cl<12>(r, evenOdd, carry);
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void accumulate_xyz_simple2(PointXYZ<Field1>& acc, const PointXYZ<Field1>& pt) {
|
||||
modmul2(acc.x, acc.x, pt.x); // t00 ← X1 · X2 < 2
|
||||
modmul2(acc.y, acc.y, pt.y); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(acc.z, acc.z, pt.z); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(acc.x, acc.x, pt.x); // t00 ← X1 · X2 < 2
|
||||
modmul2(acc.y, acc.y, pt.y); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(acc.z, acc.z, pt.z); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(acc.x, acc.x, pt.x); // t00 ← X1 · X2 < 2
|
||||
modmul2(acc.y, acc.y, pt.y); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(acc.z, acc.z, pt.z); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(acc.x, acc.x, pt.x); // t00 ← X1 · X2 < 2
|
||||
modmul2(acc.y, acc.y, pt.y); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(acc.z, acc.z, pt.z); // t02 ← Z1 · Z2 < 2
|
||||
}
|
||||
|
||||
__device__ __forceinline__ void accumulate_xyz_simple(Field1 &x1,Field1 &y1,Field1 &z1, const Field1& x2, const Field1& y2, const Field1& z2)
|
||||
{
|
||||
modmul2(x1,x1,x2); // t00 ← X1 · X2 < 2
|
||||
modmul2(y1,y1,y2); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(z1,z1,z2); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(x1,x1,x2); // t00 ← X1 · X2 < 2
|
||||
modmul2(y1,y1,y2); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(z1,z1,z2); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(x1,x1,x2); // t00 ← X1 · X2 < 2
|
||||
modmul2(y1,y1,y2); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(z1,z1,z2); // t02 ← Z1 · Z2 < 2
|
||||
modmul2(x1,x1,x2); // t00 ← X1 · X2 < 2
|
||||
modmul2(y1,y1,y2); // t01 ← Y1 · Y2 < 2
|
||||
modmul2(z1,z1,z2); // t02 ← Z1 · Z2 < 2
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
namespace msm {
|
||||
|
||||
namespace {
|
||||
@@ -141,6 +267,228 @@ namespace msm {
|
||||
if (threadIdx.x == 0) { v_r[blockIdx.x] = v[tid]; }
|
||||
}
|
||||
|
||||
|
||||
template <class T1, class T2, int N_REP>
|
||||
__global__ void add_elements_kernel(const T1* x, const T2* y, T1* result, const unsigned count)
|
||||
{
|
||||
const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (gid >= count) return;
|
||||
T1 res = x[gid];
|
||||
T2 y_gid = y[gid];
|
||||
for (int i = 0; i < N_REP; i++)
|
||||
{
|
||||
T1::accumulate_proj(res,y_gid);
|
||||
// res = ec_add(res,y_gid);
|
||||
// res = res + y_gid;
|
||||
}
|
||||
result[gid] = res;
|
||||
}
|
||||
|
||||
template <class F, int N_REP>
|
||||
__global__ void mul_elements_kernel(const F* x, const F* y, F* result, const unsigned count)
|
||||
{
|
||||
const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (gid >= count) return;
|
||||
F x_gid = x[gid];
|
||||
F res = y[gid];
|
||||
for (int i = 0; i < N_REP; i++)
|
||||
res = res * x_gid;
|
||||
result[gid] = res;
|
||||
}
|
||||
|
||||
template<typename P>
|
||||
__launch_bounds__(256,1)
|
||||
__global__ void segments_reduce_kernel(P* results, P* buckets, uint32_t nof_bms, uint32_t bucketBits) {
|
||||
const uint32_t BUCKET_COUNT=1<<bucketBits;
|
||||
|
||||
// uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x;
|
||||
uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x, warpThread=threadIdx.x & 0x1F;
|
||||
uint32_t threads_per_bm=(gridDim.x*blockDim.x)/nof_bms;
|
||||
uint32_t bucketsPerThread=(BUCKET_COUNT+threads_per_bm-1)/threads_per_bm;
|
||||
uint32_t bm_index = globalTID/threads_per_bm;
|
||||
uint32_t bm_tid = globalTID%threads_per_bm;
|
||||
uint32_t loadIndex, storeIndex, scale;
|
||||
bool infinity, bit;
|
||||
|
||||
if (globalTID >= threads_per_bm*nof_bms) return;
|
||||
|
||||
P sum = P::zero();
|
||||
P sumOfSums = P::zero();
|
||||
P pt = P::zero();
|
||||
|
||||
if (globalTID == 0) printf("bucket count %d\n", BUCKET_COUNT);
|
||||
|
||||
// loadIndex=bm_tid*bucketsPerThread;
|
||||
// pt = buckets[loadIndex];
|
||||
// for(int32_t i=194;i>=0;i--) {
|
||||
|
||||
for(int32_t i=bucketsPerThread-1;i>=0;i--) {
|
||||
loadIndex=bm_tid*bucketsPerThread + i;
|
||||
if(loadIndex<BUCKET_COUNT) {
|
||||
pt = buckets[bm_index*BUCKET_COUNT + loadIndex];
|
||||
// sum = ec_add(sum,pt);
|
||||
accumulate_xyz_simple(sum.x.limbs_storage.limbs, sum.y.limbs_storage.limbs, sum.z.limbs_storage.limbs,pt.x.limbs_storage.limbs, pt.y.limbs_storage.limbs, pt.z.limbs_storage.limbs);
|
||||
// P::accumulate_simple(sum,pt);
|
||||
// P::accumulate_simple_zpmul(sum,pt);
|
||||
// P::accumulate_proj(sum,pt);
|
||||
// sum = sum+pt;
|
||||
// sum = sum + sum;
|
||||
}
|
||||
// accumulate_xyz_simple(sumOfSums.x.limbs_storage.limbs, sumOfSums.y.limbs_storage.limbs, sumOfSums.z.limbs_storage.limbs,sum.x.limbs_storage.limbs, sum.y.limbs_storage.limbs, sum.z.limbs_storage.limbs);
|
||||
// accumulate_xyz(sumOfSums.x, sumOfSums.y, sumOfSums.z,sum.x, sum.y, sum.z);
|
||||
// sumOfSums = ec_add(sumOfSums,sum);
|
||||
// P::accumulate_proj(sumOfSums,sum);
|
||||
// P::accumulate_simple(sumOfSums,sum);
|
||||
// sumOfSums = sumOfSums + sum;
|
||||
}
|
||||
|
||||
// scale=bm_tid * bucketsPerThread;
|
||||
|
||||
// // compute sum=scale*sum
|
||||
// pt = sum;
|
||||
// infinity=true;
|
||||
// for(int32_t i=31;i>=0;i--) {
|
||||
// if(!infinity)
|
||||
// sum = sum + sum;
|
||||
// bit=(scale & (1<<i))!=0;
|
||||
// if(infinity && bit)
|
||||
// infinity=false;
|
||||
// else if(!infinity && bit)
|
||||
// sum = sum + pt;
|
||||
|
||||
// __syncwarp(0xFFFFFFFF);
|
||||
// }
|
||||
|
||||
// if(infinity)
|
||||
// sum = sumOfSums;
|
||||
// else
|
||||
// sum = sum + sumOfSums;
|
||||
|
||||
// sum across the warp
|
||||
// #pragma unroll 1
|
||||
// for(int32_t i=1;i<32;i=i+i) {
|
||||
// P::warpShuffleProjective(pt, sum, threadIdx.x ^ i); //TODO: implement
|
||||
// sum = sum + pt;
|
||||
// }
|
||||
|
||||
// since all threads in the warp have the same sum, we can use a trick!
|
||||
// buckets[globalTID] = sum;
|
||||
storeIndex=globalTID>>5;
|
||||
if(warpThread==0) {
|
||||
results[storeIndex] = sum;
|
||||
// results[storeIndex] = sumOfSums;
|
||||
}
|
||||
}
|
||||
|
||||
template<typename P>
|
||||
__global__ void point_format_switch(PointXYZ<Field1>* zip_buckets, P* ice_buckets, uint32_t bucketBits){
|
||||
const uint32_t BUCKET_COUNT=1<<bucketBits;
|
||||
|
||||
// uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x;
|
||||
uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x, warpThread=threadIdx.x & 0x1F;
|
||||
|
||||
if (globalTID >= BUCKET_COUNT) return;
|
||||
|
||||
for (int i = 0; i < 12; i++)
|
||||
{
|
||||
zip_buckets[globalTID].x[i] = ice_buckets[globalTID].x.limbs_storage.limbs[i];
|
||||
zip_buckets[globalTID].y[i] = ice_buckets[globalTID].y.limbs_storage.limbs[i];
|
||||
zip_buckets[globalTID].z[i] = ice_buckets[globalTID].z.limbs_storage.limbs[i];
|
||||
}
|
||||
}
|
||||
|
||||
// template<typename P>
|
||||
__launch_bounds__(256,1)
|
||||
__global__ void segments_reduce_kernel2(PointXYZ<Field1>* results, PointXYZ<Field1>* buckets, uint32_t nof_bms, uint32_t bucketBits) {
|
||||
const uint32_t BUCKET_COUNT=1<<bucketBits;
|
||||
|
||||
uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x, warpThread=threadIdx.x & 0x1F;
|
||||
uint32_t bucketsPerThread=(BUCKET_COUNT+gridDim.x*blockDim.x-1)/(gridDim.x*blockDim.x);
|
||||
uint32_t loadIndex, storeIndex, scale;
|
||||
bool infinity, bit;
|
||||
// uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x;
|
||||
// uint32_t globalTID=blockIdx.x*blockDim.x+threadIdx.x, warpThread=threadIdx.x & 0x1F;
|
||||
// uint32_t threads_per_bm=(gridDim.x*blockDim.x)/nof_bms;
|
||||
// uint32_t bucketsPerThread=(BUCKET_COUNT+threads_per_bm-1)/threads_per_bm;
|
||||
// uint32_t bm_index = globalTID/threads_per_bm;
|
||||
// uint32_t bm_tid = globalTID%threads_per_bm;
|
||||
// uint32_t loadIndex, storeIndex, scale;
|
||||
// bool infinity, bit;
|
||||
|
||||
// if (globalTID >= threads_per_bm*nof_bms) return;
|
||||
|
||||
PointXYZ<Field1> sum, sumOfSums, pt;
|
||||
initializeXYZ(sum);
|
||||
initializeXYZ(sumOfSums);
|
||||
initializeXYZ(pt);
|
||||
|
||||
// if (globalTID == 0) printf("bucket count %d\n", BUCKET_COUNT);
|
||||
|
||||
// loadIndex=bm_tid*bucketsPerThread;
|
||||
// pt = buckets[loadIndex];
|
||||
// for(int32_t i=194;i>=0;i--) {
|
||||
|
||||
for(int32_t i=bucketsPerThread-1;i>=0;i--) {
|
||||
loadIndex=globalTID*bucketsPerThread + i;
|
||||
if(loadIndex<BUCKET_COUNT) {
|
||||
pt = buckets[loadIndex];
|
||||
// sum = ec_add(sum,pt);
|
||||
// accumulate_xyz_simple2(sum, pt);
|
||||
// P::accumulate_simple(sum,pt);
|
||||
// P::accumulate_simple_zpmul(sum,pt);
|
||||
// P::accumulate_proj(sum,pt);
|
||||
// sum = sum+pt;
|
||||
// sum = sum + sum;
|
||||
}
|
||||
// accumulate_xyz_simple2(sumOfSums, sum);
|
||||
// accumulate_xyz(sumOfSums.x, sumOfSums.y, sumOfSums.z,sum.x, sum.y, sum.z);
|
||||
// sumOfSums = ec_add(sumOfSums,sum);
|
||||
// P::accumulate_proj(sumOfSums,sum);
|
||||
// P::accumulate_simple(sumOfSums,sum);
|
||||
// sumOfSums = sumOfSums + sum;
|
||||
}
|
||||
|
||||
// scale=bm_tid * bucketsPerThread;
|
||||
|
||||
// // compute sum=scale*sum
|
||||
// pt = sum;
|
||||
// infinity=true;
|
||||
// for(int32_t i=31;i>=0;i--) {
|
||||
// if(!infinity)
|
||||
// sum = sum + sum;
|
||||
// bit=(scale & (1<<i))!=0;
|
||||
// if(infinity && bit)
|
||||
// infinity=false;
|
||||
// else if(!infinity && bit)
|
||||
// sum = sum + pt;
|
||||
|
||||
// __syncwarp(0xFFFFFFFF);
|
||||
// }
|
||||
|
||||
// if(infinity)
|
||||
// sum = sumOfSums;
|
||||
// else
|
||||
// sum = sum + sumOfSums;
|
||||
|
||||
// sum across the warp
|
||||
// #pragma unroll 1
|
||||
// for(int32_t i=1;i<32;i=i+i) {
|
||||
// P::warpShuffleProjective(pt, sum, threadIdx.x ^ i); //TODO: implement
|
||||
// sum = sum + pt;
|
||||
// }
|
||||
|
||||
// since all threads in the warp have the same sum, we can use a trick!
|
||||
// buckets[globalTID] = sum;
|
||||
storeIndex=globalTID>>5;
|
||||
if(warpThread==0) {
|
||||
// results[storeIndex] = sum;
|
||||
results[storeIndex] = sumOfSums;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
// this kernel initializes the buckets with zero points
|
||||
// each thread initializes a different bucket
|
||||
template <typename P>
|
||||
@@ -371,6 +719,7 @@ namespace msm {
|
||||
bool are_points_montgomery_form,
|
||||
bool are_results_on_device,
|
||||
bool is_big_triangle,
|
||||
bool segments_reduction,
|
||||
int large_bucket_factor,
|
||||
int precompute_factor,
|
||||
bool is_async,
|
||||
@@ -378,17 +727,20 @@ namespace msm {
|
||||
{
|
||||
CHK_INIT_IF_RETURN();
|
||||
|
||||
printf("c value = %d\n", c);
|
||||
printf("precompute_factor = %d\n", precompute_factor);
|
||||
|
||||
const unsigned nof_scalars = batch_size * single_msm_size; // assuming scalars not shared between batch elements
|
||||
const bool is_nof_points_valid = ((single_msm_size * batch_size) % nof_points == 0);
|
||||
if (!is_nof_points_valid) {
|
||||
THROW_ICICLE_ERR(
|
||||
IcicleError_t::InvalidArgument, "bucket_method_msm: #points must be divisible by single_msm_size*batch_size");
|
||||
}
|
||||
if ((precompute_factor & (precompute_factor - 1)) != 0) {
|
||||
THROW_ICICLE_ERR(
|
||||
IcicleError_t::InvalidArgument,
|
||||
"bucket_method_msm: precompute factors that are not powers of 2 currently unsupported");
|
||||
}
|
||||
// if ((precompute_factor & (precompute_factor - 1)) != 0) {
|
||||
// THROW_ICICLE_ERR(
|
||||
// IcicleError_t::InvalidArgument,
|
||||
// "bucket_method_msm: precompute factors that are not powers of 2 currently unsupported");
|
||||
// }
|
||||
|
||||
const S* d_scalars;
|
||||
S* d_allocated_scalars = nullptr;
|
||||
@@ -617,6 +969,7 @@ namespace msm {
|
||||
cudaEvent_t event_large_buckets_accumulated;
|
||||
// ---------------- This is where handling of large buckets happens (if there are any) -------------
|
||||
if (h_nof_large_buckets > 0 && bucket_th > 0) {
|
||||
printf("large buckets found\n");
|
||||
CHK_IF_RETURN(cudaStreamCreate(&stream_large_buckets));
|
||||
CHK_IF_RETURN(cudaEventCreateWithFlags(&event_large_buckets_accumulated, cudaEventDisableTiming));
|
||||
|
||||
@@ -721,9 +1074,41 @@ namespace msm {
|
||||
if (!are_results_on_device)
|
||||
CHK_IF_RETURN(cudaMallocAsync(&d_allocated_final_result, sizeof(P) * batch_size, stream));
|
||||
|
||||
return CHK_LAST();
|
||||
printf("starting reduction phase...\n");
|
||||
|
||||
// --- Reduction of buckets happens here, after this we'll get a single sum for each bucket module/window ---
|
||||
unsigned nof_empty_bms_per_batch = 0; // for non-triangle accumluation this may be >0
|
||||
P* final_results;
|
||||
PointXYZ<Field1> *final_results_temp, *buckets_temp;
|
||||
|
||||
if (segments_reduction){
|
||||
int32_t ec, smCount;
|
||||
ec=cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, 0);
|
||||
if(ec!=0) return CHK_LAST();
|
||||
// CHK_IF_RETURN(cudaMallocAsync(&final_results, sizeof(P) * nof_bms_in_batch * 8 * smCount, stream));
|
||||
CHK_IF_RETURN(cudaMallocAsync(&final_results_temp, sizeof(PointXYZ<Field1>) * nof_bms_in_batch * 8 * smCount, stream));
|
||||
CHK_IF_RETURN(cudaMallocAsync(&buckets_temp, sizeof(PointXYZ<Field1>) * (total_nof_buckets + nof_bms_in_batch), stream));
|
||||
printf("nof_total buckets %d\n", total_nof_buckets + nof_bms_in_batch);
|
||||
// segments_reduce_kernel<<<smCount, 256, 0, stream>>>(final_results, buckets, nof_bms_in_batch, c);
|
||||
|
||||
|
||||
point_format_switch<<<(1<<14), 256, 0, stream>>>(buckets_temp, buckets, c);
|
||||
cudaDeviceSynchronize();
|
||||
printf("cuda err f %d\n", cudaGetLastError());
|
||||
segments_reduce_kernel2<<<smCount, 256, 0, stream>>>(final_results_temp, buckets_temp, nof_bms_in_batch, c);
|
||||
cudaDeviceSynchronize();
|
||||
printf("cuda err s %d\n", cudaGetLastError());
|
||||
return CHK_LAST();
|
||||
// add_elements_kernel<P,P,256><<<4096, 256, 0, stream>>>(buckets, buckets, buckets, 4096*256);
|
||||
// S* scalars_temp;
|
||||
// CHK_IF_RETURN(cudaMallocAsync(&scalars_temp, sizeof(S) * nof_scalars, stream));
|
||||
// printf("nof scalars %d\n", nof_scalars);
|
||||
// mul_elements_kernel<S,256><<<4096, 256, 0, stream>>>(scalars, scalars, scalars_temp, 4096*256);
|
||||
}
|
||||
|
||||
|
||||
else{
|
||||
if (is_big_triangle || c == 1) {
|
||||
CHK_IF_RETURN(cudaMallocAsync(&final_results, sizeof(P) * nof_bms_in_batch, stream));
|
||||
// launch the bucket module sum kernel - a thread for each bucket module
|
||||
@@ -801,6 +1186,7 @@ namespace msm {
|
||||
source_buckets_count = target_buckets_count;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// ------- This is the final stage where bucket modules/window sums get added up with appropriate weights
|
||||
// -------
|
||||
@@ -835,18 +1221,20 @@ namespace msm {
|
||||
cudaStream_t& stream = config.ctx.stream;
|
||||
|
||||
unsigned c = config.batch_size > 1 ? ((config.c == 0) ? get_optimal_c(msm_size) : config.c) : 16;
|
||||
// unsigned c = config.c;
|
||||
// reduce c to closest power of two (from below) if not using big_triangle reduction logic
|
||||
// TODO: support arbitrary values of c
|
||||
if (!config.is_big_triangle) {
|
||||
while ((c & (c - 1)) != 0)
|
||||
c &= (c - 1);
|
||||
}
|
||||
c = config.c;
|
||||
|
||||
return CHK_STICKY(bucket_method_msm(
|
||||
bitsize, c, scalars, points, config.batch_size, msm_size,
|
||||
(config.points_size == 0) ? msm_size : config.points_size, results, config.are_scalars_on_device,
|
||||
config.are_scalars_montgomery_form, config.are_points_on_device, config.are_points_montgomery_form,
|
||||
config.are_results_on_device, config.is_big_triangle, config.large_bucket_factor, config.precompute_factor,
|
||||
config.are_results_on_device, config.is_big_triangle, config.segments_reduction, config.large_bucket_factor, config.precompute_factor,
|
||||
config.is_async, stream));
|
||||
}
|
||||
|
||||
@@ -868,7 +1256,9 @@ namespace msm {
|
||||
output_bases, bases, sizeof(A) * bases_size,
|
||||
are_bases_on_device ? cudaMemcpyDeviceToDevice : cudaMemcpyHostToDevice, stream));
|
||||
|
||||
unsigned c = 16;
|
||||
// unsigned c = 16;
|
||||
unsigned c = _c;
|
||||
printf("precomp c value = %d\n", c);
|
||||
unsigned total_nof_bms = (P::SCALAR_FF_NBITS - 1) / c + 1;
|
||||
unsigned shift = c * ((total_nof_bms - 1) / precompute_factor + 1);
|
||||
|
||||
|
||||
@@ -1,3 +1,10 @@
|
||||
#include "fields/id.h"
|
||||
// #define FIELD_ID 2
|
||||
#define CURVE_ID 3
|
||||
#include "curves/curve_config.cuh"
|
||||
// #include "fields/field_config.cuh"
|
||||
|
||||
|
||||
#include "msm.cu"
|
||||
|
||||
#include <chrono>
|
||||
@@ -6,10 +13,13 @@
|
||||
|
||||
#include "curves/params/bn254.cuh"
|
||||
#include "fields/field.cuh"
|
||||
// #include "fields/asm.cu"
|
||||
// #include "fields/Chain.cu"
|
||||
// #include "fields/MP.cu"
|
||||
#include "curves/projective.cuh"
|
||||
#include "gpu-utils/device_context.cuh"
|
||||
|
||||
using namespace bn254;
|
||||
// using namespace bn254;
|
||||
|
||||
class Dummy_Scalar
|
||||
{
|
||||
@@ -111,20 +121,33 @@ public:
|
||||
|
||||
// switch between dummy and real:
|
||||
|
||||
typedef scalar_t test_scalar;
|
||||
typedef projective_t test_projective;
|
||||
typedef affine_t test_affine;
|
||||
// typedef scalar_t test_scalar;
|
||||
// typedef projective_t test_projective;
|
||||
// typedef affine_t test_affine;
|
||||
|
||||
typedef curve_config::scalar_t test_scalar;
|
||||
typedef curve_config::projective_t test_projective;
|
||||
typedef curve_config::affine_t test_affine;
|
||||
|
||||
// typedef Dummy_Scalar test_scalar;
|
||||
// typedef Dummy_Projective test_projective;
|
||||
// typedef Dummy_Projective test_affine;
|
||||
|
||||
int main()
|
||||
int main(int argc, char** argv)
|
||||
{
|
||||
int batch_size = 1;
|
||||
|
||||
cudaEvent_t start, stop;
|
||||
float msm_time;
|
||||
|
||||
int msm_log_size = (argc > 1) ? atoi(argv[1]) : 18;
|
||||
int msm_size = 1<<msm_log_size;
|
||||
int batch_size = (argc > 2) ? atoi(argv[2]) : 1;
|
||||
// unsigned msm_size = 1<<21;
|
||||
int msm_size = 12180757;
|
||||
int N = batch_size * msm_size;
|
||||
int precomp_factor = (argc > 3) ? atoi(argv[3]) : 1;
|
||||
int user_c = (argc > 4) ? atoi(argv[4]) : 16;
|
||||
|
||||
printf("running msm curve=%d, 2^%d, batch_size=%d, precomp_factor=%d, c=%d\n",CURVE_ID,msm_log_size, batch_size, precomp_factor, user_c);
|
||||
|
||||
test_scalar* scalars = new test_scalar[N];
|
||||
test_affine* points = new test_affine[N];
|
||||
@@ -136,7 +159,8 @@ int main()
|
||||
|
||||
// projective_t *short_res = (projective_t*)malloc(sizeof(projective_t));
|
||||
// test_projective *large_res = (test_projective*)malloc(sizeof(test_projective));
|
||||
test_projective large_res[2];
|
||||
test_projective res[batch_size];
|
||||
test_projective ref[batch_size];
|
||||
// test_projective batched_large_res[batch_size];
|
||||
// fake_point *large_res = (fake_point*)malloc(sizeof(fake_point));
|
||||
// fake_point batched_large_res[256];
|
||||
@@ -149,11 +173,15 @@ int main()
|
||||
|
||||
test_scalar* scalars_d;
|
||||
test_affine* points_d;
|
||||
test_projective* large_res_d;
|
||||
test_affine* precomp_points_d;
|
||||
test_projective* res_d;
|
||||
test_projective* ref_d;
|
||||
|
||||
cudaMalloc(&scalars_d, sizeof(test_scalar) * msm_size);
|
||||
cudaMalloc(&points_d, sizeof(test_affine) * msm_size);
|
||||
cudaMalloc(&large_res_d, sizeof(test_projective));
|
||||
cudaMalloc(&precomp_points_d, sizeof(test_affine) * msm_size * precomp_factor);
|
||||
cudaMalloc(&res_d, sizeof(test_projective));
|
||||
cudaMalloc(&ref_d, sizeof(test_projective));
|
||||
cudaMemcpy(scalars_d, scalars, sizeof(test_scalar) * msm_size, cudaMemcpyHostToDevice);
|
||||
cudaMemcpy(points_d, points, sizeof(test_affine) * msm_size, cudaMemcpyHostToDevice);
|
||||
|
||||
@@ -172,63 +200,88 @@ int main()
|
||||
msm::MSMConfig config = {
|
||||
ctx, // DeviceContext
|
||||
0, // points_size
|
||||
1, // precompute_factor
|
||||
0, // c
|
||||
precomp_factor, // precompute_factor
|
||||
user_c, // c
|
||||
0, // bitsize
|
||||
10, // large_bucket_factor
|
||||
1, // batch_size
|
||||
100, // large_bucket_factor
|
||||
batch_size, // batch_size
|
||||
false, // are_scalars_on_device
|
||||
false, // are_scalars_montgomery_form
|
||||
false, // are_points_on_device
|
||||
true, // are_points_on_device
|
||||
false, // are_points_montgomery_form
|
||||
true, // are_results_on_device
|
||||
false, // is_big_triangle
|
||||
true, // is_async
|
||||
false, // segments_reduction
|
||||
};
|
||||
|
||||
auto begin1 = std::chrono::high_resolution_clock::now();
|
||||
msm::msm<test_scalar, test_affine, test_projective>(scalars, points, msm_size, config, large_res_d);
|
||||
cudaEvent_t msm_end_event;
|
||||
cudaEventCreate(&msm_end_event);
|
||||
auto end1 = std::chrono::high_resolution_clock::now();
|
||||
auto elapsed1 = std::chrono::duration_cast<std::chrono::nanoseconds>(end1 - begin1);
|
||||
printf("No Big Triangle : %.3f seconds.\n", elapsed1.count() * 1e-9);
|
||||
config.is_big_triangle = true;
|
||||
config.are_results_on_device = false;
|
||||
cudaMemcpy(&large_res[1], large_res_d, sizeof(test_projective), cudaMemcpyDeviceToHost);
|
||||
std::cout << test_projective::to_affine(large_res[1]) << " " << test_projective::is_on_curve(large_res[1])
|
||||
<< std::endl;
|
||||
auto begin = std::chrono::high_resolution_clock::now();
|
||||
msm::msm<test_scalar, test_affine, test_projective>(scalars_d, points_d, msm_size, config, large_res);
|
||||
cudaEventCreate(&start);
|
||||
cudaEventCreate(&stop);
|
||||
|
||||
if (precomp_factor > 1) msm::precompute_msm_bases<test_affine, test_projective>(points_d, msm_size, precomp_factor, user_c, false, ctx, precomp_points_d);
|
||||
|
||||
|
||||
// warm up
|
||||
msm::msm<test_scalar, test_affine, test_projective>(scalars, precomp_factor > 1? precomp_points_d : points_d, msm_size, config, res_d);
|
||||
cudaDeviceSynchronize();
|
||||
|
||||
// auto begin1 = std::chrono::high_resolution_clock::now();
|
||||
cudaEventRecord(start, stream);
|
||||
msm::msm<test_scalar, test_affine, test_projective>(scalars, precomp_factor > 1? precomp_points_d : points_d, msm_size, config, res_d);
|
||||
cudaEventRecord(stop, stream);
|
||||
cudaStreamSynchronize(stream);
|
||||
cudaEventElapsedTime(&msm_time, start, stop);
|
||||
// cudaEvent_t msm_end_event;
|
||||
// cudaEventCreate(&msm_end_event);
|
||||
// auto end1 = std::chrono::high_resolution_clock::now();
|
||||
// auto elapsed1 = std::chrono::duration_cast<std::chrono::nanoseconds>(end1 - begin1);
|
||||
printf("msm time : %.3f ms.\n", msm_time);
|
||||
|
||||
//reference
|
||||
config.c = 16;
|
||||
config.precompute_factor = 1;
|
||||
config.segments_reduction = false;
|
||||
msm::msm<test_scalar, test_affine, test_projective>(scalars, points_d, msm_size, config, ref_d);
|
||||
|
||||
// config.is_big_triangle = true;
|
||||
// config.are_results_on_device = false;
|
||||
// std::cout << test_projective::to_affine(large_res[0]) << std::endl;
|
||||
// auto begin = std::chrono::high_resolution_clock::now();
|
||||
// msm::MSM<test_scalar, test_affine, test_projective>(scalars_d, points_d, msm_size, config, large_res);
|
||||
// test_reduce_triangle(scalars);
|
||||
// test_reduce_rectangle(scalars);
|
||||
// test_reduce_single(scalars);
|
||||
// test_reduce_var(scalars);
|
||||
auto end = std::chrono::high_resolution_clock::now();
|
||||
auto elapsed = std::chrono::duration_cast<std::chrono::nanoseconds>(end - begin);
|
||||
printf("Big Triangle: %.3f seconds.\n", elapsed.count() * 1e-9);
|
||||
// auto end = std::chrono::high_resolution_clock::now();
|
||||
// auto elapsed = std::chrono::duration_cast<std::chrono::nanoseconds>(end - begin);
|
||||
// printf("Big Triangle: %.3f seconds.\n", elapsed.count() * 1e-9);
|
||||
cudaStreamSynchronize(stream);
|
||||
cudaStreamDestroy(stream);
|
||||
|
||||
// std::cout << test_projective::to_affine(large_res[0]) << std::endl;
|
||||
|
||||
cudaMemcpy(res, res_d, sizeof(test_projective) * batch_size, cudaMemcpyDeviceToHost);
|
||||
cudaMemcpy(ref, ref_d, sizeof(test_projective) * batch_size, cudaMemcpyDeviceToHost);
|
||||
|
||||
// reference_msm<test_affine, test_scalar, test_projective>(scalars, points, msm_size);
|
||||
|
||||
// std::cout<<"final results batched large"<<std::endl;
|
||||
// bool success = true;
|
||||
// for (unsigned i = 0; i < batch_size; i++)
|
||||
// {
|
||||
// std::cout<<test_projective::to_affine(batched_large_res[i])<<std::endl;
|
||||
// if (test_projective::to_affine(large_res[i])==test_projective::to_affine(batched_large_res[i])){
|
||||
// std::cout<<"good"<<std::endl;
|
||||
// }
|
||||
// else{
|
||||
// std::cout<<"miss"<<std::endl;
|
||||
// std::cout<<test_projective::to_affine(large_res[i])<<std::endl;
|
||||
// success = false;
|
||||
// }
|
||||
// }
|
||||
// if (success){
|
||||
// std::cout<<"success!"<<std::endl;
|
||||
// }
|
||||
bool success = true;
|
||||
for (unsigned i = 0; i < batch_size; i++)
|
||||
{
|
||||
std::cout<<test_projective::to_affine(res[i])<<std::endl;
|
||||
if (test_projective::to_affine(res[i])==test_projective::to_affine(ref[i])){
|
||||
std::cout<<"good"<<std::endl;
|
||||
}
|
||||
else{
|
||||
std::cout<<"miss"<<std::endl;
|
||||
std::cout<<test_projective::to_affine(ref[i])<<std::endl;
|
||||
success = false;
|
||||
}
|
||||
}
|
||||
if (success){
|
||||
std::cout<<"success!"<<std::endl;
|
||||
}
|
||||
|
||||
// std::cout<<batched_large_res[0]<<std::endl;
|
||||
// std::cout<<batched_large_res[1]<<std::endl;
|
||||
@@ -239,4 +292,4 @@ int main()
|
||||
// std::cout<<pr<<std::endl;
|
||||
|
||||
return 0;
|
||||
}
|
||||
}
|
||||
@@ -134,14 +134,23 @@ namespace mxntt {
|
||||
int n_scalars,
|
||||
uint32_t log_size,
|
||||
eRevType rev_type,
|
||||
bool dit,
|
||||
bool fast_tw,
|
||||
E* out_vec)
|
||||
{
|
||||
int tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
if (tid >= size * batch_size) return;
|
||||
int64_t scalar_id = (tid / columns_batch_size) % size;
|
||||
if (rev_type != eRevType::None)
|
||||
scalar_id = generalized_rev((tid / columns_batch_size) & ((1 << log_size) - 1), log_size, dit, false, rev_type);
|
||||
if (rev_type != eRevType::None) {
|
||||
// Note: when we multiply an in_vec that is mixed (by DIF (I)NTT), we want to shuffle the
|
||||
// scalars the same way (then multiply element-wise). This would be a DIT-digit-reverse shuffle. (this is
|
||||
// confusing but) BUT to avoid shuffling the scalars, we instead want to ask which element in the non-shuffled
|
||||
// vec is now placed at index tid, which is the opposite of a DIT-digit-reverse --> this is the DIF-digit-reverse.
|
||||
// Therefore we use the DIF-digit-reverse to know which element moved to index tid and use it to access the
|
||||
// corresponding element in scalars vec.
|
||||
const bool dif = rev_type == eRevType::NaturalToMixedRev;
|
||||
scalar_id =
|
||||
generalized_rev((tid / columns_batch_size) & ((1 << log_size) - 1), log_size, !dif, fast_tw, rev_type);
|
||||
}
|
||||
out_vec[tid] = *(scalar_vec + ((scalar_id * step) % n_scalars)) * in_vec[tid];
|
||||
}
|
||||
|
||||
@@ -903,6 +912,7 @@ namespace mxntt {
|
||||
S* external_twiddles,
|
||||
S* internal_twiddles,
|
||||
S* basic_twiddles,
|
||||
S* linear_twiddle, // twiddles organized as [1,w,w^2,...] for coset-eval in fast-tw mode
|
||||
int ntt_size,
|
||||
int max_logn,
|
||||
int batch_size,
|
||||
@@ -958,8 +968,8 @@ namespace mxntt {
|
||||
if (is_on_coset && !is_inverse) {
|
||||
batch_elementwise_mul_with_reorder_kernel<<<NOF_BLOCKS, NOF_THREADS, 0, cuda_stream>>>(
|
||||
d_input, ntt_size, columns_batch, batch_size, columns_batch ? batch_size : 1,
|
||||
arbitrary_coset ? arbitrary_coset : external_twiddles, arbitrary_coset ? 1 : coset_gen_index, n_twiddles, logn,
|
||||
reverse_coset, dit, d_output);
|
||||
arbitrary_coset ? arbitrary_coset : linear_twiddle, arbitrary_coset ? 1 : coset_gen_index, n_twiddles, logn,
|
||||
reverse_coset, fast_tw, d_output);
|
||||
|
||||
d_input = d_output;
|
||||
}
|
||||
@@ -991,8 +1001,8 @@ namespace mxntt {
|
||||
if (is_on_coset && is_inverse) {
|
||||
batch_elementwise_mul_with_reorder_kernel<<<NOF_BLOCKS, NOF_THREADS, 0, cuda_stream>>>(
|
||||
d_output, ntt_size, columns_batch, batch_size, columns_batch ? batch_size : 1,
|
||||
arbitrary_coset ? arbitrary_coset : external_twiddles + n_twiddles, arbitrary_coset ? 1 : -coset_gen_index,
|
||||
n_twiddles, logn, reverse_coset, dit, d_output);
|
||||
arbitrary_coset ? arbitrary_coset : linear_twiddle + n_twiddles, arbitrary_coset ? 1 : -coset_gen_index,
|
||||
n_twiddles, logn, reverse_coset, fast_tw, d_output);
|
||||
}
|
||||
|
||||
return CHK_LAST();
|
||||
@@ -1021,6 +1031,8 @@ namespace mxntt {
|
||||
scalar_t* external_twiddles,
|
||||
scalar_t* internal_twiddles,
|
||||
scalar_t* basic_twiddles,
|
||||
scalar_t* linear_twiddles,
|
||||
|
||||
int ntt_size,
|
||||
int max_logn,
|
||||
int batch_size,
|
||||
@@ -1039,6 +1051,8 @@ namespace mxntt {
|
||||
scalar_t* external_twiddles,
|
||||
scalar_t* internal_twiddles,
|
||||
scalar_t* basic_twiddles,
|
||||
scalar_t* linear_twiddles,
|
||||
|
||||
int ntt_size,
|
||||
int max_logn,
|
||||
int batch_size,
|
||||
|
||||
@@ -717,8 +717,7 @@ namespace ntt {
|
||||
d_input, d_output, domain.twiddles, size, domain.max_size, batch_size, is_inverse, config.ordering, coset,
|
||||
coset_index, stream));
|
||||
} else {
|
||||
const bool is_on_coset = (coset_index != 0) || coset;
|
||||
const bool is_fast_twiddles_enabled = (domain.fast_external_twiddles != nullptr) && !is_on_coset;
|
||||
const bool is_fast_twiddles_enabled = (domain.fast_external_twiddles != nullptr);
|
||||
S* twiddles = is_fast_twiddles_enabled
|
||||
? (is_inverse ? domain.fast_external_twiddles_inv : domain.fast_external_twiddles)
|
||||
: domain.twiddles;
|
||||
@@ -728,9 +727,11 @@ namespace ntt {
|
||||
S* basic_twiddles = is_fast_twiddles_enabled
|
||||
? (is_inverse ? domain.fast_basic_twiddles_inv : domain.fast_basic_twiddles)
|
||||
: domain.basic_twiddles;
|
||||
S* linear_twiddles = domain.twiddles; // twiddles organized as [1,w,w^2,...]
|
||||
CHK_IF_RETURN(mxntt::mixed_radix_ntt(
|
||||
d_input, d_output, twiddles, internal_twiddles, basic_twiddles, size, domain.max_log_size, batch_size,
|
||||
config.columns_batch, is_inverse, is_fast_twiddles_enabled, config.ordering, coset, coset_index, stream));
|
||||
d_input, d_output, twiddles, internal_twiddles, basic_twiddles, linear_twiddles, size, domain.max_log_size,
|
||||
batch_size, config.columns_batch, is_inverse, is_fast_twiddles_enabled, config.ordering, coset, coset_index,
|
||||
stream));
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -53,6 +53,12 @@ A set corresponding to each field contains headers that shouldn't be included.
|
||||
FIELDS_CONFIG = {
|
||||
"babybear": {
|
||||
"poseidon.h",
|
||||
},
|
||||
"stark252": {
|
||||
"poseidon.h",
|
||||
"field_ext.h",
|
||||
"vec_ops_ext.h",
|
||||
"ntt_ext.h",
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -91,7 +91,7 @@ This is normally fixed by exporting the path to the shared library location in t
|
||||
### cuda_runtime.h: No such file or directory
|
||||
|
||||
```
|
||||
# github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381
|
||||
# github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381
|
||||
In file included from wrappers/golang/curves/bls12381/curve.go:5:
|
||||
wrappers/golang/curves/bls12381/include/curve.h:1:10: fatal error: cuda_runtime.h: No such file or directory
|
||||
1 | #include <cuda_runtime.h>
|
||||
|
||||
@@ -7,9 +7,11 @@ DEVMODE=OFF
|
||||
EXT_FIELD=OFF
|
||||
BUILD_CURVES=( )
|
||||
BUILD_FIELDS=( )
|
||||
BUILD_HASHES=( )
|
||||
|
||||
SUPPORTED_CURVES=("bn254" "bls12_377" "bls12_381" "bw6_761", "grumpkin")
|
||||
SUPPORTED_FIELDS=("babybear")
|
||||
SUPPORTED_HASHES=("keccak")
|
||||
|
||||
if [[ $1 == "-help" ]]; then
|
||||
echo "Build script for building ICICLE cpp libraries"
|
||||
@@ -67,6 +69,15 @@ do
|
||||
-field-ext)
|
||||
EXT_FIELD=ON
|
||||
;;
|
||||
-hash*)
|
||||
hash=$(echo "$arg_lower" | cut -d'=' -f2)
|
||||
if [[ $hash == "all" ]]
|
||||
then
|
||||
BUILD_HASHES=("${SUPPORTED_HASHES[@]}")
|
||||
else
|
||||
BUILD_HASHES=( $hash )
|
||||
fi
|
||||
;;
|
||||
-devmode)
|
||||
DEVMODE=ON
|
||||
;;
|
||||
@@ -105,3 +116,10 @@ do
|
||||
cmake --build build -j8 && rm build_config.txt
|
||||
done
|
||||
|
||||
for HASH in "${BUILD_HASHES[@]}"
|
||||
do
|
||||
echo "HASH=${HASH_DEFINED}" > build_config.txt
|
||||
echo "DEVMODE=${DEVMODE}" >> build_config.txt
|
||||
cmake -DCMAKE_CUDA_COMPILER=$CUDA_COMPILER_PATH -DBUILD_HASH=$HASH -DDEVMODE=$DEVMODE -DCMAKE_BUILD_TYPE=Release -S . -B build
|
||||
cmake --build build -j8 && rm build_config.txt
|
||||
done
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
package core
|
||||
|
||||
import (
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type IcicleErrorCode int
|
||||
|
||||
@@ -4,7 +4,7 @@ import (
|
||||
"fmt"
|
||||
"unsafe"
|
||||
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type MSMConfig struct {
|
||||
|
||||
@@ -3,8 +3,8 @@ package core
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core/internal"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core/internal"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
@@ -4,7 +4,7 @@ import (
|
||||
"fmt"
|
||||
"unsafe"
|
||||
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type NTTDir int8
|
||||
|
||||
@@ -3,8 +3,8 @@ package core
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core/internal"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core/internal"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -3,7 +3,7 @@ package core
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type HostOrDeviceSlice interface {
|
||||
|
||||
@@ -5,7 +5,7 @@ import (
|
||||
"testing"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core/internal"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core/internal"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -4,7 +4,7 @@ import (
|
||||
"fmt"
|
||||
"unsafe"
|
||||
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type VecOps int
|
||||
|
||||
@@ -3,7 +3,7 @@ package core
|
||||
import (
|
||||
"testing"
|
||||
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type Projective struct {
|
||||
|
||||
@@ -5,8 +5,8 @@ package ecntt
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func ECNtt[T any](points core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type G2Projective struct {
|
||||
|
||||
@@ -5,9 +5,10 @@ package g2
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func G2GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -5,9 +5,10 @@ package msm
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -4,14 +4,12 @@ package ntt
|
||||
// #include "ntt.h"
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
)
|
||||
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
)
|
||||
|
||||
func Ntt[T any](scalars core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
)
|
||||
|
||||
type PolynomialHandle = C.struct_PolynomialInst
|
||||
|
||||
@@ -6,9 +6,10 @@ import "C"
|
||||
import (
|
||||
"encoding/binary"
|
||||
"fmt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestAffineZero(t *testing.T) {
|
||||
|
||||
@@ -3,10 +3,10 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
ecntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
ecntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestG2AffineZero(t *testing.T) {
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -8,14 +8,14 @@ import (
|
||||
"github.com/stretchr/testify/assert"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377"
|
||||
bls12377 "github.com/consensys/gnark-crypto/ecc/bls12-377"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/g2"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffineG2(p g2.G2Projective) bls12377.G2Affine {
|
||||
|
||||
@@ -4,10 +4,10 @@ import (
|
||||
"os"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/polynomial"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fr/fft"
|
||||
)
|
||||
|
||||
@@ -8,14 +8,14 @@ import (
|
||||
"github.com/stretchr/testify/assert"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377"
|
||||
bls12377 "github.com/consensys/gnark-crypto/ecc/bls12-377"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/msm"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/msm"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffine(p icicleBls12_377.Projective) bls12377.G1Affine {
|
||||
|
||||
@@ -6,11 +6,11 @@ import (
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fr"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-377/fr/fft"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -3,11 +3,12 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
// "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/vecOps"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
|
||||
// "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/vecOps"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,11 +1,12 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -3,10 +3,10 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12377/vecOps"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_377 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12377/vecOps"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -5,8 +5,8 @@ package vecOps
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
|
||||
"unsafe"
|
||||
)
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type Projective struct {
|
||||
|
||||
@@ -5,8 +5,8 @@ package ecntt
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func ECNtt[T any](points core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type G2Projective struct {
|
||||
|
||||
@@ -5,9 +5,10 @@ package g2
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func G2GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -5,9 +5,10 @@ package msm
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -4,14 +4,12 @@ package ntt
|
||||
// #include "ntt.h"
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
)
|
||||
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
)
|
||||
|
||||
func Ntt[T any](scalars core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
)
|
||||
|
||||
type PolynomialHandle = C.struct_PolynomialInst
|
||||
|
||||
@@ -6,9 +6,10 @@ import "C"
|
||||
import (
|
||||
"encoding/binary"
|
||||
"fmt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestAffineZero(t *testing.T) {
|
||||
|
||||
@@ -3,10 +3,10 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
ecntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
ecntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestG2AffineZero(t *testing.T) {
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -8,14 +8,14 @@ import (
|
||||
"github.com/stretchr/testify/assert"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381"
|
||||
bls12381 "github.com/consensys/gnark-crypto/ecc/bls12-381"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/g2"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffineG2(p g2.G2Projective) bls12381.G2Affine {
|
||||
|
||||
@@ -4,10 +4,10 @@ import (
|
||||
"os"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/polynomial"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fr/fft"
|
||||
)
|
||||
|
||||
@@ -8,14 +8,14 @@ import (
|
||||
"github.com/stretchr/testify/assert"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381"
|
||||
bls12381 "github.com/consensys/gnark-crypto/ecc/bls12-381"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/msm"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/msm"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffine(p icicleBls12_381.Projective) bls12381.G1Affine {
|
||||
|
||||
@@ -6,11 +6,11 @@ import (
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fr"
|
||||
"github.com/consensys/gnark-crypto/ecc/bls12-381/fr/fft"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -3,11 +3,12 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
// "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/vecOps"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
|
||||
// "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/vecOps"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,11 +1,12 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -3,10 +3,10 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bls12381/vecOps"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bls12_381 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bls12381/vecOps"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -5,8 +5,8 @@ package vecOps
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
|
||||
"unsafe"
|
||||
)
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type Projective struct {
|
||||
|
||||
@@ -5,8 +5,8 @@ package ecntt
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func ECNtt[T any](points core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
type G2Projective struct {
|
||||
|
||||
@@ -5,9 +5,10 @@ package g2
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func G2GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -5,9 +5,10 @@ package msm
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
func GetDefaultMSMConfig() core.MSMConfig {
|
||||
|
||||
@@ -4,14 +4,12 @@ package ntt
|
||||
// #include "ntt.h"
|
||||
import "C"
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
func Ntt[T any](scalars core.HostOrDeviceSlice, dir core.NTTDir, cfg *core.NTTConfig[T], results core.HostOrDeviceSlice) core.IcicleError {
|
||||
|
||||
@@ -7,8 +7,8 @@ import "C"
|
||||
import (
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
)
|
||||
|
||||
type PolynomialHandle = C.struct_PolynomialInst
|
||||
|
||||
@@ -6,9 +6,10 @@ import "C"
|
||||
import (
|
||||
"encoding/binary"
|
||||
"fmt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
"unsafe"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestAffineZero(t *testing.T) {
|
||||
|
||||
@@ -3,10 +3,10 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
ecntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
ecntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/ecntt"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
func TestG2AffineZero(t *testing.T) {
|
||||
|
||||
@@ -1,10 +1,11 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
@@ -12,10 +12,10 @@ import (
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/g2"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/g2"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffineG2(p g2.G2Projective) bn254.G2Affine {
|
||||
|
||||
@@ -4,10 +4,10 @@ import (
|
||||
"os"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/ntt"
|
||||
poly "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/polynomial"
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr/fft"
|
||||
)
|
||||
|
||||
@@ -12,10 +12,10 @@ import (
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fp"
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
icicleBn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/msm"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
icicleBn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/msm"
|
||||
)
|
||||
|
||||
func projectiveToGnarkAffine(p icicleBn254.Projective) bn254.G1Affine {
|
||||
|
||||
@@ -6,11 +6,11 @@ import (
|
||||
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr"
|
||||
"github.com/consensys/gnark-crypto/ecc/bn254/fr/fft"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
ntt "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
cr "github.com/ingonyama-zk/icicle/v2/wrappers/golang/cuda_runtime"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
ntt "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -3,11 +3,12 @@ package tests
|
||||
import (
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
// "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254/vecOps"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
|
||||
// "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/ntt"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/polynomial"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254/vecOps"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
|
||||
@@ -1,11 +1,12 @@
|
||||
package tests
|
||||
|
||||
import (
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
"testing"
|
||||
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/core"
|
||||
bn254 "github.com/ingonyama-zk/icicle/v2/wrappers/golang/curves/bn254"
|
||||
"github.com/ingonyama-zk/icicle/v2/wrappers/golang/test_helpers"
|
||||
"github.com/stretchr/testify/assert"
|
||||
)
|
||||
|
||||
const (
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Reference in New Issue
Block a user