-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathcommitment.cuh
127 lines (104 loc) · 4.88 KB
/
commitment.cuh
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
#ifndef COMMITMENT_CUH
#define COMMITMENT_CUH
#include "fr-tensor.cuh"
#include "g1-tensor.cuh"
#include "proof.cuh"
class Commitment: public G1TensorJacobian
{
public:
using G1TensorJacobian::G1TensorJacobian;
using G1TensorJacobian::operator+;
using G1TensorJacobian::operator-;
using G1TensorJacobian::operator*;
using G1TensorJacobian::operator*=;
G1TensorJacobian commit(const FrTensor& t) const;
static Commitment random_generators(uint size);
Fr_t open(const FrTensor& t, const G1TensorJacobian& c, const vector<Fr_t>& u, uint& proof_size_count) const;
static Fr_t me_open(const FrTensor& t, const Commitment& generators, vector<Fr_t>::const_iterator begin, vector<Fr_t>::const_iterator end, vector<G1Jacobian_t>& proof);
};
KERNEL void sum_axis_n_optimized(GLOBAL G1Jacobian_t* arr, GLOBAL G1Jacobian_t* arr_out, uint n, uint m) {
const uint gid = GET_GLOBAL_ID();
if (gid >= m) return;
__shared__ G1Jacobian_t shared_data[G1NumThread];
uint local_id = threadIdx.x;
G1Jacobian_t sum = arr[gid * n + local_id];
for (uint i = local_id + G1NumThread; i < n; i += G1NumThread) {
sum = blstrs__g1__G1Affine_add(sum, arr[gid * n + i]);
}
shared_data[local_id] = sum;
__syncthreads();
// Reduce shared_data to a single value
for (uint s = G1NumThread / 2; s > 0; s >>= 1) {
if (local_id < s) {
shared_data[local_id] = blstrs__g1__G1Affine_add(shared_data[local_id], shared_data[local_id + s]);
}
__syncthreads();
}
if (local_id == 0) arr_out[gid] = shared_data[0];
}
G1TensorJacobian Commitment::commit(const FrTensor& t) const
{
if (t.size % size != 0) throw std::runtime_error("Incompatible dimensions");
auto t_unmont = t;
t_unmont.unmont();
uint m = t.size / size;
G1TensorJacobian temp = (*this) * t_unmont;
G1TensorJacobian out(m);
sum_axis_n_optimized<<<(m+G1NumThread-1)/G1NumThread,G1NumThread>>>(temp.gpu_data, out.gpu_data, size, m);
cudaDeviceSynchronize();
return out;
}
KERNEL void me_open_step(GLOBAL Fr_t* scalars, GLOBAL G1Jacobian_t* generators, Fr_t u, // always assume that scalars and u is in mont form
GLOBAL Fr_t* new_scalars, GLOBAL G1Jacobian_t* new_generators,
GLOBAL G1Jacobian_t* temp_out, GLOBAL G1Jacobian_t* temp_out0, GLOBAL G1Jacobian_t* temp_out1,
uint new_size) // always assume that old size is even, as will be implemented
{
const uint gid = GET_GLOBAL_ID();
if (gid >= new_size) return;
uint gid0 = 2 * gid;
uint gid1 = 2 * gid + 1;
Fr_t u_unmont = blstrs__scalar__Scalar_unmont(u);
new_scalars[gid] = blstrs__scalar__Scalar_add(scalars[gid0], blstrs__scalar__Scalar_mul(u, blstrs__scalar__Scalar_sub(scalars[gid1], scalars[gid0])));
new_generators[gid] = blstrs__g1__G1Affine_add(generators[gid1], G1Jacobian_mul(blstrs__g1__G1Affine_add(generators[gid0], G1Jacobian_minus(generators[gid1])), u_unmont));
temp_out[gid] = blstrs__g1__G1Affine_add(G1Jacobian_mul(generators[gid0], scalars[gid0]), G1Jacobian_mul(generators[gid1], scalars[gid1]));
temp_out0[gid] = G1Jacobian_mul(generators[gid1], scalars[gid0]);
temp_out1[gid] = G1Jacobian_mul(generators[gid0], scalars[gid1]);
}
Fr_t Commitment::me_open(const FrTensor& t, const Commitment& generators, vector<Fr_t>::const_iterator begin, vector<Fr_t>::const_iterator end, vector<G1Jacobian_t>& proof)
{
if (t.size != generators.size) throw std::runtime_error("Incompatible dimensions");
if (begin >= end)
{
proof.push_back(generators(0));
return t(0);
}
uint new_size = t.size / 2;
FrTensor new_scalars(new_size);
Commitment new_generators(new_size);
G1TensorJacobian temp(new_size), temp0(new_size), temp1(new_size);
me_open_step<<<(new_size+G1NumThread-1)/G1NumThread,G1NumThread>>>(t.gpu_data, generators.gpu_data, *begin,
new_scalars.gpu_data, new_generators.gpu_data, temp.gpu_data, temp0.gpu_data, temp1.gpu_data, new_size);
cudaDeviceSynchronize();
proof.push_back(temp.sum());
proof.push_back(temp0.sum());
proof.push_back(temp1.sum());
return me_open(new_scalars, new_generators, begin + 1, end, proof);
}
Fr_t Commitment::open(const FrTensor& t, const G1TensorJacobian& com, const vector<Fr_t>& u, uint& proof_size_count) const
{
const vector<Fr_t> u_out(u.end() - ceilLog2(com.size), u.end());
const vector<Fr_t> u_in(u.begin(), u.end() - ceilLog2(com.size));
auto g_temp = com(u_out);
if (size != (1 << u_in.size())) throw std::runtime_error("Incompatible dimensions");
vector<G1Jacobian_t> proof;
auto val = me_open(t.partial_me(u_out, 1 << u_in.size()), *this, u_in.begin(), u_in.end(), proof);
proof_size_count += proof.size() * 36;
return val;
}
Commitment Commitment::random_generators(uint size)
{
Commitment out(size, G1Jacobian_generator);
out *= FrTensor::random(size);
return out;
}
#endif