Skip to content

Commit 0ce719c

Browse files
unzvfuunzvfu
authored
Implement REDC-based modnum (#57)
* Rename files. * Implement REDC mulmod; factor out common Monty operations. * Move normalise to internal monty code. * Add comment; minor code reformatting. * Fix bug with digit->fixnum conversion of cy; disable hypothesis code. * Test REDC and CIOS modexp code. * Bench CIOS and REDC modexp separately. * Move code chunk.
1 parent 04084d5 commit 0ce719c

File tree

12 files changed

+388
-151
lines changed

12 files changed

+388
-151
lines changed

‎bench/bench.cu‎

Lines changed: 39 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,9 @@
55
#include "fixnum/warp_fixnum.cu"
66
#include "array/fixnum_array.h"
77
#include "functions/modexp.cu"
8-
#include "modnum/monty_mul.cu"
8+
#include "functions/multi_modexp.cu"
9+
#include "modnum/modnum_monty_redc.cu"
10+
#include "modnum/modnum_monty_cios.cu"
911

1012
using namespace std;
1113
using namespace cuFIXNUM;
@@ -37,17 +39,30 @@ struct sqr_wide {
3739
}
3840
};
3941

40-
template< typename fixnum >
42+
template< typename modnum >
4143
struct my_modexp {
44+
typedef typename modnum::fixnum fixnum;
45+
4246
__device__ void operator()(fixnum &z, fixnum x) {
43-
typedef modnum_monty_cios<fixnum> modnum;
4447
modexp<modnum> me(x, x);
4548
fixnum zz;
4649
me(zz, x);
4750
z = zz;
4851
};
4952
};
5053

54+
template< typename modnum >
55+
struct my_multi_modexp {
56+
typedef typename modnum::fixnum fixnum;
57+
58+
__device__ void operator()(fixnum &z, fixnum x) {
59+
multi_modexp<modnum> mme(x);
60+
fixnum zz;
61+
mme(zz, x, x);
62+
z = zz;
63+
};
64+
};
65+
5166
template< int fn_bytes, typename word_fixnum, template <typename> class Func >
5267
void bench(int nelts) {
5368
typedef warp_fixnum<fn_bytes, word_fixnum> fixnum;
@@ -101,6 +116,18 @@ void bench_func(const char *fn_name, int nelts) {
101116
puts("");
102117
}
103118

119+
template< typename fixnum >
120+
using modexp_redc = my_modexp< modnum_monty_redc<fixnum> >;
121+
122+
template< typename fixnum >
123+
using modexp_cios = my_modexp< modnum_monty_cios<fixnum> >;
124+
125+
template< typename fixnum >
126+
using multi_modexp_redc = my_multi_modexp< modnum_monty_redc<fixnum> >;
127+
128+
template< typename fixnum >
129+
using multi_modexp_cios = my_multi_modexp< modnum_monty_cios<fixnum> >;
130+
104131
int main(int argc, char *argv[]) {
105132
long m = 1;
106133
if (argc > 1)
@@ -112,7 +139,15 @@ int main(int argc, char *argv[]) {
112139
puts("");
113140
bench_func<sqr_wide>("sqr_wide", m);
114141
puts("");
115-
bench_func<my_modexp>("modexp", m / 100);
142+
bench_func<modexp_redc>("modexp redc", m / 100);
143+
puts("");
144+
bench_func<modexp_cios>("modexp cios", m / 100);
145+
puts("");
146+
147+
bench_func<modexp_redc>("multi modexp redc", m / 100);
148+
puts("");
149+
bench_func<modexp_cios>("multi modexp cios", m / 100);
150+
puts("");
116151

117152
return 0;
118153
}

‎src/fixnum/warp_fixnum.cu‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,8 @@ public:
5151
__device__ __forceinline__
5252
warp_fixnum() { }
5353

54+
// TODO: Shouldn't this be equivalent to the digit_to_fixnum() function
55+
// below?
5456
__device__ __forceinline__
5557
warp_fixnum(digit z) : x(z) { }
5658

‎src/functions/chinese.cu‎

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

33
#include "functions/quorem_preinv.cu"
44
#include "functions/multi_modexp.cu"
5-
#include "modnum/monty_mul.cu"
5+
#include "modnum/modnum_monty_cios.cu"
66

77
namespace cuFIXNUM {
88

‎src/functions/modexp.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#pragma once
22

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

66
namespace cuFIXNUM {
77

‎src/functions/multi_modexp.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#pragma once
22

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

66
namespace cuFIXNUM {
77

‎src/functions/paillier_decrypt.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,7 +4,7 @@
44
#include "functions/divexact.cu"
55
#include "functions/chinese.cu"
66
#include "functions/multi_modexp.cu"
7-
#include "modnum/monty_mul.cu"
7+
#include "modnum/modnum_monty_cios.cu"
88

99
namespace cuFIXNUM {
1010

‎src/functions/paillier_encrypt.cu‎

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

33
#include "functions/quorem_preinv.cu"
44
#include "functions/multi_modexp.cu"
5-
#include "modnum/monty_mul.cu"
5+
#include "modnum/modnum_monty_cios.cu"
66

77
namespace cuFIXNUM {
88

‎src/modnum/internal/monty.cu‎

Lines changed: 115 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,115 @@
1+
#pragma once
2+
3+
#include "functions/quorem_preinv.cu"
4+
5+
namespace cuFIXNUM {
6+
7+
namespace internal {
8+
9+
template< typename fixnum_ >
10+
class monty {
11+
public:
12+
typedef fixnum_ fixnum;
13+
typedef fixnum modnum;
14+
15+
__device__ monty(fixnum modulus);
16+
17+
__device__ void add(modnum &z, modnum x, modnum y) const {
18+
fixnum::add(z, x, y);
19+
if (fixnum::cmp(z, mod) >= 0)
20+
fixnum::sub(z, z, mod);
21+
}
22+
23+
__device__ void neg(modnum &z, modnum x) const {
24+
fixnum::sub(z, mod, x);
25+
}
26+
27+
__device__ void sub(modnum &z, modnum x, modnum y) const {
28+
fixnum my;
29+
neg(my, y);
30+
fixnum::add(z, x, my);
31+
if (fixnum::cmp(z, mod) >= 0)
32+
fixnum::sub(z, z, mod);
33+
}
34+
35+
/*
36+
* Return the Montgomery image of one.
37+
*/
38+
__device__ modnum one() const {
39+
return R_mod;
40+
}
41+
42+
/*
43+
* Return the Montgomery image of one.
44+
*/
45+
__device__ modnum zero() const {
46+
return fixnum::zero();
47+
}
48+
49+
// FIXME: Get rid of this hack
50+
int is_valid;
51+
52+
// Modulus for Monty arithmetic
53+
fixnum mod;
54+
// R_mod = 2^fixnum::BITS % mod
55+
modnum R_mod;
56+
// Rsqr = R^2 % mod
57+
modnum Rsqr_mod;
58+
59+
// TODO: We save this after using it in the constructor; work out
60+
// how to make it available for later use. For example, it could
61+
// be used to reduce arguments to modexp prior to the main
62+
// iteration.
63+
quorem_preinv<fixnum> modrem;
64+
65+
__device__ void normalise(modnum &x, int msb) const;
66+
};
67+
68+
69+
template< typename fixnum >
70+
__device__
71+
monty<fixnum>::monty(fixnum modulus)
72+
: mod(modulus), modrem(modulus)
73+
{
74+
// mod must be odd > 1 in order to calculate R^-1 mod "mod".
75+
// FIXME: Handle these errors properly
76+
if (fixnum::two_valuation(modulus) != 0 //fixnum::get(modulus, 0) & 1 == 0
77+
|| fixnum::cmp(modulus, fixnum::one()) == 0) {
78+
is_valid = 0;
79+
return;
80+
}
81+
is_valid = 1;
82+
83+
fixnum Rsqr_hi, Rsqr_lo;
84+
85+
// R_mod = R % mod
86+
modrem(R_mod, fixnum::one(), fixnum::zero());
87+
fixnum::sqr_wide(Rsqr_hi, Rsqr_lo, R_mod);
88+
// Rsqr_mod = R^2 % mod
89+
modrem(Rsqr_mod, Rsqr_hi, Rsqr_lo);
90+
}
91+
92+
/*
93+
* Let X = x + msb * 2^64. Then return X -= m if X > m.
94+
*
95+
* Assumes X < 2*m, i.e. msb = 0 or 1, and if msb = 1, then x < m.
96+
*/
97+
template< typename fixnum >
98+
__device__ void
99+
monty<fixnum>::normalise(modnum &x, int msb) const {
100+
typedef typename fixnum::digit digit;
101+
modnum r;
102+
digit br;
103+
104+
// br = 0 ==> x >= mod
105+
fixnum::sub_br(r, br, x, mod);
106+
if (msb || digit::is_zero(br)) {
107+
// If the msb was set, then we must have had to borrow.
108+
assert(!msb || msb == br);
109+
x = r;
110+
}
111+
}
112+
113+
} // End namespace internal
114+
115+
} // End namespace cuFIXNUM

0 commit comments

Comments
 (0)