commit f7d71405a6d636b67b0d34a32c9afdd2edfc8a9b
parent 6413fa3a63ce472796f0a534d223f815e25af678
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Thu, 16 Nov 2023 18:05:47 +0100
Prepare code for batch processing
At the moment, I am not moving forward with montgomery multiplication on the GPU – instead focusing on improving code that already works, and providing tests
Diffstat:
12 files changed, 574 insertions(+), 192 deletions(-)
diff --git a/.DS_Store b/.DS_Store
Binary files differ.
diff --git a/source/big-int-test.c b/source/big-int-test.c
@@ -787,30 +787,23 @@ void mpPrintHex(const char *prefix, const DIGIT_T *a, size_t len, const char *su
}
-int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, size_t ndigits)
-{ /* Computes y = x^e mod m */
+int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, size_t ndigits, size_t edigits)
+{
+ /* Computes y = x^e mod m */
/* "Classic" binary left-to-right method */
- /* [v2.2] removed const restriction on m[] to avoid using an extra alloc'd var
- (m is changed in-situ during the divide operation then restored) */
+
DIGIT_T mask;
size_t n;
size_t nn = ndigits * 2;
- /* Create some double-length temps */
-//#ifdef NO_ALLOCS
- DIGIT_T t1[MAX_FIXED_DIGITS * 2];
- DIGIT_T t2[MAX_FIXED_DIGITS * 2];
- DIGIT_T y[MAX_FIXED_DIGITS * 2];
+
+ DIGIT_T t1[nn];
+ DIGIT_T t2[nn];
+ DIGIT_T y[nn];
+
assert(ndigits <= MAX_FIXED_DIGITS);
-/*#else
- DIGIT_T *t1, *t2, *y;
- t1 = mpAlloc(nn);
- t2 = mpAlloc(nn);
- y = mpAlloc(nn);
-#endif
- */
assert(ndigits != 0);
- n = mpSizeof(e, ndigits);
+ n = mpSizeof(e, edigits);
/* Catch e==0 => x^0=1 */
if (0 == n)
{
@@ -847,9 +840,7 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
/* Move to next bit */
mpNEXTBITMASK(mask, n);
-
- // ctr++;
-
+
}
/* Return y */
diff --git a/source/big-int-test.h b/source/big-int-test.h
@@ -62,8 +62,6 @@ typedef uint16_t HALF_DIGIT_T;
#define mpMODSQUARETEMP(y,m,n,t1,t2) do{mpSquare(t1,y,n);mpDivide(t2,y,t1,n*2,m,n);}while(0)
/* Mult: y = (y * x) mod m */
#define mpMODMULTTEMP(y,x,m,n,t1,t2) do{mpMultiply(t1,x,y,n);mpDivide(t2,y,t1,n*2,m,n);}while(0)
-/* Mult: w = (y * x) mod m */
-#define mpMODMULTXYTEMP(w,y,x,m,n,t1,t2) do{mpMultiply(t1,x,y,(n));mpDivide(t2,w,t1,(n)*2,m,(n));}while(0)
#define mpNEXTBITMASK(mask, n) do{if(mask==1){mask=HIBITMASK;n--;}else{mask>>=1;}}while(0)
@@ -118,7 +116,7 @@ void mpFail(char *msg);
void mpPrintHex(const char *prefix, const DIGIT_T *a, size_t len, const char *suffix);
-int mpModExpO(DIGIT_T yout[], const DIGIT_T x[], const DIGIT_T e[], DIGIT_T m[], size_t ndigits);
+int mpModExpO(DIGIT_T yout[], const DIGIT_T x[], const DIGIT_T e[], DIGIT_T m[], size_t ndigits, size_t edigits);
static size_t conv_to_base(const DIGIT_T a[], size_t ndigits, char *s, size_t smax, int base);
diff --git a/source/lib-gpu-verify.c b/source/lib-gpu-verify.c
@@ -9,6 +9,8 @@
int main(int argc, char** argv)
{
+ //mont_prepare("07", "0A", "0D");
+
//opencl_tests();
rsa_tests();
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -18,21 +18,14 @@
//
//#include "RSA-Montgomery.h"
//
-#include "run-mmul.h"
+//#include "run-mmul.h"
#define NEED_LIBGCRYPT_VERSION "1.10.2"
#define DATA_SIZE (1024)
-int rsa_tests(void) {
-
-
-
- // MARK: UNSAFE init
+void setup_gcry(void) {
- // consider disabling optimizations, since they dont make for a fair comparison
-
- gcry_control (GCRYCTL_DISABLE_HWF, "intel-cpu", NULL);
/* Version check should be the very first call because it
makes sure that important subsystems are initialized.
@@ -48,6 +41,149 @@ int rsa_tests(void) {
/* Tell Libgcrypt that initialization has completed. */
gcry_control (GCRYCTL_INITIALIZATION_FINISHED, 0);
+}
+
+void generate_random_pairs(DIGIT_T *bases, size_t *b_len,
+ DIGIT_T *exponents, size_t *e_len,
+ DIGIT_T *moduli, size_t *m_len,
+ DIGIT_T *signatures, size_t *s_len,
+ const unsigned int n) {
+
+ int i;
+
+ int sz = 2048 / sizeof(char);
+
+ char *template = "(genkey(rsa(nbits 4:2048)))";
+ gcry_sexp_t parms;
+
+ gcry_sexp_new(&parms, template, strlen(template), 1);
+
+
+ for (i = 0; i < n; i++) {
+
+ gcry_sexp_t key;
+
+ gcry_pk_genkey(&key,parms);
+
+ char *val = "1234567890ABCDEF"; // MARK: try random values as well
+ gcry_mpi_t m_mpi = gcry_mpi_new((int)strlen(val) * 8);
+ size_t scanned = 0;
+
+ gcry_mpi_scan(&m_mpi, GCRYMPI_FMT_HEX, val, 0, &scanned);
+
+ gcry_sexp_t toSign;
+ size_t errOff = 0;
+ char *dataformat = "(data (flags raw) (value %m))";
+
+ gcry_sexp_build(&toSign,&errOff,dataformat,m_mpi);
+
+ gcry_sexp_t resSign;
+
+ gcry_pk_sign(&resSign, toSign, key);
+
+ // these must be freed manually
+ gcry_mpi_t n_mpi;
+ gcry_mpi_t e_mpi;
+
+ gcry_sexp_extract_param(key,NULL,"n e",&n_mpi, &e_mpi, NULL);
+
+ gcry_mpi_t sig_mpi;
+
+ gcry_sexp_extract_param(resSign,NULL,"s",&sig_mpi, NULL);
+
+ char *bb = malloc(sz);
+ char *ee = malloc(sz);
+ char *ss = malloc(sz);
+ char *mm = malloc(sz);
+
+ size_t nL = 0;
+
+ // check returns
+ gcry_mpi_print(GCRYMPI_FMT_HEX,(unsigned char *)mm,sz,&nL,n_mpi); // MARK: don't use magic numbers
+ gcry_mpi_print(GCRYMPI_FMT_HEX,(unsigned char *)ee,sz,&nL,e_mpi);
+ gcry_mpi_print(GCRYMPI_FMT_HEX,(unsigned char *)ss,sz,&nL,sig_mpi);
+ gcry_mpi_print(GCRYMPI_FMT_HEX,(unsigned char *)bb,sz,&nL,m_mpi);
+
+
+ DIGIT_T base [MAX_ALLOC_SIZE*2];
+ DIGIT_T exponent [MAX_ALLOC_SIZE*2];
+ DIGIT_T modulus [MAX_ALLOC_SIZE*2];
+ DIGIT_T signature [MAX_ALLOC_SIZE*2];
+
+ mpSetZero(base, MAX_ALLOC_SIZE*2);
+ mpSetZero(exponent, MAX_ALLOC_SIZE*2);
+ mpSetZero(modulus, MAX_ALLOC_SIZE*2);
+ mpSetZero(signature, MAX_ALLOC_SIZE*2);
+
+ mpConvFromHex(base, strlen(bb), bb);
+ mpConvFromHex(exponent, strlen(ee), ee);
+ mpConvFromHex(modulus, strlen(mm), mm);
+ mpConvFromHex(signature, strlen(ss), ss);
+
+ size_t max_len = max( max( mpSizeof(base, MAX_ALLOC_SIZE*2), mpSizeof(modulus, MAX_ALLOC_SIZE*2) ), mpSizeof(signature, MAX_ALLOC_SIZE*2) );
+
+ b_len[i] += max_len;
+ e_len[i] += mpSizeof(exponent, MAX_ALLOC_SIZE*2);
+ m_len[i] += max_len;
+ s_len[i] += max_len;
+
+ memcpy(&bases[i == 0 ? 0 : b_len[i - 1]], &base, ( b_len[i] - (i == 0 ? 0 : b_len[i - 1]) ) * sizeof(DIGIT_T));
+ memcpy(&exponents[i == 0 ? 0 : e_len[i - 1]], &exponent, ( e_len[i] - (i == 0 ? 0 : e_len[i - 1]) ) * sizeof(DIGIT_T));
+ memcpy(&moduli[i == 0 ? 0 : m_len[i - 1]], &modulus, ( m_len[i] - (i == 0 ? 0 : m_len[i - 1]) ) * sizeof(DIGIT_T));
+ memcpy(&signatures[i == 0 ? 0 : s_len[i - 1]], &signature, ( s_len[i] - (i == 0 ? 0 : s_len[i - 1]) ) * sizeof(DIGIT_T));
+
+ gcry_free(n_mpi);
+ gcry_free(e_mpi);
+ gcry_free(m_mpi);
+ gcry_free(sig_mpi);
+
+ free(bb);
+ free(ee);
+ free(ss);
+ free(mm);
+
+ }
+
+}
+
+int rsa_tests(void) {
+
+
+ setup_gcry();
+
+ int gen_n_pairs = 2; // MARK: it won't work with several yet, because in modexpO, they will read over the bounds of a number
+
+ DIGIT_T *q = malloc(2048);
+ DIGIT_T *r = malloc(2048);
+ DIGIT_T *s = malloc(2048);
+ DIGIT_T *t = malloc(2048);
+
+ size_t *u = malloc(gen_n_pairs * sizeof(size_t));
+ size_t *v = malloc(gen_n_pairs * sizeof(size_t));
+ size_t *w = malloc(gen_n_pairs * sizeof(size_t));
+ size_t *x = malloc(gen_n_pairs * sizeof(size_t));
+
+ generate_random_pairs(q, u,
+ r, v,
+ s, w,
+ t, x, gen_n_pairs);
+
+ DIGIT_T *y = malloc(2048);
+
+ DIGIT_T *s_window = &s[0];
+ DIGIT_T *r_window = &r[0];
+ DIGIT_T *t_window = &t[0];
+
+ mpModExpO(y, t_window, r_window, s_window, x[0], v[0]);
+
+ size_t sz_y = x[0];
+
+ char comp[sz_y];
+
+ mpConvToHex(y, sz_y, comp, sz_y);
+
+ printf("%s",comp);
+
char *template = "(genkey(rsa(nbits 4:2048)))";
gcry_sexp_t parms;
@@ -94,7 +230,7 @@ int rsa_tests(void) {
gcry_mpi_t e_mpi;
gcry_mpi_t d_mpi;
- gcry_sexp_extract_param(key,NULL,"n e d",&n_mpi, &e_mpi, &d_mpi, NULL);
+ gcry_sexp_extract_param(key,NULL,"n e",&n_mpi, &e_mpi, NULL);
gcry_mpi_t sig_mpi;
@@ -104,13 +240,13 @@ int rsa_tests(void) {
// may be a lot shorter – these will contain the numbers in HEX string form – for use in my bigNum
unsigned char *n = malloc(2048);
unsigned char *e = malloc(2048);
- unsigned char *d = malloc(2048);
+ //unsigned char *d = malloc(2048);
size_t nL = 0;
// check returns
gcry_mpi_print(GCRYMPI_FMT_HEX,n,2048,&nL,n_mpi);
gcry_mpi_print(GCRYMPI_FMT_HEX,e,2048,&nL,e_mpi);
- gcry_mpi_print(GCRYMPI_FMT_HEX,d,2048,&nL,d_mpi);
+ // gcry_mpi_print(GCRYMPI_FMT_HEX,d,2048,&nL,d_mpi);
unsigned char *sgn = malloc(2048);
gcry_mpi_print(GCRYMPI_FMT_HEX,sgn,2048,&nL,sig_mpi);
@@ -127,7 +263,7 @@ int rsa_tests(void) {
start = clock();
- main_mmul();
+ // main_mmul();
end = clock();
cpu_time_used = ((double) (end - start)) / CLOCKS_PER_SEC;
@@ -276,7 +412,7 @@ int rsa_tests(void) {
// Create the compute kernel in the program we wish to run
//
- kernel = clCreateKernel(program, "single", &err);
+ kernel = clCreateKernel(program, "several", &err);
if (!kernel || err != CL_SUCCESS)
{
printf("Error: Failed to create compute kernel!\n");
@@ -524,7 +660,7 @@ int verify(unsigned char* sign, unsigned char* ee, unsigned char* nn, unsigned c
//mpModMult(res, e, d, N, max(sz_d,sz_n)); // that works :)
- mpModExpO(res, s, e, N, max(sz_s,sz_n));
+ mpModExpO(res, s, e, N, max(sz_s,sz_n), mpSizeof(e, MAX_ALLOC_SIZE*2));
size_t sz_res = mpSizeof(res, MAX_ALLOC_SIZE*2);
diff --git a/xcode/.DS_Store b/xcode/.DS_Store
Binary files differ.
diff --git a/xcode/lib-gpu-verify.xcodeproj/project.pbxproj b/xcode/lib-gpu-verify.xcodeproj/project.pbxproj
@@ -9,7 +9,6 @@
/* Begin PBXBuildFile section */
6A8A795D2A89357400116D7D /* rsa-kernel.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A8A795C2A89357400116D7D /* rsa-kernel.cl */; };
6A8A795F2A89672700116D7D /* verify.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A8A795E2A89672700116D7D /* verify.cl */; };
- 6A9F57132B02EC0F00BC1F26 /* montgomery.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A9F57122B02EC0F00BC1F26 /* montgomery.c */; };
6AD85E072AF71AD900662919 /* big-int-test.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF7487D2ADADF4500D58E08 /* big-int-test.c */; };
6AD85E0C2AFA510C00662919 /* openssl-test.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AD85E0B2AFA510C00662919 /* openssl-test.c */; };
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */; };
@@ -34,8 +33,6 @@
466E0F5F0C932E1A00ED01DB /* lib-gpu-verify */ = {isa = PBXFileReference; explicitFileType = "compiled.mach-o.executable"; includeInIndex = 0; path = "lib-gpu-verify"; sourceTree = BUILT_PRODUCTS_DIR; };
6A8A795C2A89357400116D7D /* rsa-kernel.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = "rsa-kernel.cl"; sourceTree = "<group>"; };
6A8A795E2A89672700116D7D /* verify.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = verify.cl; sourceTree = "<group>"; };
- 6A9F57112B02EC0F00BC1F26 /* montgomery.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = montgomery.h; sourceTree = "<group>"; };
- 6A9F57122B02EC0F00BC1F26 /* montgomery.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; path = montgomery.c; sourceTree = "<group>"; };
6AD85E0A2AFA510C00662919 /* openssl-test.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = "openssl-test.h"; path = "../source/openssl-test.h"; sourceTree = "<group>"; };
6AD85E0B2AFA510C00662919 /* openssl-test.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "openssl-test.c"; path = "../source/openssl-test.c"; 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>"; };
@@ -99,8 +96,6 @@
6AF7487D2ADADF4500D58E08 /* big-int-test.c */,
6AF7487F2ADADF4500D58E08 /* rsa-test.c */,
6AF748852ADADFAD00D58E08 /* opencl-test.c */,
- 6A9F57112B02EC0F00BC1F26 /* montgomery.h */,
- 6A9F57122B02EC0F00BC1F26 /* montgomery.c */,
);
name = Sources;
sourceTree = "<group>";
@@ -167,7 +162,6 @@
isa = PBXSourcesBuildPhase;
buildActionMask = 2147483647;
files = (
- 6A9F57132B02EC0F00BC1F26 /* montgomery.c in Sources */,
6AD85E0C2AFA510C00662919 /* openssl-test.c in Sources */,
6AD85E072AF71AD900662919 /* big-int-test.c in Sources */,
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */,
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
@@ -690,8 +690,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "342"
- endingLineNumber = "342"
+ startingLineNumber = "478"
+ endingLineNumber = "478"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -738,8 +738,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "356"
- endingLineNumber = "356"
+ startingLineNumber = "492"
+ endingLineNumber = "492"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -891,8 +891,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "539"
- endingLineNumber = "539"
+ startingLineNumber = "675"
+ endingLineNumber = "675"
landmarkName = "verify(sign, ee, nn, mm)"
landmarkType = "9">
</BreakpointContent>
@@ -939,9 +939,9 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "815"
- endingLineNumber = "815"
- landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ startingLineNumber = "808"
+ endingLineNumber = "808"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
@@ -955,9 +955,9 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "838"
- endingLineNumber = "838"
- landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ startingLineNumber = "831"
+ endingLineNumber = "831"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
<Locations>
<Location
@@ -1018,9 +1018,9 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "856"
- endingLineNumber = "856"
- landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ startingLineNumber = "847"
+ endingLineNumber = "847"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
<Locations>
<Location
@@ -1081,9 +1081,9 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "839"
- endingLineNumber = "839"
- landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ startingLineNumber = "832"
+ endingLineNumber = "832"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
@@ -1400,205 +1400,397 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "566"
- endingLineNumber = "566"
+ startingLineNumber = "702"
+ endingLineNumber = "702"
landmarkName = "unknown"
landmarkType = "0">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "6AF0F685-479D-4405-BACC-8368C49802BF"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/RSA-Montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "117"
+ endingLineNumber = "117"
+ landmarkName = "modExpLUT(x, e, eBits, m, mBits, r2m, out)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "985780EE-603E-4B6C-BF80-1BB11F65F6BA"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "58"
+ endingLineNumber = "58"
+ landmarkName = "mont_prepare(base, exponent, modulus)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "714B2C00-5AA0-419D-8983-A6D8DF8F77EE"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "38"
+ endingLineNumber = "38"
+ landmarkName = "mont_prepare(base, exponent, modulus)"
+ landmarkType = "9">
<Locations>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c34b8"
+ uuid = "714B2C00-5AA0-419D-8983-A6D8DF8F77EE - 4382d64135f421be"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "39"
+ endingLineNumber = "39"
+ offsetFromSymbolStart = "201">
+ </Location>
+ <Location
+ uuid = "714B2C00-5AA0-419D-8983-A6D8DF8F77EE - 4382d64135f421be"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/montgomery.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "39"
+ endingLineNumber = "39"
+ offsetFromSymbolStart = "206">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "E4BA3895-109B-4936-ADE7-D8A141D7FA55"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "68"
+ endingLineNumber = "68"
+ landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "20DDCB70-9665-44F5-ABC4-C2D9C1BE45B7"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "59"
+ endingLineNumber = "59"
+ landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ landmarkName = "rsa_tests()"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "rsa_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "622">
+ </Location>
+ <Location
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770cafb9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "rsa_tests"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "185"
+ endingLineNumber = "185"
+ offsetFromSymbolStart = "562">
+ </Location>
+ <Location
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770cafb9"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "539"
- endingLineNumber = "539"
- offsetFromSymbolStart = "98">
+ startingLineNumber = "185"
+ endingLineNumber = "185"
+ offsetFromSymbolStart = "582">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c34b8"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca8ff"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "539"
- endingLineNumber = "539"
- offsetFromSymbolStart = "84">
+ startingLineNumber = "191"
+ endingLineNumber = "191"
+ offsetFromSymbolStart = "684">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3459"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca8ff"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "540"
- endingLineNumber = "540"
- offsetFromSymbolStart = "102">
+ startingLineNumber = "191"
+ endingLineNumber = "191"
+ offsetFromSymbolStart = "681">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3459"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "540"
- endingLineNumber = "540"
- offsetFromSymbolStart = "98">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "681">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c341f"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "542"
- endingLineNumber = "542"
- offsetFromSymbolStart = "102">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "677">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3459"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "540"
- endingLineNumber = "540"
- offsetFromSymbolStart = "88">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "514">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3459"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "540"
- endingLineNumber = "540"
- offsetFromSymbolStart = "110">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "569">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3459"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "540"
- endingLineNumber = "540"
- offsetFromSymbolStart = "115">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "572">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c349b"
+ uuid = "BFF56279-A16A-4556-9919-058156F61FD5 - b0b9078e770ca85a"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "538"
- endingLineNumber = "538"
- offsetFromSymbolStart = "88">
+ startingLineNumber = "188"
+ endingLineNumber = "188"
+ offsetFromSymbolStart = "560">
</Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "570B7C5C-BA80-46F3-A54C-51A33A4B0A75"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "663"
+ endingLineNumber = "663"
+ landmarkName = "verify(sign, ee, nn, mm)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "177"
+ endingLineNumber = "177"
+ landmarkName = "rsa_tests()"
+ landmarkType = "9">
+ <Locations>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c34fa"
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D - b0b9078e770caf52"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "537"
- endingLineNumber = "537"
- offsetFromSymbolStart = "79">
+ startingLineNumber = "180"
+ endingLineNumber = "180"
+ offsetFromSymbolStart = "562">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c36ab"
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D - b0b9078e770caf52"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "554"
- endingLineNumber = "554"
- offsetFromSymbolStart = "191">
+ startingLineNumber = "180"
+ endingLineNumber = "180"
+ offsetFromSymbolStart = "353">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c368a"
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D - b0b9078e770caeb1"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "553"
- endingLineNumber = "553"
- offsetFromSymbolStart = "191">
+ startingLineNumber = "177"
+ endingLineNumber = "177"
+ offsetFromSymbolStart = "353">
</Location>
<Location
- uuid = "1C1C50D0-346F-426D-BD58-8F706A2DD395 - 1c095a72436c3192"
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D - b0b9078e770caeb1"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- symbolName = "verify_gmp"
+ symbolName = "rsa_tests"
moduleName = "lib-gpu-verify"
usesParentBreakpointCondition = "Yes"
urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "561"
- endingLineNumber = "561"
- offsetFromSymbolStart = "191">
+ startingLineNumber = "177"
+ endingLineNumber = "177"
+ offsetFromSymbolStart = "371">
</Location>
</Locations>
</BreakpointContent>
@@ -1606,16 +1798,112 @@
<BreakpointProxy
BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
<BreakpointContent
- uuid = "6AF0F685-479D-4405-BACC-8368C49802BF"
- shouldBeEnabled = "Yes"
+ uuid = "4D5FA34D-C545-4A34-BB8B-8A2BA213FF3B"
+ shouldBeEnabled = "No"
ignoreCount = "0"
continueAfterRunningActions = "No"
- filePath = "../source/RSA-Montgomery.c"
+ filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "117"
- endingLineNumber = "117"
- landmarkName = "modExpLUT(x, e, eBits, m, mBits, r2m, out)"
+ startingLineNumber = "108"
+ endingLineNumber = "108"
+ landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "B481E43C-61F8-4041-948A-4170903D293D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "648"
+ endingLineNumber = "648"
+ landmarkName = "verify(sign, ee, nn, mm)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "73F6BB7B-689A-4F98-A7F8-5693DAF77398"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "135"
+ endingLineNumber = "135"
+ landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "46E9552D-DB7D-44F0-8A9F-7973AA1C8D61"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "130"
+ endingLineNumber = "130"
+ landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "46E9552D-DB7D-44F0-8A9F-7973AA1C8D61 - 6f45f8d7a4a135dc"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "generate_random_pairs"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "131"
+ endingLineNumber = "131"
+ offsetFromSymbolStart = "1134">
+ </Location>
+ <Location
+ uuid = "46E9552D-DB7D-44F0-8A9F-7973AA1C8D61 - 6f45f8d7a4a135ff"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "generate_random_pairs"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "130"
+ endingLineNumber = "130"
+ offsetFromSymbolStart = "1123">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "AE6F2C60-36E2-4F5D-94EA-115E01CF5285"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "166"
+ endingLineNumber = "166"
+ landmarkName = "rsa_tests()"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
diff --git a/xcode/montgomery.c b/xcode/montgomery.c
@@ -1,9 +0,0 @@
-//
-// montgomery.c
-// lib-gpu-verify
-//
-// Created by Cedric Zwahlen on 14.11.2023.
-//
-
-#include "montgomery.h"
-
diff --git a/xcode/montgomery.h b/xcode/montgomery.h
@@ -1,15 +0,0 @@
-//
-// montgomery.h
-// lib-gpu-verify
-//
-// Created by Cedric Zwahlen on 14.11.2023.
-//
-
-#ifndef montgomery_h
-#define montgomery_h
-
-#include <stdio.h>
-
-
-
-#endif /* montgomery_h */
diff --git a/xcode/verify.cl b/xcode/verify.cl
@@ -1200,20 +1200,16 @@ void assert(bool precondition) {
}
-
-//int mpModExpO(__global DIGIT_T *yout, __global DIGIT_T *x, __global DIGIT_T *e, __global DIGIT_T *m, size_t ndigits)
-
-// some might be constants
-__kernel void single(__global DIGIT_T* x, const unsigned int s_len,
- __global DIGIT_T* e, const unsigned int e_len,
- __global DIGIT_T* m, const unsigned int n_len,
- __global DIGIT_T *mm, const unsigned int mm_len,
- //global DIGIT_T* comp, const unsigned int comp_len,
- const unsigned int ndigits,
- __global int8* valid
- //const unsigned int count
- )
-{
+__kernel void several(__global DIGIT_T* x, __global const unsigned int *s_len,
+ __global DIGIT_T* e, __global const unsigned int *e_len,
+ __global DIGIT_T* m, __global const unsigned int *n_len,
+ __global DIGIT_T *mm, __global const unsigned int *mm_len,
+ const unsigned int ndigits,
+ const unsigned int count,
+ __global int8* valid
+ ) {
+
+ int index = get_global_id(0);
// the result is copied in here, compare it to mm
DIGIT_T yout[MAX_FIXED_DIGITS * 2];
@@ -1222,15 +1218,20 @@ __kernel void single(__global DIGIT_T* x, const unsigned int s_len,
size_t n;
size_t nn = ndigits * 2;
/* Create some double-length temps */
-
+
+ __global DIGIT_T *window_x = &x[s_len[index]];
+ __global DIGIT_T *window_e = &e[e_len[index]];
+ __global DIGIT_T *window_m = &m[n_len[index]];
+ __global DIGIT_T *window_mm = &mm[mm_len[index]];
+
DIGIT_T t1[MAX_FIXED_DIGITS * 2];
DIGIT_T t2[MAX_FIXED_DIGITS * 2];
DIGIT_T y[MAX_FIXED_DIGITS * 2];
assert(ndigits <= MAX_FIXED_DIGITS);
-
+
assert(ndigits != 0);
-
- n = mpSizeof_g(e, ndigits);
+
+ n = mpSizeof_g(window_e, ndigits);
/* Catch e==0 => x^0=1 */
if (0 == n)
{
@@ -1240,45 +1241,41 @@ __kernel void single(__global DIGIT_T* x, const unsigned int s_len,
/* Find second-most significant bit in e */
for (mask = HIBITMASK; mask > 0; mask >>= 1)
{
- if (e[n-1] & mask)
+ if (window_e[n-1] & mask)
break;
}
mpNEXTBITMASK(mask, n);
-
+
/* Set y = x */
- mpSetEqual_lg(y, x, ndigits);
-
+ mpSetEqual_lg(y, window_x, ndigits);
+
/* For bit j = k-2 downto 0 */
while (n) // I think it just goes the bit length of e
{
/* Square y = y * y mod n */
- mpMODSQUARETEMP(y, m, ndigits, t1, t2);
+ mpMODSQUARETEMP(y, window_mm, ndigits, t1, t2);
if (e[n-1] & mask)
{ /* if e(j) == 1 then multiply
- y = y * x mod n */
- mpMODMULTTEMP(y, x, m, ndigits, t1, t2);
-
+ y = y * x mod n */
+ mpMODMULTTEMP(y, window_x, window_m, ndigits, t1, t2);
+
}
/* Move to next bit */
- mpNEXTBITMASK(mask, n);
+ mpNEXTBITMASK(mask, n);
}
-
mpSetEqual(yout, y, ndigits);
+ int len = ( mm_len[index] - (index == 0 ? 0 : mm_len[index]) );
// equal
- *valid = mpCompare_lg(yout,mm,mm_len) == 0 ? 1 : 0;
+ *valid = mpCompare_lg(yout,window_mm,len) == 0 ? 1 : 0;
-
done:
mpDESTROY(t1, nn);
mpDESTROY(t2, nn);
mpDESTROY(y, ndigits);
-
-
-
}
