Skip to content

Commit 04084d5

Browse files
unzvfuunzvfu
authored
Implement independent MODNUM concept (#56)
* Move Monty functions to modnum directory. * Initial support for a modnum concept.
1 parent ab40ae6 commit 04084d5

File tree

9 files changed

+106
-73
lines changed

9 files changed

+106
-73
lines changed

‎bench/bench.cu‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -5,6 +5,7 @@
55
#include "fixnum/warp_fixnum.cu"
66
#include "array/fixnum_array.h"
77
#include "functions/modexp.cu"
8+
#include "modnum/monty_mul.cu"
89

910
using namespace std;
1011
using namespace cuFIXNUM;
@@ -39,8 +40,9 @@ struct sqr_wide {
3940
template< typename fixnum >
4041
struct my_modexp {
4142
__device__ void operator()(fixnum &z, fixnum x) {
43+
typedef modnum_monty_cios<fixnum> modnum;
44+
modexp<modnum> me(x, x);
4245
fixnum zz;
43-
modexp<fixnum> me(x, x);
4446
me(zz, x);
4547
z = zz;
4648
};

‎src/functions/chinese.cu‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22

33
#include "functions/quorem_preinv.cu"
44
#include "functions/multi_modexp.cu"
5+
#include "modnum/monty_mul.cu"
56

67
namespace cuFIXNUM {
78

@@ -25,6 +26,8 @@ __device__
2526
chinese<fixnum>::chinese(fixnum p_, fixnum q_)
2627
: p(p_), q(q_), mod_q(q)
2728
{
29+
typedef modnum_monty_cios<fixnum> modnum;
30+
2831
// TODO: q is now stored here and in mod_q; need to work out how
2932
// to share q between them. Probably best just to provide quorem_preinv
3033
// with an accessor to the divisor.
@@ -33,7 +36,7 @@ chinese<fixnum>::chinese(fixnum p_, fixnum q_)
3336
// Use a^(q-2) = 1 (mod q)
3437
fixnum qm2, two = fixnum::two();
3538
fixnum::sub(qm2, q, two);
36-
multi_modexp<fixnum> minv(q);
39+
multi_modexp<modnum> minv(q);
3740
minv(c, p, qm2);
3841
}
3942

‎src/functions/modexp.cu‎

Lines changed: 26 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -1,12 +1,13 @@
11
#pragma once
22

33
#include "functions/internal/modexp_impl.cu"
4-
#include "functions/monty_mul.cu"
4+
#include "modnum/monty_mul.cu"
55

66
namespace cuFIXNUM {
77

8-
template< typename fixnum >
8+
template< typename modnum_tp >
99
class modexp {
10+
typedef typename modnum_tp::fixnum fixnum;
1011
typedef typename fixnum::digit digit;
1112

1213
// Decomposition of the exponent for use in the constant-width sliding-window
@@ -17,9 +18,7 @@ class modexp {
1718
int exp_wins_len;
1819
int window_size;
1920

20-
// TODO: Generalise modexp so that it can work with any modular
21-
// multiplication algorithm.
22-
const monty_mul<fixnum> monty;
21+
const modnum_tp modnum;
2322

2423
// Helper functions for decomposing the exponent into windows.
2524
__device__ uint32_t
@@ -44,9 +43,9 @@ public:
4443
};
4544

4645

47-
template< typename fixnum >
46+
template< typename modnum_tp >
4847
__device__ uint32_t
49-
modexp<fixnum>::scan_nonzero_window(int &hi_idx, fixnum &n, int max_window_bits) {
48+
modexp<modnum_tp>::scan_nonzero_window(int &hi_idx, fixnum &n, int max_window_bits) {
5049
uint32_t bits_remaining = hi_idx + 1, win_bits;
5150
digit w, lsd = fixnum::bottom_digit(n);
5251

@@ -59,19 +58,19 @@ modexp<fixnum>::scan_nonzero_window(int &hi_idx, fixnum &n, int max_window_bits)
5958
}
6059

6160

62-
template< typename fixnum >
61+
template< typename modnum_tp >
6362
__device__ int
64-
modexp<fixnum>::scan_zero_window(int &hi_idx, fixnum &n) {
63+
modexp<modnum_tp>::scan_zero_window(int &hi_idx, fixnum &n) {
6564
int nzeros = fixnum::two_valuation(n);
6665
fixnum::rshift(n, n, nzeros);
6766
hi_idx -= nzeros;
6867
return nzeros;
6968
}
7069

7170

72-
template< typename fixnum >
71+
template< typename modnum_tp >
7372
__device__ uint32_t
74-
modexp<fixnum>::scan_window(int &hi_idx, fixnum &n, int max_window_bits) {
73+
modexp<modnum_tp>::scan_window(int &hi_idx, fixnum &n, int max_window_bits) {
7574
int nzeros;
7675
uint32_t window;
7776
nzeros = scan_zero_window(hi_idx, n);
@@ -82,10 +81,10 @@ modexp<fixnum>::scan_window(int &hi_idx, fixnum &n, int max_window_bits) {
8281
}
8382

8483

85-
template< typename fixnum >
84+
template< typename modnum_tp >
8685
__device__
87-
modexp<fixnum>::modexp(fixnum mod, fixnum exp)
88-
: monty(mod)
86+
modexp<modnum_tp>::modexp(fixnum mod, fixnum exp)
87+
: modnum(mod)
8988
{
9089
// sliding window decomposition
9190
int hi_idx;
@@ -115,18 +114,18 @@ modexp<fixnum>::modexp(fixnum mod, fixnum exp)
115114
}
116115

117116

118-
template< typename fixnum >
117+
template< typename modnum_tp >
119118
__device__
120-
modexp<fixnum>::~modexp()
119+
modexp<modnum_tp>::~modexp()
121120
{
122121
if (fixnum::layout::laneIdx() == 0)
123122
free(exp_wins);
124123
}
125124

126125

127-
template< typename fixnum >
126+
template< typename modnum_tp >
128127
__device__ void
129-
modexp<fixnum>::operator()(fixnum &z, fixnum x) const
128+
modexp<modnum_tp>::operator()(fixnum &z, fixnum x) const
130129
{
131130
static constexpr int WINDOW_MAX_BITS = 16;
132131
static constexpr int WINDOW_LEN_MASK = (1UL << WINDOW_MAX_BITS) - 1UL;
@@ -143,7 +142,7 @@ modexp<fixnum>::operator()(fixnum &z, fixnum x) const
143142
//z = fixnum::one();
144143
// TODO: This complicated way of producing a 1 is to
145144
// accommodate the possibility that monty.is_valid is false.
146-
monty.from_monty(z, monty.one());
145+
modnum.from_modnum(z, modnum.one());
147146
return;
148147
}
149148

@@ -152,13 +151,13 @@ modexp<fixnum>::operator()(fixnum &z, fixnum x) const
152151
int window_max = 1U << window_size;
153152
/* G[t] = z^(2t + 1) t >= 0 (odd powers of z) */
154153
fixnum G[WINDOW_MAX_VAL_REDUCED / 2];
155-
monty.to_monty(z, x);
154+
modnum.to_modnum(z, x);
156155
G[0] = z;
157156
if (window_size > 1) {
158-
monty(z, z);
157+
modnum.sqr(z, z);
159158
for (int t = 1; t < window_max / 2; ++t) {
160159
G[t] = G[t - 1];
161-
monty(G[t], G[t], z);
160+
modnum.mul(G[t], G[t], z);
162161
}
163162
}
164163

@@ -171,22 +170,22 @@ modexp<fixnum>::operator()(fixnum &z, fixnum x) const
171170

172171
z = G[e / 2];
173172
while (two_val-- > 0)
174-
monty(z, z);
173+
modnum.sqr(z, z);
175174

176175
while (windows >= exp_wins) {
177176
two_val = window_size;
178177
while (two_val-- > 0)
179-
monty(z, z);
178+
modnum.sqr(z, z);
180179

181180
win = *windows--;
182181
two_val = win & WINDOW_LEN_MASK;
183182
e = win >> WINDOW_MAX_BITS;
184183

185-
monty(z, z, G[e / 2]);
184+
modnum.mul(z, z, G[e / 2]);
186185
while (two_val-- > 0)
187-
monty(z, z);
186+
modnum.sqr(z, z);
188187
}
189-
monty.from_monty(z, z);
188+
modnum.from_modnum(z, z);
190189
}
191190

192191
} // End namespace cuFIXNUM

‎src/functions/multi_modexp.cu‎

Lines changed: 19 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,24 +1,26 @@
11
#pragma once
22

33
#include "functions/internal/modexp_impl.cu"
4-
#include "functions/monty_mul.cu"
4+
#include "modnum/monty_mul.cu"
55

66
namespace cuFIXNUM {
77

88
template<
9-
typename fixnum,
10-
int WINDOW_SIZE = internal::bytes_to_k_ary_window_size(fixnum::BYTES) >
9+
typename modnum_tp,
10+
int WINDOW_SIZE = internal::bytes_to_k_ary_window_size(modnum_tp::fixnum::BYTES) >
1111
class multi_modexp {
12-
static_assert(WINDOW_SIZE >= 1 && WINDOW_SIZE < fixnum::digit::BITS,
12+
static_assert(WINDOW_SIZE >= 1 && WINDOW_SIZE < modnum_tp::fixnum::digit::BITS,
1313
"Invalid window size.");
1414

1515
// TODO: Generalise multi_modexp so that it can work with any modular
1616
// multiplication algorithm.
17-
const monty_mul<fixnum> monty;
17+
const modnum_tp modnum;
1818

1919
public:
20+
typedef typename modnum_tp::fixnum fixnum;
21+
2022
__device__ multi_modexp(fixnum mod)
21-
: monty(mod) { }
23+
: modnum(mod) { }
2224

2325
__device__ void operator()(fixnum &z, fixnum x, fixnum e) const;
2426
};
@@ -54,11 +56,11 @@ public:
5456
* 14.83] since there the number of squarings depends on the 2-adic valuation of
5557
* the window value.
5658
*/
57-
template< typename fixnum, int WINDOW_SIZE >
59+
template< typename modnum_tp, int WINDOW_SIZE >
5860
__device__ void
59-
multi_modexp<fixnum, WINDOW_SIZE>::operator()(fixnum &z, fixnum x, fixnum e) const
61+
multi_modexp<modnum_tp, WINDOW_SIZE>::operator()(fixnum &z, fixnum x, fixnum e) const
6062
{
61-
typedef typename fixnum::digit digit;
63+
typedef typename modnum_tp::fixnum::digit digit;
6264
static constexpr int WIDTH = fixnum::SLOT_WIDTH;
6365

6466
// Window decomposition: digit::BITS = q * WINDOW_SIZE + r.
@@ -67,11 +69,11 @@ multi_modexp<fixnum, WINDOW_SIZE>::operator()(fixnum &z, fixnum x, fixnum e) con
6769

6870
/* G[t] = z^t, t >= 0 */
6971
fixnum G[WINDOW_MAX];
70-
monty.to_monty(z, x);
71-
G[0] = monty.one();
72+
modnum.to_modnum(z, x);
73+
G[0] = modnum.one();
7274
for (int t = 1; t < WINDOW_MAX; ++t) {
7375
G[t] = G[t - 1];
74-
monty(G[t], G[t], z);
76+
modnum.mul(G[t], G[t], z);
7577
}
7678

7779
z = G[0];
@@ -87,22 +89,22 @@ multi_modexp<fixnum, WINDOW_SIZE>::operator()(fixnum &z, fixnum x, fixnum e) con
8789
// statements manually. Idem for the remainder below.
8890
// Investigate how this is even possible!
8991
for (int k = 0; k < WINDOW_SIZE; ++k)
90-
monty(z, z);
92+
modnum.sqr(z, z);
9193
digit fj;
9294
// win = (f >> j) & WINDOW_MAIN_MASK;
9395
digit::rshift(fj, f, j);
9496
digit::rem_2exp(win, fj, WINDOW_SIZE);
95-
monty(z, z, G[win]);
97+
modnum.mul(z, z, G[win]);
9698
}
9799

98100
// Remainder
99101
for (int k = 0; k < WINDOW_REM_BITS; ++k)
100-
monty(z, z);
102+
modnum.sqr(z, z);
101103
//win = f & WINDOW_REM_MASK;
102104
digit::rem_2exp(win, f, WINDOW_REM_BITS);
103-
monty(z, z, G[win]);
105+
modnum.mul(z, z, G[win]);
104106
}
105-
monty.from_monty(z, z);
107+
modnum.from_modnum(z, z);
106108
}
107109

108110
} // End namespace cuFIXNUM

‎src/functions/paillier_decrypt.cu‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@
44
#include "functions/divexact.cu"
55
#include "functions/chinese.cu"
66
#include "functions/multi_modexp.cu"
7+
#include "modnum/monty_mul.cu"
78

89
namespace cuFIXNUM {
910

@@ -90,7 +91,8 @@ private:
9091
quorem_preinv<fixnum> mod_p2;
9192

9293
// Modexp for x |--> x^(p - 1) (mod p^2)
93-
modexp<fixnum> pow;
94+
typedef modnum_monty_cios<fixnum> modnum;
95+
modexp<modnum> pow;
9496

9597
// TODO: It is flipping stupid that these are necessary.
9698
__device__ fixnum square(fixnum p) {
@@ -135,7 +137,7 @@ paillier_decrypt_mod<fixnum>::paillier_decrypt_mod(fixnum p, fixnum n)
135137
// Use a^(p-2) = 1 (mod p)
136138
fixnum pm2;
137139
fixnum::sub(pm2, p, fixnum::two());
138-
multi_modexp<fixnum> minv(p);
140+
multi_modexp<modnum> minv(p);
139141
minv(h, t, pm2);
140142
}
141143

‎src/functions/paillier_encrypt.cu‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22

33
#include "functions/quorem_preinv.cu"
44
#include "functions/multi_modexp.cu"
5+
#include "modnum/monty_mul.cu"
56

67
namespace cuFIXNUM {
78

@@ -27,9 +28,11 @@ public:
2728
}
2829

2930
private:
31+
typedef modnum_monty_cios<fixnum> modnum;
32+
3033
fixnum n;
3134
fixnum n_sqr;
32-
modexp<fixnum> pow;
35+
modexp<modnum> pow;
3336
quorem_preinv<fixnum> mod_n2;
3437

3538
// TODO: It is flipping stupid that this is necessary.

0 commit comments

Comments
 (0)