commit 8b93b6635c36c5a5c3902af59bc0e614628b009d
parent f8cbe4d8d06f4be7d5d02a280492f7b34a9167d0
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Sat, 23 Dec 2023 18:47:57 +0100
Fix Montgomery
Though far from usable still, some progress is being made
Diffstat:
21 files changed, 3626 insertions(+), 782 deletions(-)
diff --git a/.DS_Store b/.DS_Store
Binary files differ.
diff --git a/source/.DS_Store b/source/.DS_Store
Binary files differ.
diff --git a/source/Makefile b/source/Makefile
@@ -1,3 +1,3 @@
all:
- gcc -g -O0 -D CL_TARGET_OPENCL_VERSION=100 -o gpu-verify lib-gpu-verify.c rsa-test.c montgomery-test.c reference-test.c util.c gmp.c -lgcrypt -lOpenCL -lm
+ gcc -g -O0 -D CL_TARGET_OPENCL_VERSION=100 -o gpu-verify lib-gpu-verify.c rsa-test.c montgomery-test.c reference-test.c util.c gmp.c -lgcrypt -lgmp -lOpenCL -lm
diff --git a/source/gmp.c b/source/gmp.c
@@ -1,7 +1,3 @@
-#include "gmp.h"
-
-
-
/* mini-gmp, a minimalistic implementation of a GNU GMP subset.
Contributed to the GNU project by Niels Möller
@@ -53,7 +49,7 @@ see https://www.gnu.org/licenses/. */
#include <stdlib.h>
#include <string.h>
-//#include "mini-gmp.h"
+#include "gmp.h"
#if !defined(MINI_GMP_DONT_USE_FLOAT_H)
#include <float.h>
@@ -61,7 +57,24 @@ see https://www.gnu.org/licenses/. */
/* Macros */
+#define GMP_LIMB_BITS (sizeof(mp_limb_t) * CHAR_BIT)
+
+#define GMP_LIMB_MAX ((mp_limb_t) ~ (mp_limb_t) 0)
+#define GMP_LIMB_HIGHBIT ((mp_limb_t) 1 << (GMP_LIMB_BITS - 1))
+
+#define GMP_HLIMB_BIT ((mp_limb_t) 1 << (GMP_LIMB_BITS / 2))
+#define GMP_LLIMB_MASK (GMP_HLIMB_BIT - 1)
+#define GMP_ULONG_BITS (sizeof(unsigned long) * CHAR_BIT)
+#define GMP_ULONG_HIGHBIT ((unsigned long) 1 << (GMP_ULONG_BITS - 1))
+
+#define GMP_ABS(x) ((x) >= 0 ? (x) : -(x))
+#define GMP_NEG_CAST(T,x) (-((T)((x) + 1) - 1))
+
+#define GMP_MIN(a, b) ((a) < (b) ? (a) : (b))
+#define GMP_MAX(a, b) ((a) > (b) ? (a) : (b))
+
+#define GMP_CMP(a,b) (((a) > (b)) - ((a) < (b)))
#if defined(DBL_MANT_DIG) && FLT_RADIX == 2
#define GMP_DBL_MANT_BITS DBL_MANT_DIG
@@ -72,91 +85,91 @@ see https://www.gnu.org/licenses/. */
/* Return non-zero if xp,xsize and yp,ysize overlap.
If xp+xsize<=yp there's no overlap, or if yp+ysize<=xp there's no
overlap. If both these are false, there's an overlap. */
-#define GMP_MPN_OVERLAP_P(xp, xsize, yp, ysize) \
+#define GMP_MPN_OVERLAP_P(xp, xsize, yp, ysize) \
((xp) + (xsize) > (yp) && (yp) + (ysize) > (xp))
#define gmp_assert_nocarry(x) do { \
- mp_limb_t __cy = (x); \
- assert (__cy == 0); \
- (void) (__cy); \
+ mp_limb_t __cy = (x); \
+ assert (__cy == 0); \
+ (void) (__cy); \
} while (0)
-#define gmp_clz(count, x) do { \
- mp_limb_t __clz_x = (x); \
- unsigned __clz_c = 0; \
- int LOCAL_SHIFT_BITS = 8; \
- if (GMP_LIMB_BITS > LOCAL_SHIFT_BITS) \
- for (; \
- (__clz_x & ((mp_limb_t) 0xff << (GMP_LIMB_BITS - 8))) == 0; \
- __clz_c += 8) \
- { __clz_x <<= LOCAL_SHIFT_BITS; } \
- for (; (__clz_x & GMP_LIMB_HIGHBIT) == 0; __clz_c++) \
- __clz_x <<= 1; \
- (count) = __clz_c; \
+#define gmp_clz(count, x) do { \
+ mp_limb_t __clz_x = (x); \
+ unsigned __clz_c = 0; \
+ int LOCAL_SHIFT_BITS = 8; \
+ if (GMP_LIMB_BITS > LOCAL_SHIFT_BITS) \
+ for (; \
+ (__clz_x & ((mp_limb_t) 0xff << (GMP_LIMB_BITS - 8))) == 0; \
+ __clz_c += 8) \
+ { __clz_x <<= LOCAL_SHIFT_BITS; } \
+ for (; (__clz_x & GMP_LIMB_HIGHBIT) == 0; __clz_c++) \
+ __clz_x <<= 1; \
+ (count) = __clz_c; \
} while (0)
-#define gmp_ctz(count, x) do { \
- mp_limb_t __ctz_x = (x); \
- unsigned __ctz_c = 0; \
- gmp_clz (__ctz_c, __ctz_x & - __ctz_x); \
- (count) = GMP_LIMB_BITS - 1 - __ctz_c; \
+#define gmp_ctz(count, x) do { \
+ mp_limb_t __ctz_x = (x); \
+ unsigned __ctz_c = 0; \
+ gmp_clz (__ctz_c, __ctz_x & - __ctz_x); \
+ (count) = GMP_LIMB_BITS - 1 - __ctz_c; \
} while (0)
#define gmp_add_ssaaaa(sh, sl, ah, al, bh, bl) \
- do { \
- mp_limb_t __x; \
- __x = (al) + (bl); \
- (sh) = (ah) + (bh) + (__x < (al)); \
- (sl) = __x; \
+ do { \
+ mp_limb_t __x; \
+ __x = (al) + (bl); \
+ (sh) = (ah) + (bh) + (__x < (al)); \
+ (sl) = __x; \
} while (0)
#define gmp_sub_ddmmss(sh, sl, ah, al, bh, bl) \
- do { \
- mp_limb_t __x; \
- __x = (al) - (bl); \
- (sh) = (ah) - (bh) - ((al) < (bl)); \
- (sl) = __x; \
+ do { \
+ mp_limb_t __x; \
+ __x = (al) - (bl); \
+ (sh) = (ah) - (bh) - ((al) < (bl)); \
+ (sl) = __x; \
} while (0)
-#define gmp_umul_ppmm(w1, w0, u, v) \
- do { \
- int LOCAL_GMP_LIMB_BITS = GMP_LIMB_BITS; \
- if (sizeof(unsigned int) * CHAR_BIT >= 2 * GMP_LIMB_BITS) \
- { \
- unsigned int __ww = (unsigned int) (u) * (v); \
- w0 = (mp_limb_t) __ww; \
- w1 = (mp_limb_t) (__ww >> LOCAL_GMP_LIMB_BITS); \
- } \
- else if (GMP_ULONG_BITS >= 2 * GMP_LIMB_BITS) \
- { \
- unsigned long int __ww = (unsigned long int) (u) * (v); \
- w0 = (mp_limb_t) __ww; \
- w1 = (mp_limb_t) (__ww >> LOCAL_GMP_LIMB_BITS); \
- } \
- else { \
- mp_limb_t __x0, __x1, __x2, __x3; \
- unsigned __ul, __vl, __uh, __vh; \
- mp_limb_t __u = (u), __v = (v); \
- assert (sizeof (unsigned) * 2 >= sizeof (mp_limb_t)); \
- \
- __ul = __u & GMP_LLIMB_MASK; \
- __uh = __u >> (GMP_LIMB_BITS / 2); \
- __vl = __v & GMP_LLIMB_MASK; \
- __vh = __v >> (GMP_LIMB_BITS / 2); \
- \
- __x0 = (mp_limb_t) __ul * __vl; \
- __x1 = (mp_limb_t) __ul * __vh; \
- __x2 = (mp_limb_t) __uh * __vl; \
- __x3 = (mp_limb_t) __uh * __vh; \
- \
- __x1 += __x0 >> (GMP_LIMB_BITS / 2);/* this can't give carry */ \
- __x1 += __x2; /* but this indeed can */ \
- if (__x1 < __x2) /* did we get it? */ \
- __x3 += GMP_HLIMB_BIT; /* yes, add it in the proper pos. */ \
- \
- (w1) = __x3 + (__x1 >> (GMP_LIMB_BITS / 2)); \
- (w0) = (__x1 << (GMP_LIMB_BITS / 2)) + (__x0 & GMP_LLIMB_MASK); \
- } \
+#define gmp_umul_ppmm(w1, w0, u, v) \
+ do { \
+ int LOCAL_GMP_LIMB_BITS = GMP_LIMB_BITS; \
+ if (sizeof(unsigned int) * CHAR_BIT >= 2 * GMP_LIMB_BITS) \
+ { \
+ unsigned int __ww = (unsigned int) (u) * (v); \
+ w0 = (mp_limb_t) __ww; \
+ w1 = (mp_limb_t) (__ww >> LOCAL_GMP_LIMB_BITS); \
+ } \
+ else if (GMP_ULONG_BITS >= 2 * GMP_LIMB_BITS) \
+ { \
+ unsigned long int __ww = (unsigned long int) (u) * (v); \
+ w0 = (mp_limb_t) __ww; \
+ w1 = (mp_limb_t) (__ww >> LOCAL_GMP_LIMB_BITS); \
+ } \
+ else { \
+ mp_limb_t __x0, __x1, __x2, __x3; \
+ unsigned __ul, __vl, __uh, __vh; \
+ mp_limb_t __u = (u), __v = (v); \
+ assert (sizeof (unsigned) * 2 >= sizeof (mp_limb_t)); \
+ \
+ __ul = __u & GMP_LLIMB_MASK; \
+ __uh = __u >> (GMP_LIMB_BITS / 2); \
+ __vl = __v & GMP_LLIMB_MASK; \
+ __vh = __v >> (GMP_LIMB_BITS / 2); \
+ \
+ __x0 = (mp_limb_t) __ul * __vl; \
+ __x1 = (mp_limb_t) __ul * __vh; \
+ __x2 = (mp_limb_t) __uh * __vl; \
+ __x3 = (mp_limb_t) __uh * __vh; \
+ \
+ __x1 += __x0 >> (GMP_LIMB_BITS / 2);/* this can't give carry */ \
+ __x1 += __x2; /* but this indeed can */ \
+ if (__x1 < __x2) /* did we get it? */ \
+ __x3 += GMP_HLIMB_BIT; /* yes, add it in the proper pos. */ \
+ \
+ (w1) = __x3 + (__x1 >> (GMP_LIMB_BITS / 2)); \
+ (w0) = (__x1 << (GMP_LIMB_BITS / 2)) + (__x0 & GMP_LLIMB_MASK); \
+ } \
} while (0)
/* If mp_limb_t is of size smaller than int, plain u*v implies
@@ -166,106 +179,106 @@ see https://www.gnu.org/licenses/. */
#define gmp_umullo_limb(u, v) \
((sizeof(mp_limb_t) >= sizeof(int)) ? (u)*(v) : (unsigned int)(u) * (v))
-#define gmp_udiv_qrnnd_preinv(q, r, nh, nl, d, di) \
- do { \
- mp_limb_t _qh, _ql, _r, _mask; \
- gmp_umul_ppmm (_qh, _ql, (nh), (di)); \
- gmp_add_ssaaaa (_qh, _ql, _qh, _ql, (nh) + 1, (nl)); \
- _r = (nl) - gmp_umullo_limb (_qh, (d)); \
- _mask = -(mp_limb_t) (_r > _ql); /* both > and >= are OK */ \
- _qh += _mask; \
- _r += _mask & (d); \
- if (_r >= (d)) \
- { \
- _r -= (d); \
- _qh++; \
- } \
- \
- (r) = _r; \
- (q) = _qh; \
+#define gmp_udiv_qrnnd_preinv(q, r, nh, nl, d, di) \
+ do { \
+ mp_limb_t _qh, _ql, _r, _mask; \
+ gmp_umul_ppmm (_qh, _ql, (nh), (di)); \
+ gmp_add_ssaaaa (_qh, _ql, _qh, _ql, (nh) + 1, (nl)); \
+ _r = (nl) - gmp_umullo_limb (_qh, (d)); \
+ _mask = -(mp_limb_t) (_r > _ql); /* both > and >= are OK */ \
+ _qh += _mask; \
+ _r += _mask & (d); \
+ if (_r >= (d)) \
+ { \
+ _r -= (d); \
+ _qh++; \
+ } \
+ \
+ (r) = _r; \
+ (q) = _qh; \
} while (0)
-#define gmp_udiv_qr_3by2(q, r1, r0, n2, n1, n0, d1, d0, dinv) \
- do { \
- mp_limb_t _q0, _t1, _t0, _mask; \
- gmp_umul_ppmm ((q), _q0, (n2), (dinv)); \
- gmp_add_ssaaaa ((q), _q0, (q), _q0, (n2), (n1)); \
- \
- /* Compute the two most significant limbs of n - q'd */ \
- (r1) = (n1) - gmp_umullo_limb ((d1), (q)); \
- gmp_sub_ddmmss ((r1), (r0), (r1), (n0), (d1), (d0)); \
- gmp_umul_ppmm (_t1, _t0, (d0), (q)); \
- gmp_sub_ddmmss ((r1), (r0), (r1), (r0), _t1, _t0); \
- (q)++; \
- \
- /* Conditionally adjust q and the remainders */ \
- _mask = - (mp_limb_t) ((r1) >= _q0); \
- (q) += _mask; \
+#define gmp_udiv_qr_3by2(q, r1, r0, n2, n1, n0, d1, d0, dinv) \
+ do { \
+ mp_limb_t _q0, _t1, _t0, _mask; \
+ gmp_umul_ppmm ((q), _q0, (n2), (dinv)); \
+ gmp_add_ssaaaa ((q), _q0, (q), _q0, (n2), (n1)); \
+ \
+ /* Compute the two most significant limbs of n - q'd */ \
+ (r1) = (n1) - gmp_umullo_limb ((d1), (q)); \
+ gmp_sub_ddmmss ((r1), (r0), (r1), (n0), (d1), (d0)); \
+ gmp_umul_ppmm (_t1, _t0, (d0), (q)); \
+ gmp_sub_ddmmss ((r1), (r0), (r1), (r0), _t1, _t0); \
+ (q)++; \
+ \
+ /* Conditionally adjust q and the remainders */ \
+ _mask = - (mp_limb_t) ((r1) >= _q0); \
+ (q) += _mask; \
gmp_add_ssaaaa ((r1), (r0), (r1), (r0), _mask & (d1), _mask & (d0)); \
- if ((r1) >= (d1)) \
- { \
- if ((r1) > (d1) || (r0) >= (d0)) \
- { \
- (q)++; \
- gmp_sub_ddmmss ((r1), (r0), (r1), (r0), (d1), (d0)); \
- } \
- } \
+ if ((r1) >= (d1)) \
+ { \
+ if ((r1) > (d1) || (r0) >= (d0)) \
+ { \
+ (q)++; \
+ gmp_sub_ddmmss ((r1), (r0), (r1), (r0), (d1), (d0)); \
+ } \
+ } \
} while (0)
/* Swap macros. */
-#define MP_LIMB_T_SWAP(x, y) \
- do { \
- mp_limb_t __mp_limb_t_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mp_limb_t_swap__tmp; \
+#define MP_LIMB_T_SWAP(x, y) \
+ do { \
+ mp_limb_t __mp_limb_t_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mp_limb_t_swap__tmp; \
} while (0)
-#define MP_SIZE_T_SWAP(x, y) \
- do { \
- mp_size_t __mp_size_t_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mp_size_t_swap__tmp; \
+#define MP_SIZE_T_SWAP(x, y) \
+ do { \
+ mp_size_t __mp_size_t_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mp_size_t_swap__tmp; \
} while (0)
-#define MP_BITCNT_T_SWAP(x,y) \
- do { \
- mp_bitcnt_t __mp_bitcnt_t_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mp_bitcnt_t_swap__tmp; \
+#define MP_BITCNT_T_SWAP(x,y) \
+ do { \
+ mp_bitcnt_t __mp_bitcnt_t_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mp_bitcnt_t_swap__tmp; \
} while (0)
-#define MP_PTR_SWAP(x, y) \
- do { \
- mp_ptr __mp_ptr_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mp_ptr_swap__tmp; \
+#define MP_PTR_SWAP(x, y) \
+ do { \
+ mp_ptr __mp_ptr_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mp_ptr_swap__tmp; \
} while (0)
-#define MP_SRCPTR_SWAP(x, y) \
- do { \
- mp_srcptr __mp_srcptr_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mp_srcptr_swap__tmp; \
+#define MP_SRCPTR_SWAP(x, y) \
+ do { \
+ mp_srcptr __mp_srcptr_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mp_srcptr_swap__tmp; \
} while (0)
-#define MPN_PTR_SWAP(xp,xs, yp,ys) \
- do { \
- MP_PTR_SWAP (xp, yp); \
- MP_SIZE_T_SWAP (xs, ys); \
+#define MPN_PTR_SWAP(xp,xs, yp,ys) \
+ do { \
+ MP_PTR_SWAP (xp, yp); \
+ MP_SIZE_T_SWAP (xs, ys); \
} while(0)
-#define MPN_SRCPTR_SWAP(xp,xs, yp,ys) \
- do { \
- MP_SRCPTR_SWAP (xp, yp); \
- MP_SIZE_T_SWAP (xs, ys); \
+#define MPN_SRCPTR_SWAP(xp,xs, yp,ys) \
+ do { \
+ MP_SRCPTR_SWAP (xp, yp); \
+ MP_SIZE_T_SWAP (xs, ys); \
} while(0)
-#define MPZ_PTR_SWAP(x, y) \
- do { \
- mpz_ptr __mpz_ptr_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mpz_ptr_swap__tmp; \
+#define MPZ_PTR_SWAP(x, y) \
+ do { \
+ mpz_ptr __mpz_ptr_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mpz_ptr_swap__tmp; \
} while (0)
-#define MPZ_SRCPTR_SWAP(x, y) \
- do { \
- mpz_srcptr __mpz_srcptr_swap__tmp = (x); \
- (x) = (y); \
- (y) = __mpz_srcptr_swap__tmp; \
+#define MPZ_SRCPTR_SWAP(x, y) \
+ do { \
+ mpz_srcptr __mpz_srcptr_swap__tmp = (x); \
+ (x) = (y); \
+ (y) = __mpz_srcptr_swap__tmp; \
} while (0)
const int mp_bits_per_limb = GMP_LIMB_BITS;
@@ -309,7 +322,7 @@ gmp_default_realloc (void *old, size_t unused_old_size, size_t new_size)
static void
gmp_default_free (void *p, size_t unused_size)
{
- //free (p);
+ free (p);
}
static void * (*gmp_allocate_func) (size_t) = gmp_default_alloc;
@@ -318,8 +331,8 @@ static void (*gmp_free_func) (void *, size_t) = gmp_default_free;
void
mp_get_memory_functions (void *(**alloc_func) (size_t),
- void *(**realloc_func) (void *, size_t, size_t),
- void (**free_func) (void *, size_t))
+ void *(**realloc_func) (void *, size_t, size_t),
+ void (**free_func) (void *, size_t))
{
if (alloc_func)
*alloc_func = gmp_allocate_func;
@@ -333,8 +346,8 @@ mp_get_memory_functions (void *(**alloc_func) (size_t),
void
mp_set_memory_functions (void *(*alloc_func) (size_t),
- void *(*realloc_func) (void *, size_t, size_t),
- void (*free_func) (void *, size_t))
+ void *(*realloc_func) (void *, size_t, size_t),
+ void (*free_func) (void *, size_t))
{
if (!alloc_func)
alloc_func = gmp_default_alloc;
@@ -395,7 +408,7 @@ mpn_cmp (mp_srcptr ap, mp_srcptr bp, mp_size_t n)
while (--n >= 0)
{
if (ap[n] != bp[n])
- return ap[n] > bp[n] ? 1 : -1;
+ return ap[n] > bp[n] ? 1 : -1;
}
return 0;
}
@@ -704,7 +717,7 @@ mpn_rshift (mp_ptr rp, mp_srcptr up, mp_size_t n, unsigned int cnt)
static mp_bitcnt_t
mpn_common_scan (mp_limb_t limb, mp_size_t i, mp_srcptr up, mp_size_t un,
- mp_limb_t ux)
+ mp_limb_t ux)
{
unsigned cnt;
@@ -715,7 +728,7 @@ mpn_common_scan (mp_limb_t limb, mp_size_t i, mp_srcptr up, mp_size_t un,
{
i++;
if (i == un)
- return (ux == 0 ? ~(mp_bitcnt_t) 0 : un * GMP_LIMB_BITS);
+ return (ux == 0 ? ~(mp_bitcnt_t) 0 : un * GMP_LIMB_BITS);
limb = ux ^ up[i];
}
gmp_ctz (cnt, limb);
@@ -729,7 +742,7 @@ mpn_scan1 (mp_srcptr ptr, mp_bitcnt_t bit)
i = bit / GMP_LIMB_BITS;
return mpn_common_scan ( ptr[i] & (GMP_LIMB_MAX << (bit % GMP_LIMB_BITS)),
- i, ptr, i, 0);
+ i, ptr, i, 0);
}
mp_bitcnt_t
@@ -739,7 +752,7 @@ mpn_scan0 (mp_srcptr ptr, mp_bitcnt_t bit)
i = bit / GMP_LIMB_BITS;
return mpn_common_scan (~ptr[i] & (GMP_LIMB_MAX << (bit % GMP_LIMB_BITS)),
- i, ptr, i, GMP_LIMB_MAX);
+ i, ptr, i, GMP_LIMB_MAX);
}
void
@@ -756,7 +769,7 @@ mpn_neg (mp_ptr rp, mp_srcptr up, mp_size_t n)
{
*rp = 0;
if (!--n)
- return 0;
+ return 0;
++up; ++rp;
}
*rp = - *up;
@@ -794,7 +807,7 @@ mpn_invert_3by2 (mp_limb_t u1, mp_limb_t u0)
/* Adjust to get a half-limb 3/2 inverse, i.e., we want
qh' = floor( (b^3 - 1) / u) - b = floor ((b^3 - b u - 1) / u
- = floor( (b (~u) + b-1) / u),
+ = floor( (b (~u) + b-1) / u),
and the remainder
@@ -812,14 +825,14 @@ mpn_invert_3by2 (mp_limb_t u1, mp_limb_t u0)
/* Adjustment steps taken from udiv_qrnnd_c */
if (r < p)
{
- qh--;
- r += u1;
- if (r >= u1) /* i.e. we didn't get carry when adding to r */
- if (r < p)
- {
- qh--;
- r += u1;
- }
+ qh--;
+ r += u1;
+ if (r >= u1) /* i.e. we didn't get carry when adding to r */
+ if (r < p)
+ {
+ qh--;
+ r += u1;
+ }
}
r -= p;
@@ -840,14 +853,14 @@ mpn_invert_3by2 (mp_limb_t u1, mp_limb_t u0)
if (r >= (GMP_LIMB_MAX & (p << (GMP_LIMB_BITS / 2))))
{
- ql--;
- r += u1;
+ ql--;
+ r += u1;
}
m = ((mp_limb_t) qh << (GMP_LIMB_BITS / 2)) + ql;
if (r >= u1)
{
- m++;
- r -= u1;
+ m++;
+ r -= u1;
}
}
@@ -859,22 +872,22 @@ mpn_invert_3by2 (mp_limb_t u1, mp_limb_t u0)
r = ~r;
r += u0;
if (r < u0)
- {
- m--;
- if (r >= u1)
- {
- m--;
- r -= u1;
- }
- r -= u1;
- }
+ {
+ m--;
+ if (r >= u1)
+ {
+ m--;
+ r -= u1;
+ }
+ r -= u1;
+ }
gmp_umul_ppmm (th, tl, u0, m);
r += th;
if (r < th)
- {
- m--;
- m -= ((r > u1) | ((r == u1) & (tl > u0)));
- }
+ {
+ m--;
+ m -= ((r > u1) | ((r == u1) & (tl > u0)));
+ }
}
return m;
@@ -904,7 +917,7 @@ mpn_div_qr_1_invert (struct gmp_div_inverse *inv, mp_limb_t d)
static void
mpn_div_qr_2_invert (struct gmp_div_inverse *inv,
- mp_limb_t d1, mp_limb_t d0)
+ mp_limb_t d1, mp_limb_t d0)
{
unsigned shift;
@@ -923,7 +936,7 @@ mpn_div_qr_2_invert (struct gmp_div_inverse *inv,
static void
mpn_div_qr_invert (struct gmp_div_inverse *inv,
- mp_srcptr dp, mp_size_t dn)
+ mp_srcptr dp, mp_size_t dn)
{
assert (dn > 0);
@@ -942,10 +955,10 @@ mpn_div_qr_invert (struct gmp_div_inverse *inv,
gmp_clz (shift, d1);
inv->shift = shift;
if (shift > 0)
- {
- d1 = (d1 << shift) | (d0 >> (GMP_LIMB_BITS - shift));
- d0 = (d0 << shift) | (dp[dn-3] >> (GMP_LIMB_BITS - shift));
- }
+ {
+ d1 = (d1 << shift) | (d0 >> (GMP_LIMB_BITS - shift));
+ d0 = (d0 << shift) | (dp[dn-3] >> (GMP_LIMB_BITS - shift));
+ }
inv->d1 = d1;
inv->d0 = d0;
inv->di = mpn_invert_3by2 (d1, d0);
@@ -956,14 +969,12 @@ mpn_div_qr_invert (struct gmp_div_inverse *inv,
the sbpi1_div_* functions. */
static mp_limb_t
mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
- const struct gmp_div_inverse *inv)
+ const struct gmp_div_inverse *inv)
{
mp_limb_t d, di;
mp_limb_t r;
mp_ptr tp = NULL;
mp_size_t tn = 0;
-
-
if (inv->shift > 0)
{
@@ -971,11 +982,8 @@ mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
tp = qp;
if (!tp)
{
-
- printf("%ld\n",nn);
- tn = nn;
- // tp = tn->_mp_d;
- // tp = gmp_alloc_limbs (tn);
+ tn = nn;
+ tp = gmp_alloc_limbs (tn);
}
r = mpn_lshift (tp, np, nn, inv->shift);
np = tp;
@@ -991,17 +999,17 @@ mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
gmp_udiv_qrnnd_preinv (q, r, r, np[nn], d, di);
if (qp)
- qp[nn] = q;
+ qp[nn] = q;
}
- //if (tn) {}
- // gmp_free_limbs (tp, tn);
+ if (tn)
+ gmp_free_limbs (tp, tn);
return r >> inv->shift;
}
static void
mpn_div_qr_2_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
- const struct gmp_div_inverse *inv)
+ const struct gmp_div_inverse *inv)
{
unsigned shift;
mp_size_t i;
@@ -1028,7 +1036,7 @@ mpn_div_qr_2_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
gmp_udiv_qr_3by2 (q, r1, r0, r1, r0, n0, d1, d0, di);
if (qp)
- qp[i] = q;
+ qp[i] = q;
}
while (--i >= 0);
@@ -1045,9 +1053,9 @@ mpn_div_qr_2_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
static void
mpn_div_qr_pi1 (mp_ptr qp,
- mp_ptr np, mp_size_t nn, mp_limb_t n1,
- mp_srcptr dp, mp_size_t dn,
- mp_limb_t dinv)
+ mp_ptr np, mp_size_t nn, mp_limb_t n1,
+ mp_srcptr dp, mp_size_t dn,
+ mp_limb_t dinv)
{
mp_size_t i;
@@ -1074,32 +1082,32 @@ mpn_div_qr_pi1 (mp_ptr qp,
mp_limb_t n0 = np[dn-1+i];
if (n1 == d1 && n0 == d0)
- {
- q = GMP_LIMB_MAX;
- mpn_submul_1 (np+i, dp, dn, q);
- n1 = np[dn-1+i]; /* update n1, last loop's value will now be invalid */
- }
+ {
+ q = GMP_LIMB_MAX;
+ mpn_submul_1 (np+i, dp, dn, q);
+ n1 = np[dn-1+i]; /* update n1, last loop's value will now be invalid */
+ }
else
- {
- gmp_udiv_qr_3by2 (q, n1, n0, n1, n0, np[dn-2+i], d1, d0, dinv);
+ {
+ gmp_udiv_qr_3by2 (q, n1, n0, n1, n0, np[dn-2+i], d1, d0, dinv);
- cy = mpn_submul_1 (np + i, dp, dn-2, q);
+ cy = mpn_submul_1 (np + i, dp, dn-2, q);
- cy1 = n0 < cy;
- n0 = n0 - cy;
- cy = n1 < cy1;
- n1 = n1 - cy1;
- np[dn-2+i] = n0;
+ cy1 = n0 < cy;
+ n0 = n0 - cy;
+ cy = n1 < cy1;
+ n1 = n1 - cy1;
+ np[dn-2+i] = n0;
- if (cy != 0)
- {
- n1 += d1 + mpn_add_n (np + i, np + i, dp, dn - 1);
- q--;
- }
- }
+ if (cy != 0)
+ {
+ n1 += d1 + mpn_add_n (np + i, np + i, dp, dn - 1);
+ q--;
+ }
+ }
if (qp)
- qp[i] = q;
+ qp[i] = q;
}
while (--i >= 0);
@@ -1108,8 +1116,8 @@ mpn_div_qr_pi1 (mp_ptr qp,
static void
mpn_div_qr_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
- mp_srcptr dp, mp_size_t dn,
- const struct gmp_div_inverse *inv)
+ mp_srcptr dp, mp_size_t dn,
+ const struct gmp_div_inverse *inv)
{
assert (dn > 0);
assert (nn >= dn);
@@ -1129,14 +1137,14 @@ mpn_div_qr_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
shift = inv->shift;
if (shift > 0)
- nh = mpn_lshift (np, np, nn, shift);
+ nh = mpn_lshift (np, np, nn, shift);
else
- nh = 0;
+ nh = 0;
mpn_div_qr_pi1 (qp, np, nn, nh, dp, dn, inv->di);
if (shift > 0)
- gmp_assert_nocarry (mpn_rshift (np, np, dn, shift));
+ gmp_assert_nocarry (mpn_rshift (np, np, dn, shift));
}
}
@@ -1157,8 +1165,8 @@ mpn_div_qr (mp_ptr qp, mp_ptr np, mp_size_t nn, mp_srcptr dp, mp_size_t dn)
dp = tp;
}
mpn_div_qr_preinv (qp, np, nn, dp, dn, &inv);
- if (tp) {}
- //gmp_free_limbs (tp, dn);
+ if (tp)
+ gmp_free_limbs (tp, dn);
}
@@ -1222,7 +1230,7 @@ mpn_get_str_bits (unsigned char *sp, unsigned bits, mp_srcptr up, mp_size_t un)
unsigned shift;
sn = ((un - 1) * GMP_LIMB_BITS + mpn_limb_size_in_base_2 (up[un-1])
- + bits - 1) / bits;
+ + bits - 1) / bits;
mask = (1U << bits) - 1;
@@ -1233,10 +1241,10 @@ mpn_get_str_bits (unsigned char *sp, unsigned bits, mp_srcptr up, mp_size_t un)
shift += bits;
if (shift >= GMP_LIMB_BITS && ++i < un)
- {
- shift -= GMP_LIMB_BITS;
- digit |= up[i] << (bits - shift);
- }
+ {
+ shift -= GMP_LIMB_BITS;
+ digit |= up[i] << (bits - shift);
+ }
sp[j] = digit & mask;
}
return sn;
@@ -1246,7 +1254,7 @@ mpn_get_str_bits (unsigned char *sp, unsigned bits, mp_srcptr up, mp_size_t un)
the end. */
static size_t
mpn_limb_get_str (unsigned char *sp, mp_limb_t w,
- const struct gmp_div_inverse *binv)
+ const struct gmp_div_inverse *binv)
{
mp_size_t i;
for (i = 0; w > 0; i++)
@@ -1267,8 +1275,8 @@ mpn_limb_get_str (unsigned char *sp, mp_limb_t w,
static size_t
mpn_get_str_other (unsigned char *sp,
- int base, const struct mpn_base_info *info,
- mp_ptr up, mp_size_t un)
+ int base, const struct mpn_base_info *info,
+ mp_ptr up, mp_size_t un)
{
struct gmp_div_inverse binv;
size_t sn;
@@ -1284,16 +1292,16 @@ mpn_get_str_other (unsigned char *sp,
mpn_div_qr_1_invert (&bbinv, info->bb);
do
- {
- mp_limb_t w;
- size_t done;
- w = mpn_div_qr_1_preinv (up, up, un, &bbinv);
- un -= (up[un-1] == 0);
- done = mpn_limb_get_str (sp + sn, w, &binv);
-
- for (sn += done; done < info->exp; done++)
- sp[sn++] = 0;
- }
+ {
+ mp_limb_t w;
+ size_t done;
+ w = mpn_div_qr_1_preinv (up, up, un, &bbinv);
+ un -= (up[un-1] == 0);
+ done = mpn_limb_get_str (sp + sn, w, &binv);
+
+ for (sn += done; done < info->exp; done++)
+ sp[sn++] = 0;
+ }
while (un > 1);
}
sn += mpn_limb_get_str (sp + sn, up[0], &binv);
@@ -1331,7 +1339,7 @@ mpn_get_str (unsigned char *sp, int base, mp_ptr up, mp_size_t un)
static mp_size_t
mpn_set_str_bits (mp_ptr rp, const unsigned char *sp, size_t sn,
- unsigned bits)
+ unsigned bits)
{
mp_size_t rn;
mp_limb_t limb;
@@ -1342,13 +1350,13 @@ mpn_set_str_bits (mp_ptr rp, const unsigned char *sp, size_t sn,
limb |= (mp_limb_t) sp[sn] << shift;
shift += bits;
if (shift >= GMP_LIMB_BITS)
- {
- shift -= GMP_LIMB_BITS;
- rp[rn++] = limb;
- /* Next line is correct also if shift == 0,
- bits == 8, and mp_limb_t == unsigned char. */
- limb = (unsigned int) sp[sn] >> (bits - shift);
- }
+ {
+ shift -= GMP_LIMB_BITS;
+ rp[rn++] = limb;
+ /* Next line is correct also if shift == 0,
+ bits == 8, and mp_limb_t == unsigned char. */
+ limb = (unsigned int) sp[sn] >> (bits - shift);
+ }
}
if (limb != 0)
rp[rn++] = limb;
@@ -1361,7 +1369,7 @@ mpn_set_str_bits (mp_ptr rp, const unsigned char *sp, size_t sn,
case a single zero limb is written at *RP, and 1 is returned. */
static mp_size_t
mpn_set_str_other (mp_ptr rp, const unsigned char *sp, size_t sn,
- mp_limb_t b, const struct mpn_base_info *info)
+ mp_limb_t b, const struct mpn_base_info *info)
{
mp_size_t rn;
mp_limb_t w;
@@ -1385,12 +1393,12 @@ mpn_set_str_other (mp_ptr rp, const unsigned char *sp, size_t sn,
w = sp[j++];
for (k = 1; k < info->exp; k++)
- w = w * b + sp[j++];
+ w = w * b + sp[j++];
cy = mpn_mul_1 (rp, rp, rn, info->bb);
cy += mpn_add_1 (rp, rp, rn, w);
if (cy > 0)
- rp[rn++] = cy;
+ rp[rn++] = cy;
}
assert (j == sn);
@@ -1422,11 +1430,11 @@ mpn_set_str (mp_ptr rp, const unsigned char *sp, size_t sn, int base)
void
mpz_init (mpz_t r)
{
- //static const mp_limb_t dummy_limb = GMP_LIMB_MAX & 0xc1a0;
+ static const mp_limb_t dummy_limb = GMP_LIMB_MAX & 0xc1a0;
r->_mp_alloc = 0;
r->_mp_size = 0;
-// r->_mp_d = (mp_ptr) &dummy_limb;
+ r->_mp_d = (mp_ptr) &dummy_limb;
}
/* The utility of this function is a bit limited, since many functions
@@ -1436,19 +1444,19 @@ mpz_init2 (mpz_t r, mp_bitcnt_t bits)
{
mp_size_t rn;
- bits -= (bits != 0); /* Round down, except if 0 */
+ bits -= (bits != 0); /* Round down, except if 0 */
rn = 1 + bits / GMP_LIMB_BITS;
r->_mp_alloc = rn;
r->_mp_size = 0;
- // r->_mp_d = gmp_alloc_limbs (rn);
+ r->_mp_d = gmp_alloc_limbs (rn);
}
void
mpz_clear (mpz_t r)
{
- // if (r->_mp_alloc)
- // gmp_free_limbs (r->_mp_d, r->_mp_alloc);
+ if (r->_mp_alloc)
+ gmp_free_limbs (r->_mp_d, r->_mp_alloc);
}
static mp_ptr
@@ -1456,11 +1464,11 @@ mpz_realloc (mpz_t r, mp_size_t size)
{
size = GMP_MAX (size, 1);
- if (r->_mp_alloc) {}
- // r->_mp_d = gmp_realloc_limbs (r->_mp_d, r->_mp_alloc, size);
- else {}
- // r->_mp_d = gmp_alloc_limbs (size);
- //r->_mp_alloc = size;
+ if (r->_mp_alloc)
+ r->_mp_d = gmp_realloc_limbs (r->_mp_d, r->_mp_alloc, size);
+ else
+ r->_mp_d = gmp_alloc_limbs (size);
+ r->_mp_alloc = size;
if (GMP_ABS (r->_mp_size) > size)
r->_mp_size = 0;
@@ -1469,7 +1477,9 @@ mpz_realloc (mpz_t r, mp_size_t size)
}
/* Realloc for an mpz_t WHAT if it has less than NEEDED limbs. */
-#define MPZ_REALLOC(z,n) (z)->_mp_d
+#define MPZ_REALLOC(z,n) ((n) > (z)->_mp_alloc \
+ ? mpz_realloc(z,n) \
+ : (z)->_mp_d)
/* MPZ assignment and basic conversions. */
void
@@ -1480,8 +1490,8 @@ mpz_set_si (mpz_t r, signed long int x)
else /* (x < 0) */
if (GMP_LIMB_BITS < GMP_ULONG_BITS)
{
- mpz_set_ui (r, GMP_NEG_CAST (unsigned long int, x));
- mpz_neg (r, r);
+ mpz_set_ui (r, GMP_NEG_CAST (unsigned long int, x));
+ mpz_neg (r, r);
}
else
{
@@ -1498,14 +1508,14 @@ mpz_set_ui (mpz_t r, unsigned long int x)
r->_mp_size = 1;
MPZ_REALLOC (r, 1)[0] = x;
if (GMP_LIMB_BITS < GMP_ULONG_BITS)
- {
- int LOCAL_GMP_LIMB_BITS = GMP_LIMB_BITS;
- while (x >>= LOCAL_GMP_LIMB_BITS)
- {
- ++ r->_mp_size;
- MPZ_REALLOC (r, r->_mp_size)[r->_mp_size - 1] = x;
- }
- }
+ {
+ int LOCAL_GMP_LIMB_BITS = GMP_LIMB_BITS;
+ while (x >>= LOCAL_GMP_LIMB_BITS)
+ {
+ ++ r->_mp_size;
+ MPZ_REALLOC (r, r->_mp_size)[r->_mp_size - 1] = x;
+ }
+ }
}
else
r->_mp_size = 0;
@@ -1622,7 +1632,7 @@ mpz_get_ui (const mpz_t u)
mp_size_t n = GMP_ABS (u->_mp_size);
n = GMP_MIN (n, 1 + (mp_size_t) (GMP_ULONG_BITS - 1) / GMP_LIMB_BITS);
while (--n >= 0)
- r = (r << LOCAL_GMP_LIMB_BITS) + u->_mp_d[n];
+ r = (r << LOCAL_GMP_LIMB_BITS) + u->_mp_d[n];
return r;
}
@@ -1681,7 +1691,7 @@ static mpz_srcptr
mpz_roinit_normal_n (mpz_t x, mp_srcptr xp, mp_size_t xs)
{
x->_mp_alloc = 0;
- //x->_mp_d = (mp_ptr) xp;
+ x->_mp_d = (mp_ptr) xp;
x->_mp_size = xs;
return x;
}
@@ -1778,11 +1788,11 @@ mpz_get_d (const mpz_t u)
{
x = B*x;
if (m > 0) {
- l = u->_mp_d[un];
- m -= GMP_LIMB_BITS;
- if (m < 0)
- l &= GMP_LIMB_MAX << -m;
- x += l;
+ l = u->_mp_d[un];
+ m -= GMP_LIMB_BITS;
+ if (m < 0)
+ l &= GMP_LIMB_MAX << -m;
+ x += l;
}
}
@@ -1811,24 +1821,24 @@ mpz_cmpabs_d (const mpz_t x, double d)
/* Scale d so it can be compared with the top limb. */
for (i = 1; i < xn; i++)
- d *= Bi;
+ d *= Bi;
if (d >= B)
- return -1;
+ return -1;
/* Compare floor(d) to top limb, subtract and cancel when equal. */
for (i = xn; i-- > 0;)
- {
- mp_limb_t f, xl;
-
- f = (mp_limb_t) d;
- xl = x->_mp_d[i];
- if (xl > f)
- return 1;
- else if (xl < f)
- return -1;
- d = B * (d - f);
- }
+ {
+ mp_limb_t f, xl;
+
+ f = (mp_limb_t) d;
+ xl = x->_mp_d[i];
+ if (xl > f)
+ return 1;
+ else if (xl < f)
+ return -1;
+ d = B * (d - f);
+ }
}
return - (d > 0.0);
}
@@ -1839,16 +1849,16 @@ mpz_cmp_d (const mpz_t x, double d)
if (x->_mp_size < 0)
{
if (d >= 0.0)
- return -1;
+ return -1;
else
- return -mpz_cmpabs_d (x, d);
+ return -mpz_cmpabs_d (x, d);
}
else
{
if (d < 0.0)
- return 1;
+ return 1;
else
- return mpz_cmpabs_d (x, d);
+ return mpz_cmpabs_d (x, d);
}
}
@@ -1916,7 +1926,7 @@ int
mpz_cmpabs (const mpz_t u, const mpz_t v)
{
return mpn_cmp4 (u->_mp_d, GMP_ABS (u->_mp_size),
- v->_mp_d, GMP_ABS (v->_mp_size));
+ v->_mp_d, GMP_ABS (v->_mp_size));
}
void
@@ -1936,16 +1946,8 @@ mpz_neg (mpz_t r, const mpz_t u)
void
mpz_swap (mpz_t u, mpz_t v)
{
- //MP_SIZE_T_SWAP (u->_mp_alloc, v->_mp_alloc);
- //MPN_PTR_SWAP (u->_mp_d, u->_mp_size, v->_mp_d, v->_mp_size);
-
- mpz_t temp;
- mpz_init(temp);
-
- *temp = *u;
- *u = *v;
- *v = *temp;
-
+ MP_SIZE_T_SWAP (u->_mp_alloc, v->_mp_alloc);
+ MPN_PTR_SWAP (u->_mp_d, u->_mp_size, v->_mp_d, v->_mp_size);
}
@@ -2191,7 +2193,7 @@ enum mpz_div_round_mode { GMP_DIV_FLOOR, GMP_DIV_CEIL, GMP_DIV_TRUNC };
/* Allows q or r to be zero. Returns 1 iff remainder is non-zero. */
static int
mpz_div_qr (mpz_t q, mpz_t r,
- const mpz_t n, const mpz_t d, enum mpz_div_round_mode mode)
+ const mpz_t n, const mpz_t d, enum mpz_div_round_mode mode)
{
mp_size_t ns, ds, nn, dn, qs;
ns = n->_mp_size;
@@ -2203,9 +2205,9 @@ mpz_div_qr (mpz_t q, mpz_t r,
if (ns == 0)
{
if (q)
- q->_mp_size = 0;
+ q->_mp_size = 0;
if (r)
- r->_mp_size = 0;
+ r->_mp_size = 0;
return 0;
}
@@ -2217,29 +2219,29 @@ mpz_div_qr (mpz_t q, mpz_t r,
if (nn < dn)
{
if (mode == GMP_DIV_CEIL && qs >= 0)
- {
- /* q = 1, r = n - d */
- if (r)
- mpz_sub (r, n, d);
- if (q)
- mpz_set_ui (q, 1);
- }
+ {
+ /* q = 1, r = n - d */
+ if (r)
+ mpz_sub (r, n, d);
+ if (q)
+ mpz_set_ui (q, 1);
+ }
else if (mode == GMP_DIV_FLOOR && qs < 0)
- {
- /* q = -1, r = n + d */
- if (r)
- mpz_add (r, n, d);
- if (q)
- mpz_set_si (q, -1);
- }
+ {
+ /* q = -1, r = n + d */
+ if (r)
+ mpz_add (r, n, d);
+ if (q)
+ mpz_set_si (q, -1);
+ }
else
- {
- /* q = 0, r = d */
- if (r)
- mpz_set (r, n);
- if (q)
- q->_mp_size = 0;
- }
+ {
+ /* q = 0, r = d */
+ if (r)
+ mpz_set (r, n);
+ if (q)
+ q->_mp_size = 0;
+ }
return 1;
}
else
@@ -2254,46 +2256,46 @@ mpz_div_qr (mpz_t q, mpz_t r,
qn = nn - dn + 1;
if (q)
- {
- mpz_init2 (tq, qn * GMP_LIMB_BITS);
- qp = tq->_mp_d;
- }
+ {
+ mpz_init2 (tq, qn * GMP_LIMB_BITS);
+ qp = tq->_mp_d;
+ }
else
- qp = NULL;
+ qp = NULL;
mpn_div_qr (qp, np, nn, d->_mp_d, dn);
if (qp)
- {
- qn -= (qp[qn-1] == 0);
+ {
+ qn -= (qp[qn-1] == 0);
- tq->_mp_size = qs < 0 ? -qn : qn;
- }
+ tq->_mp_size = qs < 0 ? -qn : qn;
+ }
rn = mpn_normalized_size (np, dn);
tr->_mp_size = ns < 0 ? - rn : rn;
if (mode == GMP_DIV_FLOOR && qs < 0 && rn != 0)
- {
- if (q)
- mpz_sub_ui (tq, tq, 1);
- if (r)
- mpz_add (tr, tr, d);
- }
+ {
+ if (q)
+ mpz_sub_ui (tq, tq, 1);
+ if (r)
+ mpz_add (tr, tr, d);
+ }
else if (mode == GMP_DIV_CEIL && qs >= 0 && rn != 0)
- {
- if (q)
- mpz_add_ui (tq, tq, 1);
- if (r)
- mpz_sub (tr, tr, d);
- }
+ {
+ if (q)
+ mpz_add_ui (tq, tq, 1);
+ if (r)
+ mpz_sub (tr, tr, d);
+ }
if (q)
- {
- mpz_swap (tq, q);
- mpz_clear (tq);
- }
+ {
+ mpz_swap (tq, q);
+ mpz_clear (tq);
+ }
if (r)
- mpz_swap (tr, r);
+ mpz_swap (tr, r);
mpz_clear (tr);
@@ -2363,7 +2365,7 @@ mpz_mod (mpz_t r, const mpz_t n, const mpz_t d)
static void
mpz_div_q_2exp (mpz_t q, const mpz_t u, mp_bitcnt_t bit_index,
- enum mpz_div_round_mode mode)
+ enum mpz_div_round_mode mode)
{
mp_size_t un, qn;
mp_size_t limb_cnt;
@@ -2384,9 +2386,9 @@ mpz_div_q_2exp (mpz_t q, const mpz_t u, mp_bitcnt_t bit_index,
/* Note: Below, the final indexing at limb_cnt is valid because at
that point we have qn > 0. */
adjust = (qn <= 0
- || !mpn_zero_p (u->_mp_d, limb_cnt)
- || (u->_mp_d[limb_cnt]
- & (((mp_limb_t) 1 << bit_index) - 1)));
+ || !mpn_zero_p (u->_mp_d, limb_cnt)
+ || (u->_mp_d[limb_cnt]
+ & (((mp_limb_t) 1 << bit_index) - 1)));
else
adjust = 0;
@@ -2397,14 +2399,14 @@ mpz_div_q_2exp (mpz_t q, const mpz_t u, mp_bitcnt_t bit_index,
qp = MPZ_REALLOC (q, qn);
if (bit_index != 0)
- {
- mpn_rshift (qp, u->_mp_d + limb_cnt, qn, bit_index);
- qn -= qp[qn - 1] == 0;
- }
+ {
+ mpn_rshift (qp, u->_mp_d + limb_cnt, qn, bit_index);
+ qn -= qp[qn - 1] == 0;
+ }
else
- {
- mpn_copyi (qp, u->_mp_d + limb_cnt, qn);
- }
+ {
+ mpn_copyi (qp, u->_mp_d + limb_cnt, qn);
+ }
}
q->_mp_size = qn;
@@ -2417,7 +2419,7 @@ mpz_div_q_2exp (mpz_t q, const mpz_t u, mp_bitcnt_t bit_index,
static void
mpz_div_r_2exp (mpz_t r, const mpz_t u, mp_bitcnt_t bit_index,
- enum mpz_div_round_mode mode)
+ enum mpz_div_round_mode mode)
{
mp_size_t us, un, rn;
mp_ptr rp;
@@ -2440,46 +2442,46 @@ mpz_div_r_2exp (mpz_t r, const mpz_t u, mp_bitcnt_t bit_index,
if (rn > un)
{
/* Quotient (with truncation) is zero, and remainder is
- non-zero */
+ non-zero */
if (mode == ((us > 0) ? GMP_DIV_CEIL : GMP_DIV_FLOOR)) /* us != 0 here. */
- {
- /* Have to negate and sign extend. */
- mp_size_t i;
+ {
+ /* Have to negate and sign extend. */
+ mp_size_t i;
- gmp_assert_nocarry (! mpn_neg (rp, u->_mp_d, un));
- for (i = un; i < rn - 1; i++)
- rp[i] = GMP_LIMB_MAX;
+ gmp_assert_nocarry (! mpn_neg (rp, u->_mp_d, un));
+ for (i = un; i < rn - 1; i++)
+ rp[i] = GMP_LIMB_MAX;
- rp[rn-1] = mask;
- us = -us;
- }
+ rp[rn-1] = mask;
+ us = -us;
+ }
else
- {
- /* Just copy */
- if (r != u)
- mpn_copyi (rp, u->_mp_d, un);
+ {
+ /* Just copy */
+ if (r != u)
+ mpn_copyi (rp, u->_mp_d, un);
- rn = un;
- }
+ rn = un;
+ }
}
else
{
if (r != u)
- mpn_copyi (rp, u->_mp_d, rn - 1);
+ mpn_copyi (rp, u->_mp_d, rn - 1);
rp[rn-1] = u->_mp_d[rn-1] & mask;
if (mode == ((us > 0) ? GMP_DIV_CEIL : GMP_DIV_FLOOR)) /* us != 0 here. */
- {
- /* If r != 0, compute 2^{bit_count} - r. */
- mpn_neg (rp, rp, rn);
+ {
+ /* If r != 0, compute 2^{bit_count} - r. */
+ mpn_neg (rp, rp, rn);
- rp[rn-1] &= mask;
+ rp[rn-1] &= mask;
- /* us is not used for anything else, so we can modify it
- here to indicate flipped sign. */
- us = -us;
- }
+ /* us is not used for anything else, so we can modify it
+ here to indicate flipped sign. */
+ us = -us;
+ }
}
rn = mpn_normalized_size (rp, rn);
r->_mp_size = us < 0 ? -rn : rn;
@@ -2553,7 +2555,7 @@ mpz_congruent_p (const mpz_t a, const mpz_t b, const mpz_t m)
static unsigned long
mpz_div_qr_ui (mpz_t q, mpz_t r,
- const mpz_t n, unsigned long d, enum mpz_div_round_mode mode)
+ const mpz_t n, unsigned long d, enum mpz_div_round_mode mode)
{
unsigned long ret;
mpz_t rr, dd;
@@ -2687,19 +2689,19 @@ mpn_gcd_11 (mp_limb_t u, mp_limb_t v)
while (u != v)
{
if (u > v)
- {
- u -= v;
- do
- u >>= 1;
- while ( (u & 1) == 0);
- }
+ {
+ u -= v;
+ do
+ u >>= 1;
+ while ( (u & 1) == 0);
+ }
else
- {
- v -= u;
- do
- v >>= 1;
- while ( (v & 1) == 0);
- }
+ {
+ v -= u;
+ do
+ v >>= 1;
+ while ( (v & 1) == 0);
+ }
}
return u << shift;
}
@@ -2771,30 +2773,30 @@ mpz_gcd (mpz_t g, const mpz_t u, const mpz_t v)
else
for (;;)
{
- int c;
-
- mpz_make_odd (tu);
- c = mpz_cmp (tu, tv);
- if (c == 0)
- {
- mpz_swap (g, tu);
- break;
- }
- if (c < 0)
- mpz_swap (tu, tv);
-
- if (tv->_mp_size == 1)
- {
- mp_limb_t *gp;
-
- mpz_tdiv_r (tu, tu, tv);
- gp = MPZ_REALLOC (g, 1); /* gp = mpz_limbs_modify (g, 1); */
- *gp = mpn_gcd_11 (tu->_mp_d[0], tv->_mp_d[0]);
-
- g->_mp_size = *gp != 0; /* mpz_limbs_finish (g, 1); */
- break;
- }
- mpz_sub (tu, tu, tv);
+ int c;
+
+ mpz_make_odd (tu);
+ c = mpz_cmp (tu, tv);
+ if (c == 0)
+ {
+ mpz_swap (g, tu);
+ break;
+ }
+ if (c < 0)
+ mpz_swap (tu, tv);
+
+ if (tv->_mp_size == 1)
+ {
+ mp_limb_t *gp;
+
+ mpz_tdiv_r (tu, tu, tv);
+ gp = MPZ_REALLOC (g, 1); /* gp = mpz_limbs_modify (g, 1); */
+ *gp = mpn_gcd_11 (tu->_mp_d[0], tv->_mp_d[0]);
+
+ g->_mp_size = *gp != 0; /* mpz_limbs_finish (g, 1); */
+ break;
+ }
+ mpz_sub (tu, tu, tv);
}
mpz_clear (tu);
mpz_clear (tv);
@@ -2814,9 +2816,9 @@ mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v)
signed long sign = mpz_sgn (v);
mpz_abs (g, v);
if (s)
- s->_mp_size = 0;
+ s->_mp_size = 0;
if (t)
- mpz_set_si (t, sign);
+ mpz_set_si (t, sign);
return;
}
@@ -2826,9 +2828,9 @@ mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v)
signed long sign = mpz_sgn (u);
mpz_abs (g, u);
if (s)
- mpz_set_si (s, sign);
+ mpz_set_si (s, sign);
if (t)
- t->_mp_size = 0;
+ t->_mp_size = 0;
return;
}
@@ -2894,39 +2896,39 @@ mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v)
power += shift;
for (;;)
- {
- int c;
- c = mpz_cmp (tu, tv);
- if (c == 0)
- break;
-
- if (c < 0)
- {
- /* tv = tv' + tu
- *
- * u = t0 tu + t1 (tv' + tu) = (t0 + t1) tu + t1 tv'
- * v = s0 tu + s1 (tv' + tu) = (s0 + s1) tu + s1 tv' */
-
- mpz_sub (tv, tv, tu);
- mpz_add (t0, t0, t1);
- mpz_add (s0, s0, s1);
-
- shift = mpz_make_odd (tv);
- mpz_mul_2exp (t1, t1, shift);
- mpz_mul_2exp (s1, s1, shift);
- }
- else
- {
- mpz_sub (tu, tu, tv);
- mpz_add (t1, t0, t1);
- mpz_add (s1, s0, s1);
-
- shift = mpz_make_odd (tu);
- mpz_mul_2exp (t0, t0, shift);
- mpz_mul_2exp (s0, s0, shift);
- }
- power += shift;
- }
+ {
+ int c;
+ c = mpz_cmp (tu, tv);
+ if (c == 0)
+ break;
+
+ if (c < 0)
+ {
+ /* tv = tv' + tu
+ *
+ * u = t0 tu + t1 (tv' + tu) = (t0 + t1) tu + t1 tv'
+ * v = s0 tu + s1 (tv' + tu) = (s0 + s1) tu + s1 tv' */
+
+ mpz_sub (tv, tv, tu);
+ mpz_add (t0, t0, t1);
+ mpz_add (s0, s0, s1);
+
+ shift = mpz_make_odd (tv);
+ mpz_mul_2exp (t1, t1, shift);
+ mpz_mul_2exp (s1, s1, shift);
+ }
+ else
+ {
+ mpz_sub (tu, tu, tv);
+ mpz_add (t1, t0, t1);
+ mpz_add (s1, s0, s1);
+
+ shift = mpz_make_odd (tu);
+ mpz_mul_2exp (t0, t0, shift);
+ mpz_mul_2exp (s0, s0, shift);
+ }
+ power += shift;
+ }
}
else
mpz_setbit (t0, uz);
@@ -2949,10 +2951,10 @@ mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v)
{
/* s0 u + t0 v = (s0 - v/g) u - (t0 + u/g) v */
if (mpz_odd_p (s0) || mpz_odd_p (t0))
- {
- mpz_sub (s0, s0, s1);
- mpz_add (t0, t0, t1);
- }
+ {
+ mpz_sub (s0, s0, s1);
+ mpz_add (t0, t0, t1);
+ }
assert (mpz_even_p (t0) && mpz_even_p (s0));
mpz_tdiv_q_2exp (s0, s0, 1);
mpz_tdiv_q_2exp (t0, t0, 1);
@@ -3038,12 +3040,12 @@ mpz_invert (mpz_t r, const mpz_t u, const mpz_t m)
if (invertible)
{
if (tr->_mp_size < 0)
- {
- if (m->_mp_size >= 0)
- mpz_add (tr, tr, m);
- else
- mpz_sub (tr, tr, m);
- }
+ {
+ if (m->_mp_size >= 0)
+ mpz_add (tr, tr, m);
+ else
+ mpz_sub (tr, tr, m);
+ }
mpz_swap (r, tr);
}
@@ -3067,7 +3069,7 @@ mpz_pow_ui (mpz_t r, const mpz_t b, unsigned long e)
{
mpz_mul (tr, tr, tr);
if (e & bit)
- mpz_mul (tr, tr, b);
+ mpz_mul (tr, tr, b);
bit >>= 1;
}
while (bit > 0);
@@ -3095,8 +3097,7 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
mp_srcptr mp;
struct gmp_div_inverse minv;
unsigned shift;
- //mp_ptr tp = NULL;
- mpz_t tp;
+ mp_ptr tp = NULL;
en = GMP_ABS (e->_mp_size);
mn = GMP_ABS (m->_mp_size);
@@ -3116,13 +3117,12 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
if (shift > 0)
{
/* To avoid shifts, we do all our reductions, except the final
- one, using a *normalized* m. */
+ one, using a *normalized* m. */
minv.shift = 0;
- //tp = gmp_alloc_limbs (mn);
-
- gmp_assert_nocarry (mpn_lshift (tp->_mp_d, mp, mn, shift));
- mp = tp->_mp_d;
+ tp = gmp_alloc_limbs (mn);
+ gmp_assert_nocarry (mpn_lshift (tp, mp, mn, shift));
+ mp = tp;
}
mpz_init (base);
@@ -3130,7 +3130,7 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
if (e->_mp_size < 0)
{
if (!mpz_invert (base, b, m))
- gmp_die ("mpz_powm: Negative exponent and non-invertible base.");
+ gmp_die ("mpz_powm: Negative exponent and non-invertible base.");
}
else
{
@@ -3139,20 +3139,20 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
bn = base->_mp_size;
if (bn >= mn)
- {
- mpn_div_qr_preinv (NULL, base->_mp_d, base->_mp_size, mp, mn, &minv);
- bn = mn;
- }
+ {
+ mpn_div_qr_preinv (NULL, base->_mp_d, base->_mp_size, mp, mn, &minv);
+ bn = mn;
+ }
/* We have reduced the absolute value. Now take care of the
- sign. Note that we get zero represented non-canonically as
- m. */
+ sign. Note that we get zero represented non-canonically as
+ m. */
if (b->_mp_size < 0)
- {
- mp_ptr bp = MPZ_REALLOC (base, mn);
- gmp_assert_nocarry (mpn_sub (bp, mp, mn, bp, bn));
- bn = mn;
- }
+ {
+ mp_ptr bp = MPZ_REALLOC (base, mn);
+ gmp_assert_nocarry (mpn_sub (bp, mp, mn, bp, bn));
+ bn = mn;
+ }
base->_mp_size = mpn_normalized_size (base->_mp_d, bn);
}
mpz_init_set_ui (tr, 1);
@@ -3164,17 +3164,17 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
bit = GMP_LIMB_HIGHBIT;
do
- {
- mpz_mul (tr, tr, tr);
- if (w & bit)
- mpz_mul (tr, tr, base);
- if (tr->_mp_size > mn)
- {
- mpn_div_qr_preinv (NULL, tr->_mp_d, tr->_mp_size, mp, mn, &minv);
- tr->_mp_size = mpn_normalized_size (tr->_mp_d, mn);
- }
- bit >>= 1;
- }
+ {
+ mpz_mul (tr, tr, tr);
+ if (w & bit)
+ mpz_mul (tr, tr, base);
+ if (tr->_mp_size > mn)
+ {
+ mpn_div_qr_preinv (NULL, tr->_mp_d, tr->_mp_size, mp, mn, &minv);
+ tr->_mp_size = mpn_normalized_size (tr->_mp_d, mn);
+ }
+ bit >>= 1;
+ }
while (bit > 0);
}
@@ -3185,8 +3185,8 @@ mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m)
mpn_div_qr_preinv (NULL, tr->_mp_d, tr->_mp_size, mp, mn, &minv);
tr->_mp_size = mpn_normalized_size (tr->_mp_d, mn);
}
- //if (tp)
- // gmp_free_limbs (tp, mn);
+ if (tp)
+ gmp_free_limbs (tp, mn);
mpz_swap (r, tr);
mpz_clear (tr);
@@ -3232,11 +3232,11 @@ mpz_rootrem (mpz_t x, mpz_t r, const mpz_t y, unsigned long z)
if (z == 2) /* simplify sqrt loop: z-1 == 1 */
do {
- mpz_swap (u, t); /* u = x */
- mpz_tdiv_q (t, y, u); /* t = y/x */
- mpz_add (t, t, u); /* t = y/x + x */
- mpz_tdiv_q_2exp (t, t, 1); /* x'= (y/x + x)/2 */
- } while (mpz_cmpabs (t, u) < 0); /* |x'| < |x| */
+ mpz_swap (u, t); /* u = x */
+ mpz_tdiv_q (t, y, u); /* t = y/x */
+ mpz_add (t, t, u); /* t = y/x + x */
+ mpz_tdiv_q_2exp (t, t, 1); /* x'= (y/x + x)/2 */
+ } while (mpz_cmpabs (t, u) < 0); /* |x'| < |x| */
else /* z != 2 */ {
mpz_t v;
@@ -3245,13 +3245,13 @@ mpz_rootrem (mpz_t x, mpz_t r, const mpz_t y, unsigned long z)
mpz_neg (t, t);
do {
- mpz_swap (u, t); /* u = x */
- mpz_pow_ui (t, u, z - 1); /* t = x^(z-1) */
- mpz_tdiv_q (t, y, t); /* t = y/x^(z-1) */
- mpz_mul_ui (v, u, z - 1); /* v = x*(z-1) */
- mpz_add (t, t, v); /* t = y/x^(z-1) + x*(z-1) */
- mpz_tdiv_q_ui (t, t, z); /* x'=(y/x^(z-1) + x*(z-1))/z */
- } while (mpz_cmpabs (t, u) < 0); /* |x'| < |x| */
+ mpz_swap (u, t); /* u = x */
+ mpz_pow_ui (t, u, z - 1); /* t = x^(z-1) */
+ mpz_tdiv_q (t, y, t); /* t = y/x^(z-1) */
+ mpz_mul_ui (v, u, z - 1); /* v = x*(z-1) */
+ mpz_add (t, t, v); /* t = y/x^(z-1) + x*(z-1) */
+ mpz_tdiv_q_ui (t, t, z); /* x'=(y/x^(z-1) + x*(z-1))/z */
+ } while (mpz_cmpabs (t, u) < 0); /* |x'| < |x| */
mpz_clear (v);
}
@@ -3405,18 +3405,18 @@ gmp_jacobi_coprime (mp_limb_t a, mp_limb_t b)
/* (2/b) = -1 if b = 3 or 5 mod 8 */
bit ^= c & (b ^ (b >> 1));
if (a < b)
- {
- if (a == 0)
- return bit & 1 ? -1 : 1;
- bit ^= a & b;
- a = b - a;
- b -= a;
- }
+ {
+ if (a == 0)
+ return bit & 1 ? -1 : 1;
+ bit ^= a & b;
+ a = b - a;
+ b -= a;
+ }
else
- {
- a -= b;
- assert (a != 0);
- }
+ {
+ a -= b;
+ assert (a != 0);
+ }
gmp_ctz(c, a);
++c;
@@ -3441,7 +3441,7 @@ gmp_lucas_step_k_2k (mpz_t V, mpz_t Qk, const mpz_t n)
/* Returns (U_k == 0) and sets V=V_k and Qk=Q^k. */
static int
gmp_lucas_mod (mpz_t V, mpz_t Qk, long Q,
- mp_bitcnt_t b0, const mpz_t n)
+ mp_bitcnt_t b0, const mpz_t n)
{
mp_bitcnt_t bs;
mpz_t U;
@@ -3465,27 +3465,27 @@ gmp_lucas_mod (mpz_t V, mpz_t Qk, long Q,
/* Q^{2k} = (Q^k)^2 */
gmp_lucas_step_k_2k (V, Qk, n);
- /* A step k->k+1 is performed if the bit in $n$ is 1 */
- /* mpz_tstbit(n,bs) or the bit is 0 in $n$ but */
- /* should be 1 in $n+1$ (bs == b0) */
+ /* A step k->k+1 is performed if the bit in $n$ is 1 */
+ /* mpz_tstbit(n,bs) or the bit is 0 in $n$ but */
+ /* should be 1 in $n+1$ (bs == b0) */
if (b0 == bs || mpz_tstbit (n, bs))
- {
- /* Q^{k+1} <- Q^k * Q */
- mpz_mul_si (Qk, Qk, Q);
- /* U_{k+1} <- (U_k + V_k) / 2 */
- mpz_swap (U, V); /* Keep in V the old value of U_k */
- mpz_add (U, U, V);
- /* We have to compute U/2, so we need an even value, */
- /* equivalent (mod n) */
- if (mpz_odd_p (U))
- mpz_add (U, U, n);
- mpz_tdiv_q_2exp (U, U, 1);
- /* V_{k+1} <-(D*U_k + V_k) / 2 =
- U_{k+1} + (D-1)/2*U_k = U_{k+1} - 2Q*U_k */
- mpz_mul_si (V, V, -2*Q);
- mpz_add (V, U, V);
- mpz_tdiv_r (V, V, n);
- }
+ {
+ /* Q^{k+1} <- Q^k * Q */
+ mpz_mul_si (Qk, Qk, Q);
+ /* U_{k+1} <- (U_k + V_k) / 2 */
+ mpz_swap (U, V); /* Keep in V the old value of U_k */
+ mpz_add (U, U, V);
+ /* We have to compute U/2, so we need an even value, */
+ /* equivalent (mod n) */
+ if (mpz_odd_p (U))
+ mpz_add (U, U, n);
+ mpz_tdiv_q_2exp (U, U, 1);
+ /* V_{k+1} <-(D*U_k + V_k) / 2 =
+ U_{k+1} + (D-1)/2*U_k = U_{k+1} - 2Q*U_k */
+ mpz_mul_si (V, V, -2*Q);
+ mpz_add (V, U, V);
+ mpz_tdiv_r (V, V, n);
+ }
mpz_tdiv_r (U, U, n);
}
@@ -3524,11 +3524,11 @@ gmp_stronglucas (const mpz_t x, mpz_t Qk)
do
{
if (D >= maxD)
- return 1 + (D != GMP_LIMB_MAX); /* (1 + ! ~ D) */
+ return 1 + (D != GMP_LIMB_MAX); /* (1 + ! ~ D) */
D += 2;
tl = mpz_tdiv_ui (n, D);
if (tl == 0)
- return 0;
+ return 0;
}
while (gmp_jacobi_coprime (tl, D) == 1);
@@ -3541,8 +3541,8 @@ gmp_stronglucas (const mpz_t x, mpz_t Qk)
/* D= P^2 - 4Q; P = 1; Q = (1-D)/4 */
Q = (D & 2) ? (long) (D >> 2) + 1 : -(long) (D >> 2);
- if (! gmp_lucas_mod (V, Qk, Q, b0, n)) /* If Ud != 0 */
- while (V->_mp_size != 0 && --b0 != 0) /* while Vk != 0 */
+ if (! gmp_lucas_mod (V, Qk, Q, b0, n)) /* If Ud != 0 */
+ while (V->_mp_size != 0 && --b0 != 0) /* while Vk != 0 */
/* V <- V ^ 2 - 2Q^k */
/* Q^{2k} = (Q^k)^2 */
gmp_lucas_step_k_2k (V, Qk, n);
@@ -3553,7 +3553,7 @@ gmp_stronglucas (const mpz_t x, mpz_t Qk)
static int
gmp_millerrabin (const mpz_t n, const mpz_t nm1, mpz_t y,
- const mpz_t q, mp_bitcnt_t k)
+ const mpz_t q, mp_bitcnt_t k)
{
assert (k > 0);
@@ -3567,7 +3567,7 @@ gmp_millerrabin (const mpz_t n, const mpz_t nm1, mpz_t y,
{
mpz_powm_ui (y, y, 2, n);
if (mpz_cmp (y, nm1) == 0)
- return 1;
+ return 1;
}
return 0;
}
@@ -3631,12 +3631,12 @@ mpz_probab_prime_p (const mpz_t n, int reps)
{
mpz_set_ui (y, (unsigned long) j*j+j+41);
if (mpz_cmp (y, nm1) >= 0)
- {
- /* Don't try any further bases. This "early" break does not affect
- the result for any reasonable reps value (<=5000 was tested) */
- assert (j >= 30);
- break;
- }
+ {
+ /* Don't try any further bases. This "early" break does not affect
+ the result for any reasonable reps value (<=5000 was tested) */
+ assert (j >= 30);
+ break;
+ }
is_prime = gmp_millerrabin (n, nm1, y, q, k);
}
mpz_clear (nm1);
@@ -3694,12 +3694,12 @@ mpz_tstbit (const mpz_t d, mp_bitcnt_t bit_index)
if (ds < 0)
{
/* d < 0. Check if any of the bits below is set: If so, our bit
- must be complemented. */
+ must be complemented. */
if (shift > 0 && (mp_limb_t) (w << (GMP_LIMB_BITS - shift)) > 0)
- return bit ^ 1;
+ return bit ^ 1;
while (--limb_index >= 0)
- if (d->_mp_d[limb_index] > 0)
- return bit ^ 1;
+ if (d->_mp_d[limb_index] > 0)
+ return bit ^ 1;
}
return bit;
}
@@ -3720,12 +3720,12 @@ mpz_abs_add_bit (mpz_t d, mp_bitcnt_t bit_index)
{
mp_size_t i;
/* The bit should be set outside of the end of the number.
- We have to increase the size of the number. */
+ We have to increase the size of the number. */
dp = MPZ_REALLOC (d, limb_index + 1);
dp[limb_index] = bit;
for (i = dn; i < limb_index; i++)
- dp[i] = 0;
+ dp[i] = 0;
dn = limb_index + 1;
}
else
@@ -3736,10 +3736,10 @@ mpz_abs_add_bit (mpz_t d, mp_bitcnt_t bit_index)
cy = mpn_add_1 (dp + limb_index, dp + limb_index, dn - limb_index, bit);
if (cy > 0)
- {
- dp = MPZ_REALLOC (d, dn + 1);
- dp[dn++] = cy;
- }
+ {
+ dp = MPZ_REALLOC (d, dn + 1);
+ dp[dn++] = cy;
+ }
}
d->_mp_size = (d->_mp_size < 0) ? - dn : dn;
@@ -3761,7 +3761,7 @@ mpz_abs_sub_bit (mpz_t d, mp_bitcnt_t bit_index)
assert (limb_index < dn);
gmp_assert_nocarry (mpn_sub_1 (dp + limb_index, dp + limb_index,
- dn - limb_index, bit));
+ dn - limb_index, bit));
dn = mpn_normalized_size (dp, dn);
d->_mp_size = (d->_mp_size < 0) ? - dn : dn;
}
@@ -3772,9 +3772,9 @@ mpz_setbit (mpz_t d, mp_bitcnt_t bit_index)
if (!mpz_tstbit (d, bit_index))
{
if (d->_mp_size >= 0)
- mpz_abs_add_bit (d, bit_index);
+ mpz_abs_add_bit (d, bit_index);
else
- mpz_abs_sub_bit (d, bit_index);
+ mpz_abs_sub_bit (d, bit_index);
}
}
@@ -3784,9 +3784,9 @@ mpz_clrbit (mpz_t d, mp_bitcnt_t bit_index)
if (mpz_tstbit (d, bit_index))
{
if (d->_mp_size >= 0)
- mpz_abs_sub_bit (d, bit_index);
+ mpz_abs_sub_bit (d, bit_index);
else
- mpz_abs_add_bit (d, bit_index);
+ mpz_abs_add_bit (d, bit_index);
}
}
@@ -3810,7 +3810,7 @@ void
mpz_and (mpz_t r, const mpz_t u, const mpz_t v)
{
mp_size_t un, vn, rn, i;
- mp_ptr up, vp, rp;
+ mp_ptr up, vp, rp;
mp_limb_t ux, vx, rx;
mp_limb_t uc, vc, rc;
@@ -3842,8 +3842,8 @@ mpz_and (mpz_t r, const mpz_t u, const mpz_t v)
rp = MPZ_REALLOC (r, rn + (mp_size_t) rc);
- up = *(mp_ptr *)u->_mp_d;
- vp = *(mp_ptr *)v->_mp_d;
+ up = u->_mp_d;
+ vp = v->_mp_d;
i = 0;
do
@@ -3915,8 +3915,8 @@ mpz_ior (mpz_t r, const mpz_t u, const mpz_t v)
rp = MPZ_REALLOC (r, rn + (mp_size_t) rc);
- up = *(mp_ptr *)u->_mp_d;
- vp = *(mp_ptr *)v->_mp_d;
+ up = u->_mp_d;
+ vp = v->_mp_d;
i = 0;
do
@@ -3984,8 +3984,8 @@ mpz_xor (mpz_t r, const mpz_t u, const mpz_t v)
rp = MPZ_REALLOC (r, un + (mp_size_t) rc);
- up = *(mp_ptr *)u->_mp_d;
- vp = *(mp_ptr *)v->_mp_d;
+ up = u->_mp_d;
+ vp = v->_mp_d;
i = 0;
do
@@ -4035,9 +4035,9 @@ gmp_popcount_limb (mp_limb_t x)
w = ((w >> 8) & 0x000f) + (w & 0x000f);
c += w;
if (GMP_LIMB_BITS > LOCAL_SHIFT_BITS)
- x >>= LOCAL_SHIFT_BITS;
+ x >>= LOCAL_SHIFT_BITS;
else
- x = 0;
+ x = 0;
}
return c;
}
@@ -4134,18 +4134,18 @@ mpz_scan1 (const mpz_t u, mp_bitcnt_t starting_bit)
if (i >= un)
return (us >= 0 ? ~(mp_bitcnt_t) 0 : starting_bit);
- up = *(mp_ptr *)u->_mp_d;
+ up = u->_mp_d;
ux = 0;
limb = up[i];
if (starting_bit != 0)
{
if (us < 0)
- {
- ux = mpn_zero_p (up, i);
- limb = ~ limb + ux;
- ux = - (mp_limb_t) (limb >= ux);
- }
+ {
+ ux = mpn_zero_p (up, i);
+ limb = ~ limb + ux;
+ ux = - (mp_limb_t) (limb >= ux);
+ }
/* Mask to 0 all bits before starting_bit, thus ignoring them. */
limb &= GMP_LIMB_MAX << (starting_bit % GMP_LIMB_BITS);
@@ -4171,7 +4171,7 @@ mpz_scan0 (const mpz_t u, mp_bitcnt_t starting_bit)
if (i >= un)
return (ux ? starting_bit : ~(mp_bitcnt_t) 0);
- up = *(mp_ptr *)u->_mp_d;
+ up = u->_mp_d;
limb = up[i] ^ ux;
if (ux == 0)
@@ -4191,8 +4191,7 @@ mpz_sizeinbase (const mpz_t u, int base)
{
mp_size_t un, tn;
mp_srcptr up;
- //mp_ptr tp;
- mpz_t tp;
+ mp_ptr tp;
mp_bitcnt_t bits;
struct gmp_div_inverse bi;
size_t ndigits;
@@ -4220,11 +4219,11 @@ mpz_sizeinbase (const mpz_t u, int base)
case 32:
return (bits + 4) / 5;
/* FIXME: Do something more clever for the common case of base
- 10. */
+ 10. */
}
- //tp = gmp_alloc_limbs (un);
- mpn_copyi (tp->_mp_d, up, un);
+ tp = gmp_alloc_limbs (un);
+ mpn_copyi (tp, up, un);
mpn_div_qr_1_invert (&bi, base);
tn = un;
@@ -4232,12 +4231,12 @@ mpz_sizeinbase (const mpz_t u, int base)
do
{
ndigits++;
- mpn_div_qr_1_preinv (tp->_mp_d, tp->_mp_d, tn, &bi);
- tn -= (tp->_mp_d[tn-1] == 0);
+ mpn_div_qr_1_preinv (tp, tp, tn, &bi);
+ tn -= (tp[tn-1] == 0);
}
while (tn > 0);
- // gmp_free_limbs (tp, un);
+ gmp_free_limbs (tp, un);
return ndigits;
}
@@ -4253,9 +4252,9 @@ mpz_get_str (char *sp, int base, const mpz_t u)
if (base > 1)
{
if (base <= 36)
- digits = "0123456789abcdefghijklmnopqrstuvwxyz";
+ digits = "0123456789abcdefghijklmnopqrstuvwxyz";
else if (base > 62)
- return NULL;
+ return NULL;
}
else if (base >= -1)
base = 10;
@@ -4263,7 +4262,7 @@ mpz_get_str (char *sp, int base, const mpz_t u)
{
base = -base;
if (base > 36)
- return NULL;
+ return NULL;
}
sn = 1 + mpz_sizeinbase (u, base);
@@ -4316,20 +4315,6 @@ ret:
return sp;
}
-
-int strlen_c(const char *c) {
-
- // rather naive implementation – we assume a string is terminated, and is not 0 characters long.
-
- int i = 0;
- while (1) {
- if (c[i] == '\0')
- return i;
- i++;
- }
- return i;
-}
-
int
mpz_set_str (mpz_t r, const char *sp, int base)
{
@@ -4351,22 +4336,22 @@ mpz_set_str (mpz_t r, const char *sp, int base)
if (base == 0)
{
if (sp[0] == '0')
- {
- if (sp[1] == 'x' || sp[1] == 'X')
- {
- base = 16;
- sp += 2;
- }
- else if (sp[1] == 'b' || sp[1] == 'B')
- {
- base = 2;
- sp += 2;
- }
- else
- base = 8;
- }
+ {
+ if (sp[1] == 'x' || sp[1] == 'X')
+ {
+ base = 16;
+ sp += 2;
+ }
+ else if (sp[1] == 'b' || sp[1] == 'B')
+ {
+ base = 2;
+ sp += 2;
+ }
+ else
+ base = 8;
+ }
else
- base = 10;
+ base = 10;
}
if (!*sp)
@@ -4383,22 +4368,22 @@ mpz_set_str (mpz_t r, const char *sp, int base)
unsigned digit;
if (isspace ((unsigned char) *sp))
- continue;
+ continue;
else if (*sp >= '0' && *sp <= '9')
- digit = *sp - '0';
+ digit = *sp - '0';
else if (*sp >= 'a' && *sp <= 'z')
- digit = *sp - 'a' + value_of_a;
+ digit = *sp - 'a' + value_of_a;
else if (*sp >= 'A' && *sp <= 'Z')
- digit = *sp - 'A' + 10;
+ digit = *sp - 'A' + 10;
else
- digit = base; /* fail */
+ digit = base; /* fail */
if (digit >= (unsigned) base)
- {
- gmp_free (dp, sn);
- r->_mp_size = 0;
- return -1;
- }
+ {
+ gmp_free (dp, sn);
+ r->_mp_size = 0;
+ return -1;
+ }
dp[dn++] = digit;
}
@@ -4470,7 +4455,7 @@ gmp_detect_endian (void)
/* Import and export. Does not support nails. */
void
mpz_import (mpz_t r, size_t count, int order, size_t size, int endian,
- size_t nails, const void *src)
+ size_t nails, const void *src)
{
const unsigned char *p;
ptrdiff_t word_step;
@@ -4517,15 +4502,15 @@ mpz_import (mpz_t r, size_t count, int order, size_t size, int endian,
{
size_t j;
for (j = 0; j < size; j++, p -= (ptrdiff_t) endian)
- {
- limb |= (mp_limb_t) *p << (bytes++ * CHAR_BIT);
- if (bytes == sizeof(mp_limb_t))
- {
- rp[i++] = limb;
- bytes = 0;
- limb = 0;
- }
- }
+ {
+ limb |= (mp_limb_t) *p << (bytes++ * CHAR_BIT);
+ if (bytes == sizeof(mp_limb_t))
+ {
+ rp[i++] = limb;
+ bytes = 0;
+ limb = 0;
+ }
+ }
}
assert (i + (bytes > 0) == rn);
if (limb != 0)
@@ -4538,7 +4523,7 @@ mpz_import (mpz_t r, size_t count, int order, size_t size, int endian,
void *
mpz_export (void *r, size_t *countp, int order, size_t size, int endian,
- size_t nails, const mpz_t u)
+ size_t nails, const mpz_t u)
{
size_t count;
mp_size_t un;
@@ -4572,65 +4557,65 @@ mpz_export (void *r, size_t *countp, int order, size_t size, int endian,
k = (GMP_LIMB_BITS <= CHAR_BIT);
if (!k)
- {
- do {
- int LOCAL_CHAR_BIT = CHAR_BIT;
- k++; limb >>= LOCAL_CHAR_BIT;
- } while (limb != 0);
- }
+ {
+ do {
+ int LOCAL_CHAR_BIT = CHAR_BIT;
+ k++; limb >>= LOCAL_CHAR_BIT;
+ } while (limb != 0);
+ }
/* else limb = 0; */
count = (k + (un-1) * sizeof (mp_limb_t) + size - 1) / size;
if (!r)
- r = gmp_alloc (count * size);
+ r = gmp_alloc (count * size);
if (endian == 0)
- endian = gmp_detect_endian ();
+ endian = gmp_detect_endian ();
p = (unsigned char *) r;
word_step = (order != endian) ? 2 * size : 0;
/* Process bytes from the least significant end, so point p at the
- least significant word. */
+ least significant word. */
if (order == 1)
- {
- p += size * (count - 1);
- word_step = - word_step;
- }
+ {
+ p += size * (count - 1);
+ word_step = - word_step;
+ }
/* And at least significant byte of that word. */
if (endian == 1)
- p += (size - 1);
+ p += (size - 1);
for (bytes = 0, i = 0, k = 0; k < count; k++, p += word_step)
- {
- size_t j;
- for (j = 0; j < size; ++j, p -= (ptrdiff_t) endian)
- {
- if (sizeof (mp_limb_t) == 1)
- {
- if (i < un)
- *p = u->_mp_d[i++];
- else
- *p = 0;
- }
- else
- {
- int LOCAL_CHAR_BIT = CHAR_BIT;
- if (bytes == 0)
- {
- if (i < un)
- limb = u->_mp_d[i++];
- bytes = sizeof (mp_limb_t);
- }
- *p = limb;
- limb >>= LOCAL_CHAR_BIT;
- bytes--;
- }
- }
- }
+ {
+ size_t j;
+ for (j = 0; j < size; ++j, p -= (ptrdiff_t) endian)
+ {
+ if (sizeof (mp_limb_t) == 1)
+ {
+ if (i < un)
+ *p = u->_mp_d[i++];
+ else
+ *p = 0;
+ }
+ else
+ {
+ int LOCAL_CHAR_BIT = CHAR_BIT;
+ if (bytes == 0)
+ {
+ if (i < un)
+ limb = u->_mp_d[i++];
+ bytes = sizeof (mp_limb_t);
+ }
+ *p = limb;
+ limb >>= LOCAL_CHAR_BIT;
+ bytes--;
+ }
+ }
+ }
assert (i == un);
assert (k == count);
}
@@ -4640,4 +4625,3 @@ mpz_export (void *r, size_t *countp, int order, size_t size, int endian,
return r;
}
-
diff --git a/source/gmp.h b/source/gmp.h
@@ -0,0 +1,310 @@
+/* mini-gmp, a minimalistic implementation of a GNU GMP subset.
+
+Copyright 2011-2015, 2017, 2019-2021 Free Software Foundation, Inc.
+
+This file is part of the GNU MP Library.
+
+The GNU MP Library is free software; you can redistribute it and/or modify
+it under the terms of either:
+
+ * the GNU Lesser General Public License as published by the Free
+ Software Foundation; either version 3 of the License, or (at your
+ option) any later version.
+
+or
+
+ * the GNU General Public License as published by the Free Software
+ Foundation; either version 2 of the License, or (at your option) any
+ later version.
+
+or both in parallel, as here.
+
+The GNU MP Library is distributed in the hope that it will be useful, but
+WITHOUT ANY WARRANTY; without even the implied warranty of MERCHANTABILITY
+or FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
+for more details.
+
+You should have received copies of the GNU General Public License and the
+GNU Lesser General Public License along with the GNU MP Library. If not,
+see https://www.gnu.org/licenses/. */
+
+/* About mini-gmp: This is a minimal implementation of a subset of the
+ GMP interface. It is intended for inclusion into applications which
+ have modest bignums needs, as a fallback when the real GMP library
+ is not installed.
+
+ This file defines the public interface. */
+
+#ifndef __MINI_GMP_H__
+#define __MINI_GMP_H__
+
+/* For size_t */
+#include <stddef.h>
+
+#if defined (__cplusplus)
+extern "C" {
+#endif
+
+void mp_set_memory_functions (void *(*) (size_t),
+ void *(*) (void *, size_t, size_t),
+ void (*) (void *, size_t));
+
+void mp_get_memory_functions (void *(**) (size_t),
+ void *(**) (void *, size_t, size_t),
+ void (**) (void *, size_t));
+
+#ifndef MINI_GMP_LIMB_TYPE
+#define MINI_GMP_LIMB_TYPE long
+#endif
+
+typedef unsigned MINI_GMP_LIMB_TYPE mp_limb_t;
+typedef long mp_size_t;
+typedef unsigned long mp_bitcnt_t;
+
+typedef mp_limb_t *mp_ptr;
+typedef const mp_limb_t *mp_srcptr;
+
+typedef struct
+{
+ int _mp_alloc; /* Number of *limbs* allocated and pointed
+ to by the _mp_d field. */
+ int _mp_size; /* abs(_mp_size) is the number of limbs the
+ last field points to. If _mp_size is
+ negative this is a negative number. */
+ mp_limb_t *_mp_d; /* Pointer to the limbs. */
+} __mpz_struct;
+
+typedef __mpz_struct mpz_t[1];
+
+typedef __mpz_struct *mpz_ptr;
+typedef const __mpz_struct *mpz_srcptr;
+
+extern const int mp_bits_per_limb;
+
+void mpn_copyi (mp_ptr, mp_srcptr, mp_size_t);
+void mpn_copyd (mp_ptr, mp_srcptr, mp_size_t);
+void mpn_zero (mp_ptr, mp_size_t);
+
+int mpn_cmp (mp_srcptr, mp_srcptr, mp_size_t);
+int mpn_zero_p (mp_srcptr, mp_size_t);
+
+mp_limb_t mpn_add_1 (mp_ptr, mp_srcptr, mp_size_t, mp_limb_t);
+mp_limb_t mpn_add_n (mp_ptr, mp_srcptr, mp_srcptr, mp_size_t);
+mp_limb_t mpn_add (mp_ptr, mp_srcptr, mp_size_t, mp_srcptr, mp_size_t);
+
+mp_limb_t mpn_sub_1 (mp_ptr, mp_srcptr, mp_size_t, mp_limb_t);
+mp_limb_t mpn_sub_n (mp_ptr, mp_srcptr, mp_srcptr, mp_size_t);
+mp_limb_t mpn_sub (mp_ptr, mp_srcptr, mp_size_t, mp_srcptr, mp_size_t);
+
+mp_limb_t mpn_mul_1 (mp_ptr, mp_srcptr, mp_size_t, mp_limb_t);
+mp_limb_t mpn_addmul_1 (mp_ptr, mp_srcptr, mp_size_t, mp_limb_t);
+mp_limb_t mpn_submul_1 (mp_ptr, mp_srcptr, mp_size_t, mp_limb_t);
+
+mp_limb_t mpn_mul (mp_ptr, mp_srcptr, mp_size_t, mp_srcptr, mp_size_t);
+void mpn_mul_n (mp_ptr, mp_srcptr, mp_srcptr, mp_size_t);
+void mpn_sqr (mp_ptr, mp_srcptr, mp_size_t);
+int mpn_perfect_square_p (mp_srcptr, mp_size_t);
+mp_size_t mpn_sqrtrem (mp_ptr, mp_ptr, mp_srcptr, mp_size_t);
+
+mp_limb_t mpn_lshift (mp_ptr, mp_srcptr, mp_size_t, unsigned int);
+mp_limb_t mpn_rshift (mp_ptr, mp_srcptr, mp_size_t, unsigned int);
+
+mp_bitcnt_t mpn_scan0 (mp_srcptr, mp_bitcnt_t);
+mp_bitcnt_t mpn_scan1 (mp_srcptr, mp_bitcnt_t);
+
+void mpn_com (mp_ptr, mp_srcptr, mp_size_t);
+mp_limb_t mpn_neg (mp_ptr, mp_srcptr, mp_size_t);
+
+mp_bitcnt_t mpn_popcount (mp_srcptr, mp_size_t);
+
+mp_limb_t mpn_invert_3by2 (mp_limb_t, mp_limb_t);
+#define mpn_invert_limb(x) mpn_invert_3by2 ((x), 0)
+
+size_t mpn_get_str (unsigned char *, int, mp_ptr, mp_size_t);
+mp_size_t mpn_set_str (mp_ptr, const unsigned char *, size_t, int);
+
+void mpz_init (mpz_t);
+void mpz_init2 (mpz_t, mp_bitcnt_t);
+void mpz_clear (mpz_t);
+
+#define mpz_odd_p(z) (((z)->_mp_size != 0) & (int) (z)->_mp_d[0])
+#define mpz_even_p(z) (! mpz_odd_p (z))
+
+int mpz_sgn (const mpz_t);
+int mpz_cmp_si (const mpz_t, long);
+int mpz_cmp_ui (const mpz_t, unsigned long);
+int mpz_cmp (const mpz_t, const mpz_t);
+int mpz_cmpabs_ui (const mpz_t, unsigned long);
+int mpz_cmpabs (const mpz_t, const mpz_t);
+int mpz_cmp_d (const mpz_t, double);
+int mpz_cmpabs_d (const mpz_t, double);
+
+void mpz_abs (mpz_t, const mpz_t);
+void mpz_neg (mpz_t, const mpz_t);
+void mpz_swap (mpz_t, mpz_t);
+
+void mpz_add_ui (mpz_t, const mpz_t, unsigned long);
+void mpz_add (mpz_t, const mpz_t, const mpz_t);
+void mpz_sub_ui (mpz_t, const mpz_t, unsigned long);
+void mpz_ui_sub (mpz_t, unsigned long, const mpz_t);
+void mpz_sub (mpz_t, const mpz_t, const mpz_t);
+
+void mpz_mul_si (mpz_t, const mpz_t, long int);
+void mpz_mul_ui (mpz_t, const mpz_t, unsigned long int);
+void mpz_mul (mpz_t, const mpz_t, const mpz_t);
+void mpz_mul_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_addmul_ui (mpz_t, const mpz_t, unsigned long int);
+void mpz_addmul (mpz_t, const mpz_t, const mpz_t);
+void mpz_submul_ui (mpz_t, const mpz_t, unsigned long int);
+void mpz_submul (mpz_t, const mpz_t, const mpz_t);
+
+void mpz_cdiv_qr (mpz_t, mpz_t, const mpz_t, const mpz_t);
+void mpz_fdiv_qr (mpz_t, mpz_t, const mpz_t, const mpz_t);
+void mpz_tdiv_qr (mpz_t, mpz_t, const mpz_t, const mpz_t);
+void mpz_cdiv_q (mpz_t, const mpz_t, const mpz_t);
+void mpz_fdiv_q (mpz_t, const mpz_t, const mpz_t);
+void mpz_tdiv_q (mpz_t, const mpz_t, const mpz_t);
+void mpz_cdiv_r (mpz_t, const mpz_t, const mpz_t);
+void mpz_fdiv_r (mpz_t, const mpz_t, const mpz_t);
+void mpz_tdiv_r (mpz_t, const mpz_t, const mpz_t);
+
+void mpz_cdiv_q_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_fdiv_q_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_tdiv_q_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_cdiv_r_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_fdiv_r_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+void mpz_tdiv_r_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+
+void mpz_mod (mpz_t, const mpz_t, const mpz_t);
+
+void mpz_divexact (mpz_t, const mpz_t, const mpz_t);
+
+int mpz_divisible_p (const mpz_t, const mpz_t);
+int mpz_congruent_p (const mpz_t, const mpz_t, const mpz_t);
+
+unsigned long mpz_cdiv_qr_ui (mpz_t, mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_fdiv_qr_ui (mpz_t, mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_tdiv_qr_ui (mpz_t, mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_cdiv_q_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_fdiv_q_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_tdiv_q_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_cdiv_r_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_fdiv_r_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_tdiv_r_ui (mpz_t, const mpz_t, unsigned long);
+unsigned long mpz_cdiv_ui (const mpz_t, unsigned long);
+unsigned long mpz_fdiv_ui (const mpz_t, unsigned long);
+unsigned long mpz_tdiv_ui (const mpz_t, unsigned long);
+
+unsigned long mpz_mod_ui (mpz_t, const mpz_t, unsigned long);
+
+void mpz_divexact_ui (mpz_t, const mpz_t, unsigned long);
+
+int mpz_divisible_ui_p (const mpz_t, unsigned long);
+
+unsigned long mpz_gcd_ui (mpz_t, const mpz_t, unsigned long);
+void mpz_gcd (mpz_t, const mpz_t, const mpz_t);
+void mpz_gcdext (mpz_t, mpz_t, mpz_t, const mpz_t, const mpz_t);
+void mpz_lcm_ui (mpz_t, const mpz_t, unsigned long);
+void mpz_lcm (mpz_t, const mpz_t, const mpz_t);
+int mpz_invert (mpz_t, const mpz_t, const mpz_t);
+
+void mpz_sqrtrem (mpz_t, mpz_t, const mpz_t);
+void mpz_sqrt (mpz_t, const mpz_t);
+int mpz_perfect_square_p (const mpz_t);
+
+void mpz_pow_ui (mpz_t, const mpz_t, unsigned long);
+void mpz_ui_pow_ui (mpz_t, unsigned long, unsigned long);
+void mpz_powm (mpz_t, const mpz_t, const mpz_t, const mpz_t);
+void mpz_powm_ui (mpz_t, const mpz_t, unsigned long, const mpz_t);
+
+void mpz_rootrem (mpz_t, mpz_t, const mpz_t, unsigned long);
+int mpz_root (mpz_t, const mpz_t, unsigned long);
+
+void mpz_fac_ui (mpz_t, unsigned long);
+void mpz_2fac_ui (mpz_t, unsigned long);
+void mpz_mfac_uiui (mpz_t, unsigned long, unsigned long);
+void mpz_bin_uiui (mpz_t, unsigned long, unsigned long);
+
+int mpz_probab_prime_p (const mpz_t, int);
+
+int mpz_tstbit (const mpz_t, mp_bitcnt_t);
+void mpz_setbit (mpz_t, mp_bitcnt_t);
+void mpz_clrbit (mpz_t, mp_bitcnt_t);
+void mpz_combit (mpz_t, mp_bitcnt_t);
+
+void mpz_com (mpz_t, const mpz_t);
+void mpz_and (mpz_t, const mpz_t, const mpz_t);
+void mpz_ior (mpz_t, const mpz_t, const mpz_t);
+void mpz_xor (mpz_t, const mpz_t, const mpz_t);
+
+mp_bitcnt_t mpz_popcount (const mpz_t);
+mp_bitcnt_t mpz_hamdist (const mpz_t, const mpz_t);
+mp_bitcnt_t mpz_scan0 (const mpz_t, mp_bitcnt_t);
+mp_bitcnt_t mpz_scan1 (const mpz_t, mp_bitcnt_t);
+
+int mpz_fits_slong_p (const mpz_t);
+int mpz_fits_ulong_p (const mpz_t);
+int mpz_fits_sint_p (const mpz_t);
+int mpz_fits_uint_p (const mpz_t);
+int mpz_fits_sshort_p (const mpz_t);
+int mpz_fits_ushort_p (const mpz_t);
+long int mpz_get_si (const mpz_t);
+unsigned long int mpz_get_ui (const mpz_t);
+double mpz_get_d (const mpz_t);
+size_t mpz_size (const mpz_t);
+mp_limb_t mpz_getlimbn (const mpz_t, mp_size_t);
+
+void mpz_realloc2 (mpz_t, mp_bitcnt_t);
+mp_srcptr mpz_limbs_read (mpz_srcptr);
+mp_ptr mpz_limbs_modify (mpz_t, mp_size_t);
+mp_ptr mpz_limbs_write (mpz_t, mp_size_t);
+void mpz_limbs_finish (mpz_t, mp_size_t);
+mpz_srcptr mpz_roinit_n (mpz_t, mp_srcptr, mp_size_t);
+
+#define MPZ_ROINIT_N(xp, xs) {{0, (xs),(xp) }}
+
+void mpz_set_si (mpz_t, signed long int);
+void mpz_set_ui (mpz_t, unsigned long int);
+void mpz_set (mpz_t, const mpz_t);
+void mpz_set_d (mpz_t, double);
+
+void mpz_init_set_si (mpz_t, signed long int);
+void mpz_init_set_ui (mpz_t, unsigned long int);
+void mpz_init_set (mpz_t, const mpz_t);
+void mpz_init_set_d (mpz_t, double);
+
+size_t mpz_sizeinbase (const mpz_t, int);
+char *mpz_get_str (char *, int, const mpz_t);
+int mpz_set_str (mpz_t, const char *, int);
+int mpz_init_set_str (mpz_t, const char *, int);
+
+/* This long list taken from gmp.h. */
+/* For reference, "defined(EOF)" cannot be used here. In g++ 2.95.4,
+ <iostream> defines EOF but not FILE. */
+#if defined (FILE) \
+ || defined (H_STDIO) \
+ || defined (_H_STDIO) /* AIX */ \
+ || defined (_STDIO_H) /* glibc, Sun, SCO */ \
+ || defined (_STDIO_H_) /* BSD, OSF */ \
+ || defined (__STDIO_H) /* Borland */ \
+ || defined (__STDIO_H__) /* IRIX */ \
+ || defined (_STDIO_INCLUDED) /* HPUX */ \
+ || defined (__dj_include_stdio_h_) /* DJGPP */ \
+ || defined (_FILE_DEFINED) /* Microsoft */ \
+ || defined (__STDIO__) /* Apple MPW MrC */ \
+ || defined (_MSL_STDIO_H) /* Metrowerks */ \
+ || defined (_STDIO_H_INCLUDED) /* QNX4 */ \
+ || defined (_ISO_STDIO_ISO_H) /* Sun C++ */ \
+ || defined (__STDIO_LOADED) /* VMS */ \
+ || defined (_STDIO) /* HPE NonStop */ \
+ || defined (__DEFINED_FILE) /* musl */
+size_t mpz_out_str (FILE *, int, const mpz_t);
+#endif
+
+void mpz_import (mpz_t, size_t, int, size_t, int, size_t, const void *);
+void *mpz_export (void *, size_t *, int, size_t, int, size_t, const mpz_t);
+
+#if defined (__cplusplus)
+}
+#endif
+#endif /* __MINI_GMP_H__ */
diff --git a/source/lib-gpu-verify.c b/source/lib-gpu-verify.c
@@ -10,6 +10,7 @@
#include "reference-test.h"
#include "montgomery-test.h"
+#include "montmodmult.h"
int main(int argc, char** argv)
@@ -19,12 +20,12 @@ int main(int argc, char** argv)
setup_gcry();
-
+ mont_modmult_tests();
// MARK: might have memory issues
- mont_rsa_tests();
+ //mont_rsa_tests();
- //rsa_tests();
+ rsa_tests();
//reference_tests();
diff --git a/source/montgomery-test.c b/source/montgomery-test.c
@@ -7,7 +7,7 @@
#include "montgomery-test.h"
-#include "gmp.h" // has been adapted
+#include <gmp.h> // has been adapted
void mont_pairs_from_files(void *bases, unsigned long *b_off,
void *exponents, unsigned long *e_off,
diff --git a/source/montgomery.h b/source/montgomery.h
@@ -10,7 +10,7 @@
#include <stdio.h>
-#include "gmp.h"
+#include <gmp.h>
#include <assert.h>
diff --git a/source/montmodmult.c b/source/montmodmult.c
@@ -6,3 +6,619 @@
//
#include "montmodmult.h"
+#include "util.h"
+
+static unsigned long len_in_bytes = 0;
+
+#define ORDER -1 // I think we need to do this, because we want to write it in the 'wrong' way
+#define END 0
+
+#define BIT_LENGTH 2048
+
+#define BITS 64
+
+// sizes are always the same 32 units for all, except exp_buf
+void montmodmult_pairs_from_files(void *x_buf, void *m_buf,
+ void *r_1_buf,
+ void *n_buf, void *ni_buf,
+ void *msg_buf,
+ void *exp_buf,
+ void *mod_buf,
+ void *s_buf,
+ unsigned long *pks, unsigned long *n) {
+
+ FILE * pk;
+ FILE * ms;
+
+ pk = fopen("lib-gpu-generate/publickey.txt", "r");
+ ms = fopen("lib-gpu-generate/msgsig.txt", "r");
+
+ fseek (ms, 0, SEEK_END);
+ long ms_l = ftell(ms);
+ fseek (ms, 0, SEEK_SET);
+ char *ms_ptr = malloc(ms_l);
+ char *ms_ptr_rest = malloc(ms_l);
+ if (ms_ptr || ms_ptr_rest)
+ {
+ fread (ms_ptr, 1, ms_l, ms);
+ memcpy(ms_ptr_rest, ms_ptr, ms_l);
+ }
+ fclose (ms);
+
+ fseek (pk, 0, SEEK_END);
+ long pk_l = ftell(pk);
+ fseek (pk, 0, SEEK_SET);
+ char *pk_ptr = malloc(pk_l);
+ char *pk_ptr_rest = malloc(pk_l);
+ if (pk_ptr && pk_ptr_rest)
+ {
+ fread (pk_ptr, 1, pk_l, pk);
+ memcpy(pk_ptr_rest, pk_ptr, pk_l);
+ }
+ fclose (pk);
+
+ gpu_register *x_buf_t = x_buf;
+ gpu_register *m_buf_t = m_buf;
+ gpu_register *r_1_buf_t = r_1_buf;
+
+ gpu_register *ni_buf_t = ni_buf;
+ gpu_register *n_buf_t = n_buf;
+ gpu_register *msg_buf_t = msg_buf;
+ gpu_register *s_buf_t = s_buf;
+ gpu_register *mod_buf_t = mod_buf;
+ gpu_register *exp_buf_t = exp_buf;
+
+ int len = (BIT_LENGTH / 8) / sizeof(gpu_register);
+
+ char* message = strtok_r(ms_ptr, "\n", &ms_ptr_rest);
+ char* signature = strtok_r(0, "\n", &ms_ptr_rest);
+ char* modulus = strtok_r(pk_ptr, "\n", &pk_ptr_rest);
+ char* exponent = strtok_r(0, "\n", &pk_ptr_rest);
+ char* offs = strtok_r(0, "\n", &pk_ptr_rest);
+
+ int i = 0;
+ int j = 0;
+
+ mpz_t e,mod,msg,s;
+
+ mpz_init(e);
+ mpz_init(mod);
+ mpz_init(msg);
+ mpz_init(s);
+
+ mpz_t r, r_1, ni, M, x, N;
+ mpz_t NN;
+
+ mpz_init(r);
+ mpz_init(r_1);
+ mpz_init(ni);
+ mpz_init(M);
+ mpz_init(x);
+ mpz_init(N);
+ mpz_init(NN);
+
+ mpz_t one; // some helper variables
+ mpz_init_set_si(one,1);
+
+
+ // N is a constant for a given field (2048)
+ mpz_mul_2exp(N,one,2043); // r
+ mpz_mul_2exp(NN,one,44); // r
+
+ mpz_add(N, N, NN);
+ mpz_add_ui(N, N, 1); // N = 2^2044 + 2^45 + 1
+
+ while (message != NULL && signature != NULL) {
+
+ if (i == 0 || pks[j - 1] < i) {
+
+ mpz_set_str(mod,modulus,16);
+ mpz_set_str(e,exponent,16);
+
+ pks[j] = atoi(offs);
+
+ modulus = strtok_r(0, "\n", &pk_ptr_rest);
+ exponent = strtok_r(0, "\n", &pk_ptr_rest);
+ offs = strtok_r(0, "\n", &pk_ptr_rest);
+
+ // MARK: into montgomery domain
+
+
+ mpz_set_si(one, 1);
+
+
+
+ mpz_mul_2exp(r,one,BIT_LENGTH); // r
+
+ mpz_gcdext(one, r_1, ni, r, mod); // set r_1 and ni
+
+ // MARK: n' is negative – fix that
+
+ int sgn = mpz_sgn(r_1);
+
+ mpz_abs(r_1, r_1);
+ mpz_abs(ni, ni);
+
+ if (sgn == -1) {
+ mpz_sub(ni, r, ni);
+ mpz_sub(r_1, N, r_1);
+ }
+
+ // MARK: export
+
+ // exponent as is and mod -> n, ni -> ni, r_1 -> r_1
+
+ mpz_export(&n_buf_t[len * j], NULL, ORDER, sizeof(gpu_register), END, 0, mod);
+ mpz_export(&mod_buf_t[len * j], NULL, ORDER, sizeof(gpu_register), END, 0, mod);
+ mpz_export(&ni_buf_t[len * j], NULL, ORDER, sizeof(gpu_register), END, 0, ni);
+ mpz_export(&r_1_buf_t[len * j], NULL, ORDER, sizeof(gpu_register), END, 0, r_1);
+ mpz_export(&exp_buf_t[j], NULL, ORDER, sizeof(gpu_register), END, 0, e);
+
+ j++;
+
+ }
+
+ mpz_set_str(msg,message,16);
+ mpz_set_str(s,signature,16);
+
+ message = strtok_r(0, "\n",&ms_ptr_rest);
+ signature = strtok_r(0, "\n",&ms_ptr_rest);
+
+ // set x (the number to 'square' (multiply by itself))
+ mpz_mul(x, s, r);
+ mpz_mod(x, x, mod);
+
+ // message as is and 'signature' -> x , M -> m
+ mpz_export(&msg_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, msg);
+ mpz_export(&s_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, s);
+ mpz_export(&x_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, s);
+ mpz_export(&m_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, s);
+
+
+
+
+
+ i++;
+ }
+
+ mpz_clear(e);
+ mpz_clear(mod);
+ mpz_clear(msg);
+ mpz_clear(s);
+
+ mpz_clear(r);
+ mpz_clear(r_1);
+ mpz_clear(ni);
+ mpz_clear(M);
+ mpz_clear(x);
+
+ mpz_clear(N);
+ mpz_clear(NN);
+
+ mpz_clear(one);
+
+ *n = i;
+
+}
+
+
+int modmult_opencl_prepare(struct gpu_info *info, struct gpu_state_alt *state,
+ void *x_buf, void *m_buf,
+ //void *r_1_buf,
+ void *n_buf, void *ni_buf,
+ void *msg_buf,
+ void *exp_buf,
+ void *mod_buf,
+ void *s_buf,
+ unsigned long *pks, unsigned long n
+ ) {
+
+ int err; // error code returned from api calls
+
+ unsigned long pk = 0;
+
+ while (1) {
+ if (pks[pk] + 1 >= n)
+ break;
+ pk++;
+ }
+
+ unsigned long len = len_in_bytes;
+
+ state->x_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len, NULL, NULL);
+ state->m_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
+
+ //state->res_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len, NULL, NULL); // the result that is written back
+ state->n_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
+ state->ni_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
+
+ state->exp_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, n * sizeof(gpu_register), NULL, NULL);
+
+ state->tmp_1 = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len * 2, NULL, NULL);
+ state->tmp_2 = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len * 2, NULL, NULL);
+
+ state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(unsigned long) * (pk + 1),NULL, NULL);
+
+ if (!state->x_mem || !state->m_mem || !state->n_mem || !state->ni_mem || !state->exp_mem || !state->tmp_1 || !state->tmp_2)
+ {
+ printf("Error: Failed to allocate device memory!\n");
+ exit(1);
+ }
+
+ void *tmp_1 = malloc(len * 2);
+ void *tmp_2 = malloc(len * 2);
+
+ //void *res = malloc(len);
+
+ memset(tmp_1, 0, len * 2);
+ memset(tmp_2, 0, len * 2);
+ //memset(res, 0, len);
+
+ // Write our data set into the input array in device memory
+ //
+ err = clEnqueueWriteBuffer(info->commands, state->x_mem, CL_TRUE, 0, len, x_buf, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(info->commands, state->m_mem, CL_TRUE, 0, len, m_buf, 0, NULL, NULL);
+
+ //err |= clEnqueueWriteBuffer(info->commands, state->res_mem, CL_TRUE, 0, len, res, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(info->commands, state->n_mem, CL_TRUE, 0, len, n_buf, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(info->commands, state->ni_mem, CL_TRUE, 0, len, ni_buf, 0, NULL, NULL);
+
+ err |= clEnqueueWriteBuffer(info->commands, state->exp_mem, CL_TRUE, 0, n * sizeof(gpu_register), exp_buf, 0, NULL, NULL);
+
+ err |= clEnqueueWriteBuffer(info->commands, state->tmp_1, CL_TRUE, 0, len * 2, tmp_1, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(info->commands, state->tmp_2, CL_TRUE, 0, len * 2, tmp_2, 0, NULL, NULL);
+
+ err |= clEnqueueWriteBuffer(info->commands, state->pks_indices, CL_TRUE, 0, sizeof(unsigned long) * (pk + 1), pks, 0, NULL, NULL);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to write to source array!\n");
+ exit(1);
+ }
+
+ free(tmp_1);
+ free(tmp_2);
+ //free(res);
+
+ // Set the arguments to our compute kernel
+ //
+ err = 0;
+ err = clSetKernelArg(info->kernel, 0, sizeof(cl_mem), &state->x_mem);
+ err |= clSetKernelArg(info->kernel, 1, sizeof(cl_mem), &state->m_mem);
+ //err |= clSetKernelArg(info->kernel, 2, sizeof(cl_mem), &state->res_mem);
+ err |= clSetKernelArg(info->kernel, 3, sizeof(cl_mem), &state->n_mem);
+ err |= clSetKernelArg(info->kernel, 4, sizeof(cl_mem), &state->ni_mem);
+ err |= clSetKernelArg(info->kernel, 5, sizeof(cl_mem), &state->exp_mem);
+ err |= clSetKernelArg(info->kernel, 6, sizeof(cl_mem), &state->tmp_1);
+ err |= clSetKernelArg(info->kernel, 7, sizeof(cl_mem), &state->tmp_2);
+ err |= clSetKernelArg(info->kernel, 8, sizeof(cl_mem), &state->pks_indices);
+ //err |= clSetKernelArg(info->kernel, 9, sizeof(unsigned long), &n);
+
+
+ if (err != CL_SUCCESS)
+ {
+ printf("RSA-Error: Failed to set kernel arguments! %d\n", err);
+ exit(1);
+ }
+
+ state->total = n;
+
+
+ return 0;
+
+}
+
+int modmult_opencl_exec_kernel(struct gpu_info *info, struct gpu_state_alt *state) {
+
+ size_t global;
+ // size_t local;
+ int err;
+
+ // MARK: deal with size limitations
+
+ size_t modifier = (BIT_LENGTH / 8) / sizeof(gpu_register);
+
+ // MARK: testing with one kernel, and a loop that simulates work items
+
+ //global = state->total % len == 0? state->total : state->total + state->total % len;
+ global = state->total; // has to be exactly the amount of signatures we want to verify
+
+ // measure from the first call to the kernel...
+ if (state->skip) {
+ state->skip = false;
+ clock_gettime(CLOCK_REALTIME, &state->t1);
+ }
+
+ err = clEnqueueNDRangeKernel(info->commands, info->kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
+ if (err)
+ {
+ printf("Error: Failed to execute kernel!\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("KERNEL IS EXECUTING...\n");
+
+ return 0;
+
+}
+
+unsigned long modmult_opencl_results(struct gpu_info *info, struct gpu_state_alt *state, bool timed, void * msg_buf, void * r_1_buf, void * n_buf, void* s_buf, void *mod_buf, unsigned long n) {
+
+ if (state->skip) {
+ // reset skip in the kernel execution
+ return 0;
+ }
+
+ int err;
+
+ // Wait for the command commands to get serviced before reading back results
+ //
+ err = clFinish(info->commands);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Kernel failure! %d\n", err);
+ exit(1);
+ }
+
+
+
+ void *results = malloc(len_in_bytes);
+ memset(results, 0, len_in_bytes);
+
+
+ // Read back the results from the device to verify the output
+ err = clEnqueueReadBuffer(info->commands, state->res_mem, CL_TRUE, 0, len_in_bytes, results, 0, NULL, NULL );
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to read output array! %d\n", err);
+ exit(1);
+ }
+
+ // MARK: convert out of montgomery
+
+ mpz_t ret,a,r_1,mod,N;
+
+ mpz_init(ret);
+ mpz_init(a);
+ mpz_init(r_1);
+ mpz_init(mod);
+ mpz_init(N);
+
+ //mpz_export(&n_buf_t[len * j], NULL, 1, sizeof(gpu_register), 0, 0, mod);
+
+ int len = (BIT_LENGTH / 8) / sizeof(gpu_register);
+
+ for (int i = 0; i < n; i++) {
+
+ mpz_import(a, len, ORDER, sizeof(gpu_register), END, 0, &results[i * len]);
+ mpz_import(r_1, len, ORDER, sizeof(gpu_register), END, 0, &r_1_buf[i * len]);
+ mpz_import(N, len, ORDER, sizeof(gpu_register), END, 0, &n_buf[i * len]);
+ mpz_import(mod, len, ORDER, sizeof(gpu_register), END, 0, &mod_buf[i * len]);
+
+ // mpz_mul(ret, a, r_1);
+ // mpz_mod(ret, ret, mod);
+
+ }
+
+ gmp_printf ("ret in decimal: %Zd\n", a);
+
+ // test reference
+
+ mpz_t sig, sig_tmp, one, r;
+ mpz_init(sig);
+ mpz_init(sig_tmp);
+
+ mpz_init(one);
+ mpz_init(r);
+
+ mpz_import(sig, len, ORDER, sizeof(gpu_register), END, 0, s_buf);
+
+ mpz_set_si(one, 1);
+ mpz_mul_2exp(r,one,BIT_LENGTH); // r
+
+ mpz_mul(sig, sig, sig);
+ mpz_mod(sig, sig, r);
+
+ gmp_printf ("ref: %Zd\n", sig);
+
+
+ printf("");
+
+ mpz_clear(ret);
+ mpz_clear(a);
+ mpz_clear(r_1);
+ mpz_clear(mod);
+
+ mpz_clear(one);
+
+
+
+
+
+
+
+
+ if (timed) {
+ // stop measuring after the last command has been read
+ clock_gettime(CLOCK_REALTIME, &state->t2);
+
+ printf("GPU verification took %ld.%06ld s\n", ( state->t2.tv_nsec < state->t1.tv_nsec ? state->t2.tv_sec - (state->t1.tv_sec + 1) : state->t2.tv_sec - state->t1.tv_sec ), ( state->t2.tv_nsec < state->t1.tv_nsec ? ((999999999 - state->t1.tv_nsec) + state->t2.tv_nsec) : (state->t2.tv_nsec - state->t1.tv_nsec) ) / 1000);
+
+ }
+
+
+ return 0;
+
+}
+
+void modmult_opencl_cleanup(struct gpu_info *info) {
+
+ clReleaseProgram(info->program);
+ clReleaseKernel(info->kernel);
+ clReleaseCommandQueue(info->commands);
+ clReleaseContext(info->context);
+
+}
+
+void modmult_opencl_release(struct gpu_state_alt *state) {
+
+ clReleaseMemObject(state->x_mem);
+ clReleaseMemObject(state->m_mem);
+ //clReleaseMemObject(state->res_mem);
+ clReleaseMemObject(state->n_mem);
+ clReleaseMemObject(state->ni_mem);
+ clReleaseMemObject(state->exp_mem);
+
+ clReleaseMemObject(state->tmp_1);
+ clReleaseMemObject(state->tmp_2);
+
+ clReleaseMemObject(state->pks_indices);
+}
+
+// MARK: for library
+
+void modmult_gpu_init(struct gpu_info *info, struct gpu_state_alt *state) {
+
+ info->platform = select_platform(0, false);
+ info->device_id = select_device (info->platform);
+ info->context = create_compute_context (info->device_id);
+ info->commands = create_command_queue (info->device_id, info->context);
+ info->program = compile_program (info->device_id, info->context, "montmodmult.cl");
+ info->kernel = create_kernel (info->program, "mont");
+
+ state->result = 0;
+ state->total = 0;
+ state->skip = true;
+
+ int err = 0;
+
+ void *results = malloc(len_in_bytes);
+ memset(results, 0, len_in_bytes);
+
+ state->res_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len_in_bytes ,NULL, NULL);
+
+ err |= clEnqueueWriteBuffer(info->commands, state->res_mem, CL_TRUE, 0, len_in_bytes, results, 0, NULL, NULL);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to write to source array!\n");
+ exit(1);
+ }
+
+ err |= clSetKernelArg(info->kernel, 2, sizeof(cl_mem), &state->res_mem);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to set kernel arguments! %d\n", err);
+ exit(1);
+ }
+}
+
+void modmult_gpu_execute(struct gpu_info *info, struct gpu_state_alt *state,
+ void *x_buf, void *m_buf,
+ void *r_1_buf,
+ void *n_buf, void *ni_buf,
+ void *msg_buf,
+ void *exp_buf,
+ void *mod_buf,
+ void *s_buf,
+ unsigned long *pks, unsigned long n) {
+
+ modmult_opencl_prepare(info, state,
+ x_buf, m_buf,
+ //r_1_buf,
+ n_buf, ni_buf,
+ msg_buf, exp_buf, mod_buf, s_buf,
+ pks, n
+ ); // prepares the next batch of signatures on CPU, naturally blocks until it's finished
+
+ state->result += modmult_opencl_results(info, state, false, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n); // waits for kernel, if it is not ready yet
+ modmult_opencl_exec_kernel(info,state); // start kernel (returns immediately)
+ modmult_opencl_release(state); // release buffers
+
+
+
+}
+
+unsigned long modmult_gpu_finish(struct gpu_info *info, struct gpu_state_alt *state, void * msg_buf, void * r_1_buf, void * n_buf, void *mod_buf, void* s_buf, unsigned long n) {
+
+ state->result += modmult_opencl_results(info,state,true,msg_buf,r_1_buf,n_buf, s_buf, mod_buf, n);
+
+ modmult_opencl_cleanup(info);
+ //clReleaseMemObject(state->invalid);
+
+ unsigned long res = state->result;
+
+ state->result = 0; // reset result
+
+ return res;
+
+}
+
+
+int mont_modmult_tests(void) {
+
+ unsigned long pairs = number_of_pairs(); // returns an estimation of pairs
+
+ unsigned long digit_sz = (BIT_LENGTH / 8) * pairs;
+
+ len_in_bytes = digit_sz;
+
+ unsigned long arr_sz = pairs * sizeof(unsigned long);
+
+ gpu_register *x_buf = malloc(digit_sz);
+ gpu_register *m_buf = malloc(digit_sz);
+ gpu_register *r_1_buf = malloc(digit_sz);
+ gpu_register *n_buf = malloc(digit_sz);
+ gpu_register *ni_buf = malloc(digit_sz);
+ gpu_register *msg_buf = malloc(digit_sz);
+ gpu_register *s_buf = malloc(digit_sz);
+ gpu_register *mod_buf = malloc(digit_sz);
+ gpu_register *exp_buf = malloc(pairs * sizeof(gpu_register));
+
+ memset(x_buf, 0, digit_sz);
+ memset(m_buf, 0, digit_sz);
+ memset(r_1_buf, 0, digit_sz);
+ memset(n_buf, 0, digit_sz);
+ memset(ni_buf, 0, digit_sz);
+ memset(msg_buf, 0, digit_sz);
+ memset(s_buf, 0, digit_sz);
+ memset(mod_buf, 0, digit_sz);
+ memset(exp_buf, 0, digit_sz * sizeof(gpu_register));
+
+ unsigned long *pks = malloc(arr_sz);
+
+ memset(pks, 0, arr_sz);
+
+ printf("READING KEYS...\n");
+
+ montmodmult_pairs_from_files(x_buf, m_buf,
+ r_1_buf,
+ n_buf, ni_buf,
+ msg_buf, exp_buf, mod_buf, s_buf,
+ pks, &pairs);
+
+ printf("VERIFYING %lu SIGNATURES...\n", pairs);
+
+ struct gpu_info info;
+ struct gpu_state_alt state;
+
+ modmult_gpu_init(&info, &state);
+
+ modmult_gpu_execute(&info, &state,
+ x_buf, m_buf,
+ r_1_buf,
+ n_buf, ni_buf,
+ msg_buf, exp_buf, mod_buf, s_buf,
+ pks, pairs);
+
+ unsigned long res = modmult_gpu_finish(&info, &state, msg_buf, r_1_buf, n_buf, mod_buf,s_buf, pairs);
+
+ if (res == pairs) {
+ printf("VERIFICATION RESULT: %lu - OK\n\n",res);
+ } else {
+ printf("VERIFICATION RESULT: %lu - NOT OK\n\n",res);
+ }
+
+
+
+ printf("");
+
+}
diff --git a/source/montmodmult.h b/source/montmodmult.h
@@ -9,5 +9,12 @@
#define montmodmult_h
#include <stdio.h>
+#include "ctype.h"
+#include <gmp.h>
+
+int mont_modmult_tests(void);
+
+
+typedef u_int64_t gpu_register;
#endif /* montmodmult_h */
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -335,7 +335,7 @@ int opencl_prepare(struct gpu_info *info, struct gpu_state *state,
err |= clSetKernelArg(info->kernel, 5, sizeof(cl_mem), &state->mod_len);
err |= clSetKernelArg(info->kernel, 6, sizeof(cl_mem), &state->comp_mem);
err |= clSetKernelArg(info->kernel, 7, sizeof(cl_mem), &state->comp_len);
- //err |= clSetKernelArg(info->kernel, 8, sizeof(cl_mem), &state->valid);
+ //err |= clSetKernelArg(info->kernel, 8, sizeof(cl_mem), &state->valid); -> set somewhere else
err |= clSetKernelArg(info->kernel, 9, sizeof(cl_mem), &state->pks_indices);
err |= clSetKernelArg(info->kernel, 10, sizeof(unsigned long), &n);
diff --git a/source/rsa-test.h b/source/rsa-test.h
@@ -12,41 +12,6 @@
#include "ctype.h"
-struct gpu_info {
-
- cl_platform_id platform;
- cl_device_id device_id;
- cl_context context;
- cl_command_queue commands;
- cl_program program;
- cl_kernel kernel;
-
-};
-
-struct gpu_state {
-
- cl_mem invalid;
-
- cl_mem sig_mem;
- cl_mem exp_mem;
- cl_mem mod_mem;
- cl_mem comp_mem;
-
- cl_mem sig_len;
- cl_mem exp_len;
- cl_mem mod_len;
- cl_mem comp_len;
-
- cl_mem pks_indices;
-
- struct timespec t1, t2;
-
- unsigned long total; // 0
- unsigned long result; // 0
-
- bool skip;
-
-};
int rsa_tests(void);
diff --git a/source/util.h b/source/util.h
@@ -66,4 +66,66 @@ create_kernel (cl_program program,
const char *name);
+struct gpu_info {
+
+ cl_platform_id platform;
+ cl_device_id device_id;
+ cl_context context;
+ cl_command_queue commands;
+ cl_program program;
+ cl_kernel kernel;
+
+};
+
+struct gpu_state {
+
+ cl_mem invalid;
+
+ cl_mem sig_mem;
+ cl_mem exp_mem;
+ cl_mem mod_mem;
+ cl_mem comp_mem;
+
+ cl_mem sig_len;
+ cl_mem exp_len;
+ cl_mem mod_len;
+ cl_mem comp_len;
+
+ cl_mem pks_indices;
+
+ struct timespec t1, t2;
+
+ unsigned long total; // 0
+ unsigned long result; // 0
+
+ bool skip;
+
+};
+
+struct gpu_state_alt {
+
+ // cl_mem invalid;
+
+ cl_mem x_mem;
+ cl_mem m_mem;
+ cl_mem res_mem;
+ cl_mem n_mem;
+ cl_mem ni_mem;
+ //cl_mem msg_mem;
+ cl_mem exp_mem;
+
+ cl_mem tmp_1;
+ cl_mem tmp_2;
+
+ cl_mem pks_indices;
+
+ struct timespec t1, t2;
+
+ unsigned long total; // 0
+ unsigned long result; // 0
+
+ bool skip;
+
+};
+
#endif /* util_h */
diff --git a/xcode/.DS_Store b/xcode/.DS_Store
Binary files differ.
diff --git a/xcode/lib-gpu-generate/msgsig.txt b/xcode/lib-gpu-generate/msgsig.txt
@@ -1,2 +1,2 @@
-3DB82611F0AA4549
-0083ECD7379021AE99FFCF028CC1EBF96542B931D8C3F245EF64FE7423238D90D423C1F33E1D6CA24C0ADB8302D18A98DAAB565E49E81ED08937D3C7B7AF4AA1F832397EC0EFF11118AA159AF722B030EA5F614CE458DE575A4DEC70A94539B789E60F5B6CF0F146B1FB01C5F0583571BB21592E08E195488FE83AB18E40003A7285DA710157FFB64D577051B3BD88BA0B867CA402F910BDF35440F7A83024E4A261AD6E79E3B3312C0BDB23EEA6DBFA6A81CFB2635AE801815B24EED2E08A0CA4877037776CF0FB5079BF12CE721DBBB66DE87C5ACB24E01AC751E5853AEEE6BEA1CC60E1A68836108EA11B79BC4FEB2F2F47591C1A2E552C74BAA714382610EC
+9506E0DF73
+7134105D4E8706C01F2EDE5C773E2A250A269D2DA99BAB49AC8D5FF9AEA2E059D4D66F514452DEA37AB39E2D5B65281CB6566850548E5DCA87ABD5FC171D56BD1DB449A257D1CA2ECBA66C8232B4A2F1E590427CA3E451188A036A49346D8AFFA48A4F90DA6D26B22CECEE53FD8951CA101D937598C292B380CE836C8B6D992ABE0DC5AD4240C9E01BD441F359374C15A0D30BC296F56C79D85157AC00B339A8C4B2414300A131B023ED30EDFFD9CFA6FEBCBC724D5C88C3FC89815654526D28D41871AC7FC4514F4FCC6789404A3F13661A3DA4B311AB6AF547673D5FD13E544377A8960E6360E76DEF3788BD3356720EF514DCD2BBD1C4950385A2104BA984
diff --git a/xcode/lib-gpu-generate/publickey.txt b/xcode/lib-gpu-generate/publickey.txt
@@ -1,3 +1,3 @@
-00AACCE8836848B3FE12D67F155B2EF70734C8D6E374C78300167C5F3062CC76D43F0E74281142784343BE313261A50B4B23CBC45D5F8BF88025C07BD4F02FC557E370A47775820307845F179F49592EC466C6FC33A59F9E05B1B70BCC1F5F62431B5BFF0F02377A1C94E4B8E92C249402B74EAD7CAF30D4DA2144B9B1357A7A9944540E1A1D1805E767C693B0D3D3267072DCF51AE61A0AB08E453095C3A623DD50C727EF2F74788DAE6E73613F5D7329E68B9BB3BBE770E1093FCE9A43B32DFC2975B648D9A2A89A0FED9DB41C7898FA3AEFC35602A8A96C1E84C9766C06E1F58641010EF8FCB7B93709795678988A9EDA6191D43E927D13AED334B964D05F3B
-010001
+00CA5F6FD970F17AAEA1F2272368C1CE27E2252F7357C651AAC203EA012D888CED38F4F1B97DB332E005E291D845DECBD3846A3E09F0FEBCDC71D31A735874F9848974814474055FA200F6D42B715AFD9E04C4DA1FF2457020EEB4753E44318AF4E41111923B007FCBB1E2E17898E00AD500BFF849E7360B0249DA98A3D7301E12C6A2C33925078E783336263AC07A3EED6594414736E01DFA621CBB93E1C548E539AF8D0FB87AE0F8181B9C38046E01508F2865333305BDC8C177F3D7A1CF677F11F6115D45792AFF2BAFFB36CC700757AC262C29F76D0869DF5B6258D754237DBE492236251599D3ACB500F0764021881BD1F1777143F6D73B63D7B9AD67CB2B
+02
0
diff --git a/xcode/lib-gpu-verify.xcodeproj/project.pbxproj b/xcode/lib-gpu-verify.xcodeproj/project.pbxproj
@@ -9,11 +9,12 @@
/* Begin PBXBuildFile section */
6A36F8892B0F938E00AB772D /* montgomery.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A36F8882B0F938E00AB772D /* montgomery.cl */; };
6A8A795F2A89672700116D7D /* verify.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A8A795E2A89672700116D7D /* verify.cl */; };
- 6A99B06E2B1293DA0004E4B7 /* gmp.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A7914CB2B0CF320001EDCC1 /* gmp.c */; };
6AA38E5B2B0A97FC00E85243 /* main.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AA38E5A2B0A97FC00E85243 /* main.c */; };
6AB4D99D2B1645F900A686F2 /* montgomery-test.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AB4D99C2B1645F900A686F2 /* montgomery-test.c */; };
6ABC2E842B231DFF00033B90 /* util.c in Sources */ = {isa = PBXBuildFile; fileRef = 6ABC2E832B231DFF00033B90 /* util.c */; };
6ABC2E882B231E3D00033B90 /* reference-test.c in Sources */ = {isa = PBXBuildFile; fileRef = 6ABC2E862B231E3D00033B90 /* reference-test.c */; };
+ 6AC553252B2E174900046AB7 /* montmodmult.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6AC553242B2E174900046AB7 /* montmodmult.cl */; };
+ 6AC553292B2E17C800046AB7 /* montmodmult.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AC553282B2E17C800046AB7 /* montmodmult.c */; };
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */; };
6AF748832ADADF4500D58E08 /* rsa-test.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF7487F2ADADF4500D58E08 /* rsa-test.c */; };
C3770EFD0E6F1138009A5A77 /* OpenCL.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = C3770EFC0E6F1138009A5A77 /* OpenCL.framework */; };
@@ -43,10 +44,8 @@
/* Begin PBXFileReference section */
466E0F5F0C932E1A00ED01DB /* lib-gpu-verify */ = {isa = PBXFileReference; explicitFileType = "compiled.mach-o.executable"; includeInIndex = 0; path = "lib-gpu-verify"; sourceTree = BUILT_PRODUCTS_DIR; };
6A36F8882B0F938E00AB772D /* montgomery.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = montgomery.cl; sourceTree = "<group>"; };
- 6A7914CB2B0CF320001EDCC1 /* gmp.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = gmp.c; path = ../source/gmp.c; sourceTree = "<group>"; };
6A7914CC2B0CF320001EDCC1 /* montgomery.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = montgomery.h; path = ../source/montgomery.h; sourceTree = "<group>"; };
6A7914CD2B0CF320001EDCC1 /* montgomery.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = montgomery.c; path = ../source/montgomery.c; sourceTree = "<group>"; };
- 6A7914CE2B0CF320001EDCC1 /* gmp.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = gmp.h; path = ../source/gmp.h; sourceTree = "<group>"; };
6A8A795E2A89672700116D7D /* verify.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = verify.cl; sourceTree = "<group>"; };
6AA38E582B0A97FC00E85243 /* lib-gpu-generate */ = {isa = PBXFileReference; explicitFileType = "compiled.mach-o.executable"; includeInIndex = 0; path = "lib-gpu-generate"; sourceTree = BUILT_PRODUCTS_DIR; };
6AA38E5A2B0A97FC00E85243 /* main.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; path = main.c; sourceTree = "<group>"; };
@@ -57,6 +56,11 @@
6ABC2E852B231E0400033B90 /* util.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = util.h; path = ../source/util.h; sourceTree = "<group>"; };
6ABC2E862B231E3D00033B90 /* reference-test.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "reference-test.c"; path = "../source/reference-test.c"; sourceTree = "<group>"; };
6ABC2E872B231E3D00033B90 /* reference-test.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = "reference-test.h"; path = "../source/reference-test.h"; sourceTree = "<group>"; };
+ 6AC553242B2E174900046AB7 /* montmodmult.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = montmodmult.cl; sourceTree = "<group>"; };
+ 6AC553272B2E17C800046AB7 /* montmodmult.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = montmodmult.h; path = ../source/montmodmult.h; sourceTree = "<group>"; };
+ 6AC553282B2E17C800046AB7 /* montmodmult.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; name = montmodmult.c; path = ../source/montmodmult.c; sourceTree = "<group>"; };
+ 6AC5532A2B2E885200046AB7 /* gmp.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = gmp.c; path = ../source/gmp.c; sourceTree = "<group>"; };
+ 6AC5532C2B2E889100046AB7 /* gmp.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = gmp.h; path = ../source/gmp.h; sourceTree = "<group>"; };
6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "lib-gpu-verify.c"; path = "../source/lib-gpu-verify.c"; sourceTree = "<group>"; };
6AF7487F2ADADF4500D58E08 /* rsa-test.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "rsa-test.c"; path = "../source/rsa-test.c"; sourceTree = "<group>"; };
6AF748802ADADF4500D58E08 /* rsa-test.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = "rsa-test.h"; path = "../source/rsa-test.h"; sourceTree = "<group>"; };
@@ -105,8 +109,8 @@
isa = PBXGroup;
children = (
6AF748802ADADF4500D58E08 /* rsa-test.h */,
- 6A7914CE2B0CF320001EDCC1 /* gmp.h */,
6A7914CC2B0CF320001EDCC1 /* montgomery.h */,
+ 6AC5532C2B2E889100046AB7 /* gmp.h */,
6ABC2E852B231E0400033B90 /* util.h */,
6ABC2E872B231E3D00033B90 /* reference-test.h */,
6AB4D99B2B1645F900A686F2 /* montgomery-test.h */,
@@ -123,6 +127,16 @@
path = "lib-gpu-generate";
sourceTree = "<group>";
};
+ 6AC553262B2E175500046AB7 /* pingpong */ = {
+ isa = PBXGroup;
+ children = (
+ 6AC553242B2E174900046AB7 /* montmodmult.cl */,
+ 6AC553272B2E17C800046AB7 /* montmodmult.h */,
+ 6AC553282B2E17C800046AB7 /* montmodmult.c */,
+ );
+ name = pingpong;
+ sourceTree = "<group>";
+ };
C3770EF10E6F10BB009A5A77 /* Sources */ = {
isa = PBXGroup;
children = (
@@ -131,11 +145,12 @@
6A36F8882B0F938E00AB772D /* montgomery.cl */,
6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */,
6AF7487F2ADADF4500D58E08 /* rsa-test.c */,
- 6A7914CB2B0CF320001EDCC1 /* gmp.c */,
+ 6AC5532A2B2E885200046AB7 /* gmp.c */,
6A7914CD2B0CF320001EDCC1 /* montgomery.c */,
6AB4D99C2B1645F900A686F2 /* montgomery-test.c */,
6ABC2E832B231DFF00033B90 /* util.c */,
6ABC2E862B231E3D00033B90 /* reference-test.c */,
+ 6AC553262B2E175500046AB7 /* pingpong */,
);
name = Sources;
sourceTree = "<group>";
@@ -227,7 +242,8 @@
files = (
6AB4D99D2B1645F900A686F2 /* montgomery-test.c in Sources */,
6ABC2E882B231E3D00033B90 /* reference-test.c in Sources */,
- 6A99B06E2B1293DA0004E4B7 /* gmp.c in Sources */,
+ 6AC553252B2E174900046AB7 /* montmodmult.cl in Sources */,
+ 6AC553292B2E17C800046AB7 /* montmodmult.c in Sources */,
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */,
6A8A795F2A89672700116D7D /* verify.cl in Sources */,
6ABC2E842B231DFF00033B90 /* util.c in Sources */,
@@ -345,7 +361,10 @@
);
LIBRARY_SEARCH_PATHS_QUOTED_FOR_TARGET_1 = "\"$(SRCROOT)/../lib\" \"/usr/local/lib\"";
MACOSX_DEPLOYMENT_TARGET = 13.5;
- OTHER_LDFLAGS = "-lgcrypt";
+ OTHER_LDFLAGS = (
+ "-lgcrypt",
+ "-lgmp",
+ );
PRODUCT_NAME = "lib-gpu-verify";
SYSTEM_FRAMEWORK_SEARCH_PATHS = "";
USE_HEADERMAP = NO;
@@ -372,7 +391,10 @@
);
LIBRARY_SEARCH_PATHS_QUOTED_FOR_TARGET_1 = "\"$(SRCROOT)/../lib\" \"/usr/local/lib\"";
MACOSX_DEPLOYMENT_TARGET = 13.5;
- OTHER_LDFLAGS = "-lgcrypt";
+ OTHER_LDFLAGS = (
+ "-lgcrypt",
+ "-lgmp",
+ );
PRODUCT_NAME = "lib-gpu-verify";
USE_HEADERMAP = NO;
ZERO_LINK = NO;
diff --git a/xcode/lib-gpu-verify.xcodeproj/project.xcworkspace/xcuserdata/cedriczwahlen.xcuserdatad/UserInterfaceState.xcuserstate b/xcode/lib-gpu-verify.xcodeproj/project.xcworkspace/xcuserdata/cedriczwahlen.xcuserdatad/UserInterfaceState.xcuserstate
Binary files differ.
diff --git a/xcode/lib-gpu-verify.xcodeproj/xcuserdata/cedriczwahlen.xcuserdatad/xcdebugger/Breakpoints_v2.xcbkptlist b/xcode/lib-gpu-verify.xcodeproj/xcuserdata/cedriczwahlen.xcuserdatad/xcdebugger/Breakpoints_v2.xcbkptlist
@@ -4030,5 +4030,1506 @@
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "B5F8D09B-DBD6-4D27-9F18-0AA02E970656"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "21"
+ endingLineNumber = "21"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "3947C3E7-D4B7-443E-9C64-26FA73518839"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "115"
+ endingLineNumber = "115"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "3947C3E7-D4B7-443E-9C64-26FA73518839 - 18037aea005af134"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "92"
+ endingLineNumber = "92"
+ offsetFromSymbolStart = "928">
+ </Location>
+ <Location
+ uuid = "3947C3E7-D4B7-443E-9C64-26FA73518839 - 18037aea005af115"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "91"
+ endingLineNumber = "91"
+ offsetFromSymbolStart = "1023">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "700C8EE6-BC1E-427F-B1C6-E018D10BDEC8"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "558"
+ endingLineNumber = "558"
+ landmarkName = "mont_modmult_tests()"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "700C8EE6-BC1E-427F-B1C6-E018D10BDEC8 - 3a6dd6ae57b14a5e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "462"
+ endingLineNumber = "462"
+ offsetFromSymbolStart = "23">
+ </Location>
+ <Location
+ uuid = "700C8EE6-BC1E-427F-B1C6-E018D10BDEC8 - 3a6dd6ae57b14a5e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "462"
+ endingLineNumber = "462"
+ offsetFromSymbolStart = "14">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "DB1A1F80-9BCD-4641-8418-19BE4BCBC405"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "100"
+ endingLineNumber = "100"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "605"
+ endingLineNumber = "605"
+ landmarkName = "mont_modmult_tests()"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b149b8"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "492"
+ endingLineNumber = "492"
+ offsetFromSymbolStart = "437">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b149b8"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "492"
+ endingLineNumber = "492"
+ offsetFromSymbolStart = "473">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b14e3c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "496"
+ endingLineNumber = "496"
+ offsetFromSymbolStart = "473">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b151d5"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "551"
+ endingLineNumber = "551"
+ offsetFromSymbolStart = "488">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b15742"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "562"
+ endingLineNumber = "562"
+ offsetFromSymbolStart = "488">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b15742"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "562"
+ endingLineNumber = "562"
+ offsetFromSymbolStart = "538">
+ </Location>
+ <Location
+ uuid = "D2BA3F25-8290-42F4-8C03-1F8F1A3C497D - 3a6dd6ae57b15b85"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_modmult_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "599"
+ endingLineNumber = "599"
+ offsetFromSymbolStart = "538">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "877D065B-1222-4A88-A7A1-36AA5BAAA34B"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "226"
+ endingLineNumber = "226"
+ landmarkName = "modmult_opencl_prepare(info, state, x_buf, m_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "877D065B-1222-4A88-A7A1-36AA5BAAA34B - 7ea67c6850ea9357"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "194"
+ endingLineNumber = "194"
+ offsetFromSymbolStart = "108">
+ </Location>
+ <Location
+ uuid = "877D065B-1222-4A88-A7A1-36AA5BAAA34B - 7ea67c6850ea90f0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "199"
+ endingLineNumber = "199"
+ offsetFromSymbolStart = "104">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "0919DE7F-CA68-40D2-98C2-CDC334661F60"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "216"
+ endingLineNumber = "216"
+ landmarkName = "modmult_opencl_prepare(info, state, x_buf, m_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "4B27D588-7D34-4BD1-A24D-4C6E8D39B984"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "440"
+ endingLineNumber = "440"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "4B27D588-7D34-4BD1-A24D-4C6E8D39B984 - b46ebf1cc4a0b9c0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "346"
+ endingLineNumber = "346"
+ offsetFromSymbolStart = "269">
+ </Location>
+ <Location
+ uuid = "4B27D588-7D34-4BD1-A24D-4C6E8D39B984 - b46ebf1cc4a0b9c0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "346"
+ endingLineNumber = "346"
+ offsetFromSymbolStart = "259">
+ </Location>
+ <Location
+ uuid = "4B27D588-7D34-4BD1-A24D-4C6E8D39B984 - b46ebf1cc4a0b9c0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "346"
+ endingLineNumber = "346"
+ offsetFromSymbolStart = "256">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "800CDD46-DF8C-4E7E-A8E0-2BC47B0E60D2"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "228"
+ endingLineNumber = "228"
+ landmarkName = "modmult_opencl_prepare(info, state, x_buf, m_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "800CDD46-DF8C-4E7E-A8E0-2BC47B0E60D2 - 7ea67c6850ea90d3"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "198"
+ endingLineNumber = "198"
+ offsetFromSymbolStart = "119">
+ </Location>
+ <Location
+ uuid = "800CDD46-DF8C-4E7E-A8E0-2BC47B0E60D2 - 7ea67c6850ea90be"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "201"
+ endingLineNumber = "201"
+ offsetFromSymbolStart = "116">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "388"
+ endingLineNumber = "388"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0ba5c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "358"
+ endingLineNumber = "358"
+ offsetFromSymbolStart = "361">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0ba7d"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "357"
+ endingLineNumber = "357"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0ba3f"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "359"
+ endingLineNumber = "359"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0ba5c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "358"
+ endingLineNumber = "358"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0ba1e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "360"
+ endingLineNumber = "360"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bbd0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "362"
+ endingLineNumber = "362"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bb75"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "365"
+ endingLineNumber = "365"
+ offsetFromSymbolStart = "351">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a041ce"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "408"
+ endingLineNumber = "408"
+ offsetFromSymbolStart = "369">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0be9b"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "387"
+ endingLineNumber = "387"
+ offsetFromSymbolStart = "394">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bd45"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "381"
+ endingLineNumber = "381"
+ offsetFromSymbolStart = "370">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bed9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "385"
+ endingLineNumber = "385"
+ offsetFromSymbolStart = "370">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bed9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "385"
+ endingLineNumber = "385"
+ offsetFromSymbolStart = "374">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0bed9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "385"
+ endingLineNumber = "385"
+ offsetFromSymbolStart = "398">
+ </Location>
+ <Location
+ uuid = "84A684AB-D668-4B78-9FD1-2540A08574E3 - b46ebf1cc4a0be7a"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "388"
+ endingLineNumber = "388"
+ offsetFromSymbolStart = "398">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "FAC6BA4D-8703-4276-BEAA-7EA43AF63D0F"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "326"
+ endingLineNumber = "326"
+ landmarkName = "modmult_opencl_exec_kernel(info, state)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "FAC6BA4D-8703-4276-BEAA-7EA43AF63D0F - 22ed8a6bec85ee99"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "298"
+ endingLineNumber = "298"
+ offsetFromSymbolStart = "28">
+ </Location>
+ <Location
+ uuid = "FAC6BA4D-8703-4276-BEAA-7EA43AF63D0F - 22ed8a6bec85ee99"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "298"
+ endingLineNumber = "298"
+ offsetFromSymbolStart = "32">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "323"
+ endingLineNumber = "323"
+ landmarkName = "modmult_opencl_exec_kernel(info, state)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802 - 22ed8a6bec85ee15"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "294"
+ endingLineNumber = "294"
+ offsetFromSymbolStart = "24">
+ </Location>
+ <Location
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802 - 22ed8a6bec85eed7"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "296"
+ endingLineNumber = "296"
+ offsetFromSymbolStart = "24">
+ </Location>
+ <Location
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802 - 22ed8a6bec85e8c6"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "311"
+ endingLineNumber = "311"
+ offsetFromSymbolStart = "24">
+ </Location>
+ <Location
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802 - 22ed8a6bec85e487"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "344"
+ endingLineNumber = "344"
+ offsetFromSymbolStart = "24">
+ </Location>
+ <Location
+ uuid = "5B1B4488-A19F-4A97-926D-1263BFEFF802 - 22ed8a6bec85ea71"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_exec_kernel"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "322"
+ endingLineNumber = "322"
+ offsetFromSymbolStart = "24">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "E19DB0AE-EBC4-42E6-A48E-0A5B64A9CF56"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "73"
+ endingLineNumber = "73"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "5AA79044-F9D9-4539-870E-8082CCF4F458"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "532"
+ endingLineNumber = "532"
+ landmarkName = "modmult_gpu_execute(info, state, x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "5AA79044-F9D9-4539-870E-8082CCF4F458 - 30fdc229490fff9e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_gpu_execute"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "477"
+ endingLineNumber = "477"
+ offsetFromSymbolStart = "117">
+ </Location>
+ <Location
+ uuid = "5AA79044-F9D9-4539-870E-8082CCF4F458 - 30fdc229490fffdc"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_gpu_execute"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "479"
+ endingLineNumber = "479"
+ offsetFromSymbolStart = "117">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "29FF3D84-E5F9-4429-8A53-756F9A7E5AD4"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "259"
+ endingLineNumber = "259"
+ landmarkName = "modmult_opencl_prepare(info, state, x_buf, m_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "CCF0BFD9-35C8-47D4-A9DE-665DC105BDF8"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "137"
+ endingLineNumber = "137"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "CCF0BFD9-35C8-47D4-A9DE-665DC105BDF8 - 18037aea005af767"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "137"
+ endingLineNumber = "137"
+ offsetFromSymbolStart = "1406">
+ </Location>
+ <Location
+ uuid = "CCF0BFD9-35C8-47D4-A9DE-665DC105BDF8 - 18037aea005af725"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "139"
+ endingLineNumber = "139"
+ offsetFromSymbolStart = "1406">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "854BEA05-AC22-4D30-B923-D574599C490D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "154"
+ endingLineNumber = "154"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "854BEA05-AC22-4D30-B923-D574599C490D - 18037aea005af4a9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "135"
+ endingLineNumber = "135"
+ offsetFromSymbolStart = "1465">
+ </Location>
+ <Location
+ uuid = "854BEA05-AC22-4D30-B923-D574599C490D - 18037aea005af63e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "146"
+ endingLineNumber = "146"
+ offsetFromSymbolStart = "1590">
+ </Location>
+ <Location
+ uuid = "854BEA05-AC22-4D30-B923-D574599C490D - 18037aea005af640"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "144"
+ endingLineNumber = "144"
+ offsetFromSymbolStart = "1665">
+ </Location>
+ <Location
+ uuid = "854BEA05-AC22-4D30-B923-D574599C490D - 18037aea005ac936"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "154"
+ endingLineNumber = "154"
+ offsetFromSymbolStart = "1730">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "56542C6C-ED16-48E5-A345-C0E798659638"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "378"
+ endingLineNumber = "378"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "56542C6C-ED16-48E5-A345-C0E798659638 - b46ebf1cc4a0ba5c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "358"
+ endingLineNumber = "358"
+ offsetFromSymbolStart = "315">
+ </Location>
+ <Location
+ uuid = "56542C6C-ED16-48E5-A345-C0E798659638 - b46ebf1cc4a0bcc9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "369"
+ endingLineNumber = "369"
+ offsetFromSymbolStart = "315">
+ </Location>
+ <Location
+ uuid = "56542C6C-ED16-48E5-A345-C0E798659638 - b46ebf1cc4a0bb54"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "366"
+ endingLineNumber = "366"
+ offsetFromSymbolStart = "315">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "A0B20FE3-FB60-4D60-9B51-300DE8543447"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "364"
+ endingLineNumber = "364"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "A0B20FE3-FB60-4D60-9B51-300DE8543447 - b46ebf1cc4a0b84c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "342"
+ endingLineNumber = "342"
+ offsetFromSymbolStart = "148">
+ </Location>
+ <Location
+ uuid = "A0B20FE3-FB60-4D60-9B51-300DE8543447 - b46ebf1cc4a0baf9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "353"
+ endingLineNumber = "353"
+ offsetFromSymbolStart = "148">
+ </Location>
+ <Location
+ uuid = "A0B20FE3-FB60-4D60-9B51-300DE8543447 - b46ebf1cc4a0b944"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "350"
+ endingLineNumber = "350"
+ offsetFromSymbolStart = "148">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "175"
+ endingLineNumber = "175"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005ac82e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "162"
+ endingLineNumber = "162"
+ offsetFromSymbolStart = "1892">
+ </Location>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005acb83"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "173"
+ endingLineNumber = "173"
+ offsetFromSymbolStart = "2017">
+ </Location>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005acb07"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "169"
+ endingLineNumber = "169"
+ offsetFromSymbolStart = "2078">
+ </Location>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005acb07"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "169"
+ endingLineNumber = "169"
+ offsetFromSymbolStart = "2075">
+ </Location>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005acb07"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "169"
+ endingLineNumber = "169"
+ offsetFromSymbolStart = "2072">
+ </Location>
+ <Location
+ uuid = "6B6DD1A1-289E-455E-B49C-DDF8158E3B10 - 18037aea005aca41"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "175"
+ endingLineNumber = "175"
+ offsetFromSymbolStart = "2194">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "370"
+ endingLineNumber = "370"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0ba9a"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "356"
+ endingLineNumber = "356"
+ offsetFromSymbolStart = "272">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bfd1"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "393"
+ endingLineNumber = "393"
+ offsetFromSymbolStart = "272">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bf93"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "395"
+ endingLineNumber = "395"
+ offsetFromSymbolStart = "272">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bf55"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "397"
+ endingLineNumber = "397"
+ offsetFromSymbolStart = "272">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bf55"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "397"
+ endingLineNumber = "397"
+ offsetFromSymbolStart = "290">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0be1f"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "391"
+ endingLineNumber = "391"
+ offsetFromSymbolStart = "290">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bcc9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "369"
+ endingLineNumber = "369"
+ offsetFromSymbolStart = "306">
+ </Location>
+ <Location
+ uuid = "5C0ACEE0-4246-4CA2-9B7C-A5B3E417DD3E - b46ebf1cc4a0bbb3"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "modmult_opencl_results"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "363"
+ endingLineNumber = "363"
+ offsetFromSymbolStart = "282">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "1D13B1BF-088A-44A4-B036-D2267F8385C5"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "133"
+ endingLineNumber = "133"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "1D13B1BF-088A-44A4-B036-D2267F8385C5 - 18037aea005af42d"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "131"
+ endingLineNumber = "131"
+ offsetFromSymbolStart = "1360">
+ </Location>
+ <Location
+ uuid = "1D13B1BF-088A-44A4-B036-D2267F8385C5 - 18037aea005af4eb"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montmodmult_pairs_from_files"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "133"
+ endingLineNumber = "133"
+ offsetFromSymbolStart = "1375">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "FB38C9CB-9FB8-4B84-9CE2-7184FCD7F326"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "144"
+ endingLineNumber = "144"
+ landmarkName = "mont_prepare(b, e, m, r, r_1, ni, M, x)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "EC8E40B9-5869-4CD4-A602-C7739D7BFF03"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "251"
+ endingLineNumber = "251"
+ landmarkName = "opencl_pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "E4BEB202-36D5-4E45-9886-88D8B2A6D37B"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "392"
+ endingLineNumber = "392"
+ landmarkName = "modmult_opencl_results(info, state, timed, msg_buf, r_1_buf, n_buf, s_buf, mod_buf, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "EBE3EB24-9B9A-43BD-850A-460559601A71"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "168"
+ endingLineNumber = "168"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "7E182A4C-319B-4622-9FD0-18C77CB42D7B"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montmodmult.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "155"
+ endingLineNumber = "155"
+ landmarkName = "montmodmult_pairs_from_files(x_buf, m_buf, r_1_buf, n_buf, ni_buf, msg_buf, exp_buf, mod_buf, s_buf, pks, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
</Breakpoints>
</Bucket>
diff --git a/xcode/montmodmult.cl b/xcode/montmodmult.cl
@@ -0,0 +1,375 @@
+
+
+typedef ulong gpu_register;
+
+#define BITS_PER_DIGIT (sizeof(gpu_register) * 8)
+#define HIBITMASK 0x80000000UL
+
+
+int mult(gpu_register p[2], gpu_register x, gpu_register y)
+{
+
+
+
+
+ /* Use a 64-bit temp for product */
+ //ulong t = (ulong)x * (ulong)y;
+ /* then split into two parts */
+ p[1] = mul_hi(x,y);
+ p[0] = x * y;
+
+ return 0;
+}
+
+
+int multiply(__global gpu_register *w, __global gpu_register *u, __global gpu_register *v, size_t ndigits)
+{
+ /* Computes product w = u * v
+ where u, v are multiprecision integers of ndigits each
+ and w is a multiprecision integer of 2*ndigits
+
+ Ref: Knuth Vol 2 Ch 4.3.1 p 268 Algorithm M.
+ */
+
+ gpu_register k, t[2];
+ size_t i, j, m, n;
+
+ //assert(w != u && w != v);
+
+ m = n = ndigits;
+
+ /* Step M1. Initialise */
+ for (i = 0; i < 2 * m; i++)
+ w[i] = 0;
+
+ for (j = 0; j < n; j++)
+ {
+ /* Step M2. Zero multiplier? */
+ if (v[j] == 0)
+ {
+ w[j + m] = 0;
+ }
+ else
+ {
+ /* Step M3. Initialise i */
+ k = 0;
+ for (i = 0; i < m; i++)
+ {
+ /* Step M4. Multiply and add */
+ /* t = u_i * v_j + w_(i+j) + k */
+ mult(t, u[i], v[j]);
+
+ t[0] += k;
+ if (t[0] < k)
+ t[1]++;
+ t[0] += w[i+j];
+ if (t[0] < w[i+j])
+ t[1]++;
+
+ w[i+j] = t[0];
+ k = t[1];
+ }
+ /* Step M5. Loop on i, set w_(j+m) = k */
+ w[j+m] = k;
+ }
+ } /* Step M6. Loop on j */
+
+ return 0;
+}
+
+
+int square(__global gpu_register *w, __global gpu_register *x, size_t ndigits)
+/* New in Version 2.0 */
+{
+ /* Computes square w = x * x
+ where x is a multiprecision integer of ndigits
+ and w is a multiprecision integer of 2*ndigits
+
+ Ref: Menezes p596 Algorithm 14.16 with errata.
+ */
+
+ gpu_register k, p[2], u[2], cbit, carry;
+ size_t i, j, t, i2, cpos;
+
+ t = ndigits;
+
+ /* 1. For i from 0 to (2t-1) do: w_i = 0 */
+ i2 = t << 1;
+ for (i = 0; i < i2; i++)
+ w[i] = 0;
+
+ carry = 0;
+ cpos = i2-1;
+ /* 2. For i from 0 to (t-1) do: */
+ for (i = 0; i < t; i++)
+ {
+ /* 2.1 (uv) = w_2i + x_i * x_i, w_2i = v, c = u
+ Careful, w_2i may be double-prec
+ */
+ i2 = i << 1; /* 2*i */
+ mult(p, x[i], x[i]);
+ p[0] += w[i2];
+ if (p[0] < w[i2])
+ p[1]++;
+ k = 0; /* p[1] < b, so no overflow here */
+ if (i2 == cpos && carry)
+ {
+ p[1] += carry;
+ if (p[1] < carry)
+ k++;
+ carry = 0;
+ }
+ w[i2] = p[0];
+ u[0] = p[1];
+ u[1] = k;
+
+ /* 2.2 for j from (i+1) to (t-1) do:
+ (uv) = w_{i+j} + 2x_j * x_i + c,
+ w_{i+j} = v, c = u,
+ u is double-prec
+ w_{i+j} is dbl if [i+j] == cpos
+ */
+ k = 0;
+ for (j = i+1; j < t; j++)
+ {
+ /* p = x_j * x_i */
+ mult(p, x[j], x[i]);
+ /* p = 2p <=> p <<= 1 */
+ cbit = (p[0] & HIBITMASK) != 0;
+ k = (p[1] & HIBITMASK) != 0;
+ p[0] <<= 1;
+ p[1] <<= 1;
+ p[1] |= cbit;
+ /* p = p + c */
+ p[0] += u[0];
+ if (p[0] < u[0])
+ {
+ p[1]++;
+ if (p[1] == 0)
+ k++;
+ }
+ p[1] += u[1];
+ if (p[1] < u[1])
+ k++;
+ /* p = p + w_{i+j} */
+ p[0] += w[i+j];
+ if (p[0] < w[i+j])
+ {
+ p[1]++;
+ if (p[1] == 0)
+ k++;
+ }
+ if ((i+j) == cpos && carry)
+ { /* catch overflow from last round */
+ p[1] += carry;
+ if (p[1] < carry)
+ k++;
+ carry = 0;
+ }
+ /* w_{i+j} = v, c = u */
+ w[i+j] = p[0];
+ u[0] = p[1];
+ u[1] = k;
+ }
+ /* 2.3 w_{i+t} = u */
+ w[i+t] = u[0];
+ /* remember overflow in w_{i+t} */
+ carry = u[1];
+ cpos = i+t;
+ }
+
+ /* (NB original step 3 deleted in Menezes errata) */
+
+ /* Return w */
+
+ return 0;
+}
+
+gpu_register shift_r(__global gpu_register *a, __global gpu_register *b, size_t shift, size_t ndigits)
+{ /* Computes a = b >> shift */
+ /* [v2.1] Modified to cope with shift > BITS_PERDIGIT */
+
+ gpu_register carry = 0;
+
+ while (1) {
+
+ size_t i, y, nw, bits;
+ gpu_register mask, tempCarry, nextcarry;
+
+ /* Do we shift whole digits? */
+ if (shift >= BITS_PER_DIGIT)
+ {
+ nw = shift / BITS_PER_DIGIT;
+ for (i = 0; i < ndigits; i++)
+ {
+ if ((i+nw) < ndigits)
+ a[i] = b[i+nw];
+ else
+ a[i] = 0;
+ }
+ /* Call again to shift bits inside digits */
+ bits = shift % BITS_PER_DIGIT;
+ tempCarry = b[nw-1] >> bits;
+ if (bits)
+ carry |= tempCarry;
+ return carry;
+ }
+ else
+ {
+ bits = shift;
+ }
+
+ /* Construct mask to set low bits */
+ /* (thanks to Jesse Chisholm for suggesting this improved technique) */
+ mask = ~(~(gpu_register)0 << bits);
+
+ y = BITS_PER_DIGIT - bits;
+ carry = 0;
+ i = ndigits;
+ while (i--)
+ {
+ nextcarry = (b[i] & mask) << y;
+ a[i] = b[i] >> bits | carry;
+ carry = nextcarry;
+ }
+
+ return carry;
+
+ }
+}
+
+/*
+ determine by how much to divide / shift
+ from most significant of double size number, look for first high bit.
+ */
+
+/*
+gpu_register shift_by(__global gpu_register *r, int n) {
+
+ gpu_register u = 8 * sizeof(gpu_register);
+
+ for (int i = n-1; i >= (n / 2); i--)
+ {
+ if (r[i] != 0) {
+ int k = ceil(log2((float)r[i] + (float)1));
+ gpu_register total = k + i * u;
+
+ return total - (u * (n / 2));
+ }
+ }
+
+ return 0;
+}
+*/
+gpu_register add(__global gpu_register *w, __global gpu_register *u, __global gpu_register *v, size_t ndigits)
+{
+ /* Calculates w = u + v
+ where w, u, v are multiprecision integers of ndigits each
+ Returns carry if overflow. Carry = 0 or 1.
+
+ Ref: Knuth Vol 2 Ch 4.3.1 p 266 Algorithm A.
+ */
+
+ gpu_register k;
+ size_t j;
+
+ //assert(w != v);
+
+ /* Step A1. Initialise */
+ k = 0;
+
+ for (j = 0; j < ndigits; j++)
+ {
+ /* Step A2. Add digits w_j = (u_j + v_j + k)
+ Set k = 1 if carry (overflow) occurs
+ */
+ w[j] = u[j] + k;
+ if (w[j] < k)
+ k = 1;
+ else
+ k = 0;
+
+ w[j] += v[j];
+ if (w[j] < v[j])
+ k++;
+
+ } /* Step A3. Loop on j */
+
+ return k; /* w_n = k */
+}
+
+void equal(__global gpu_register *a, __global gpu_register *b, size_t ndigits)
+{ /* Sets a = b */
+ size_t i;
+
+ for (i = 0; i < ndigits; i++)
+ {
+ a[i] = b[i];
+ }
+}
+
+void erase_above(__global gpu_register *a, size_t n)
+{
+
+ for (int i = n-1; i >= (n / 2); i--)
+ {
+ a[i] = 0;
+ }
+
+}
+
+int testbit(gpu_register e, int i) {
+
+ return (e & (0x1 << (gpu_register)i) ) > 0 ? 1 : 0;
+
+}
+
+__kernel
+ void mont(__global gpu_register *x, __global gpu_register *m,
+ __global gpu_register *res, __global gpu_register *n, // res is not needed, we write the result in x
+ __global gpu_register *ni, __global gpu_register *exp,
+ __global gpu_register *tmp_1, __global gpu_register *tmp_2,
+ __global gpu_register *pks//, unsigned long count // invalid is not needed either
+ )
+{
+
+ size_t i = get_global_id(0);
+
+ int offs = 0; // the size of one number
+
+ int pk = 0; // das funktioniert so nicht – die globale id wird grösser, ohne dass noch weitere elemente in pks sind...
+
+
+// while (1) {
+// if (pks[pk] >= i)
+// break;
+// pk++;
+// }
+
+
+ int k = ceil(log2((float)exp[pk] + (float)1));
+
+ //printf((char __constant *)"%i\n", k);
+
+ barrier(CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE);
+
+ equal(tmp_2, x, 32);
+
+ int z = 0;
+
+ for (int j = k - 1; j >= 0; j--) {
+
+ square(tmp_1,tmp_2,32);
+ erase_above(tmp_1, 64);
+ //printf((char __constant *)"sqr\n");
+
+ if (testbit(exp[pk], j)) {
+ multiply(tmp_1,tmp_2,m,32);
+ erase_above(tmp_1, 64);
+ //printf((char __constant *)"mul\n");
+ }
+ }
+
+ equal(res, tmp_1, 32);
+
+}
diff --git a/xcode/verify.cl b/xcode/verify.cl
@@ -1240,7 +1240,8 @@ __kernel void several(__global DIGIT_T* x, __global const unsigned long *s_len,
int ndigits = max( max( n_len[pk] - (pk == 0 ? 0 : n_len[pk - 1]), mm_len[index] - (index == 0 ? 0 : mm_len[index - 1]) ), s_len[index] - (index == 0 ? 0 : s_len[index - 1]) );
int edigits = e_len[pk] - (pk == 0 ? 0 : e_len[pk - 1] );
-
+
+ //printf((char __constant *)"%i\n", ndigits);
// the result is copied in here, compare it to mm
DIGIT_T yout[MAX_ALLOC_SIZE * 2];
@@ -1260,9 +1261,9 @@ __kernel void several(__global DIGIT_T* x, __global const unsigned long *s_len,
window_mm = &mm[index == 0 ? 0 : (mm_len[index - 1])];
- __private DIGIT_T t1[MAX_ALLOC_SIZE *2];
- __private DIGIT_T t2[MAX_ALLOC_SIZE *2];
- __private DIGIT_T y[MAX_ALLOC_SIZE *2];
+ __private DIGIT_T t1[MAX_ALLOC_SIZE *2]; // obsolete?
+ __private DIGIT_T t2[MAX_ALLOC_SIZE *2]; // obsolete?
+ __private DIGIT_T y[MAX_ALLOC_SIZE *2]; // obsolete?
n = mpSizeof_g(window_e, edigits);
/* Catch e==0 => x^0=1 */
