commit 421a712d4743971c39d81c71edd9beadb56f3776
parent e7b32caad0c0abf324ccd1469062046d7b2f2a3d
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Sat, 25 Nov 2023 18:06:21 +0100
Kernel compiles
but does not properly work still
Diffstat:
10 files changed, 1800 insertions(+), 122 deletions(-)
diff --git a/.DS_Store b/.DS_Store
Binary files differ.
diff --git a/source/gmp.c b/source/gmp.c
@@ -986,6 +986,8 @@ mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
mp_limb_t r;
mp_ptr tp = NULL;
mp_size_t tn = 0;
+
+
if (inv->shift > 0)
{
@@ -994,9 +996,10 @@ mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
if (!tp)
{
-
- //tn = nn;
- //tp = gmp_alloc_limbs (tn);
+ printf("%ld\n",nn);
+ tn = nn;
+ // tp = tn->_mp_d;
+ // tp = gmp_alloc_limbs (tn);
}
r = mpn_lshift (tp, np, nn, inv->shift);
np = tp;
@@ -4337,6 +4340,20 @@ 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)
{
diff --git a/source/lib-gpu-verify.c b/source/lib-gpu-verify.c
@@ -15,23 +15,23 @@ int main(int argc, char** argv)
//mont_go(res,"13", "0F", "C7",16);
- mont_go(res, "5BD6158BDE0AC0655B6FCEA57994011D18B6B3C9E5FF75C45FC1E5EC2C1F26D6AB8547A17C0BC15D40F4346CFE74CF4EB417E6850D45C3B49E9389DAF400BC5E5B3F5D8E1E45A23DD042A87E82703209F9EA9808A002FEC00C96A5F0D9B7673B4B0A224438D81C0A9CEAD0DD22802B409230072768E73688D63EAB1C9BC242FAEDCFE0C8478B38254BAAC07AD6F82A27A0C3893FCB604BB57158F9125027AECC91D55B364B5C2BB9FE07FB6AB69F5A65112A2B7D5A805CA9B2C1CB75D315DE345BA68100DD5E46FA3BA54B614C298E60EBAF95CEC738DA2513736ECE051D153CECAC29F4A432A5FEB287E2A1B8C4640C58FF9E9E7DB6889E4865D1F1C8CF4E47", "010001", "00BB5175E55C2F1BBAE52B0C1225F43385FF54B3BFEA88B42B21044328815B8742E303C843ABE76D147861AE92D563592EFD748BF2E5BE4D76793FB32FCF6B38F755D408D114C9DF89B3FAA77EDF0C9358AC3BC23C90CDAA8337927A3530DCF2AD6EFC023C96A7932F8A7935B9B3F5C84668B41FB39059A1B723A40D59A7B1BD03F56933D641409F2A49E614BBAA9F2573ED24899840585B73329A01071793332BA92A0C9033D7004B45FD01C3A850125FA2E4A40818F8E233B7B7595ABAB04B84AE88E4F7B516359EAB7C285F399A3EFF467113DDBDB17981F2F4F2DE405BA18863046570C1621AD9446CE8A3884893CEF50933CB60053B6862E2443CC8554121", 16);
+ //mont_go(res, "5BD6158BDE0AC0655B6FCEA57994011D18B6B3C9E5FF75C45FC1E5EC2C1F26D6AB8547A17C0BC15D40F4346CFE74CF4EB417E6850D45C3B49E9389DAF400BC5E5B3F5D8E1E45A23DD042A87E82703209F9EA9808A002FEC00C96A5F0D9B7673B4B0A224438D81C0A9CEAD0DD22802B409230072768E73688D63EAB1C9BC242FAEDCFE0C8478B38254BAAC07AD6F82A27A0C3893FCB604BB57158F9125027AECC91D55B364B5C2BB9FE07FB6AB69F5A65112A2B7D5A805CA9B2C1CB75D315DE345BA68100DD5E46FA3BA54B614C298E60EBAF95CEC738DA2513736ECE051D153CECAC29F4A432A5FEB287E2A1B8C4640C58FF9E9E7DB6889E4865D1F1C8CF4E47", "010001", "00BB5175E55C2F1BBAE52B0C1225F43385FF54B3BFEA88B42B21044328815B8742E303C843ABE76D147861AE92D563592EFD748BF2E5BE4D76793FB32FCF6B38F755D408D114C9DF89B3FAA77EDF0C9358AC3BC23C90CDAA8337927A3530DCF2AD6EFC023C96A7932F8A7935B9B3F5C84668B41FB39059A1B723A40D59A7B1BD03F56933D641409F2A49E614BBAA9F2573ED24899840585B73329A01071793332BA92A0C9033D7004B45FD01C3A850125FA2E4A40818F8E233B7B7595ABAB04B84AE88E4F7B516359EAB7C285F399A3EFF467113DDBDB17981F2F4F2DE405BA18863046570C1621AD9446CE8A3884893CEF50933CB60053B6862E2443CC8554121", 16);
//mont_go(res, "13", "05", "31",10);
- char str[2048];
- mpz_get_str(str, 16, res); // result is base 10!
+ //char str[2048];
+ //mpz_get_str(str, 16, res); // result is base 10!
- printf("%s\n",str);
+ //printf("%s\n",str);
- mont_go(res, "00956E3E7B09F7FECEF26CA44FFD69F19DC8DB6C3A29A707C2CDAD56994A58D6ACB8B275678D0D8670D3C716AC5C98398C8067943C7292F787F5451E8202F4C8BAEFA6CA787BC79B73A99CC4C85743EC7320E17195D560A380356A9D32AA81EF276A9DE8B9F6728647851AAD0090A458FB928BCE86884BD7CC7AC3CF226CE546E596135A948B820E1865D6A3395DF2BD5EB26FE5259B2B950CC61F887C0D5A81F77549D8F792D32552870358EC5B2B45552C35829D732CC1A08898FD2FFDFF5EBFE0BEE7D5702FCA240B377BFE7D2821E123F2A146725D01A5CF0A6C89FB7E73CA6F3B8640C44B0FA1A51B429BB3D4668495F20A25FB4185831C3B479C5041713C", "010001", "00BB5175E55C2F1BBAE52B0C1225F43385FF54B3BFEA88B42B21044328815B8742E303C843ABE76D147861AE92D563592EFD748BF2E5BE4D76793FB32FCF6B38F755D408D114C9DF89B3FAA77EDF0C9358AC3BC23C90CDAA8337927A3530DCF2AD6EFC023C96A7932F8A7935B9B3F5C84668B41FB39059A1B723A40D59A7B1BD03F56933D641409F2A49E614BBAA9F2573ED24899840585B73329A01071793332BA92A0C9033D7004B45FD01C3A850125FA2E4A40818F8E233B7B7595ABAB04B84AE88E4F7B516359EAB7C285F399A3EFF467113DDBDB17981F2F4F2DE405BA18863046570C1621AD9446CE8A3884893CEF50933CB60053B6862E2443CC8554121", 16);
+ //mont_go(res, "00956E3E7B09F7FECEF26CA44FFD69F19DC8DB6C3A29A707C2CDAD56994A58D6ACB8B275678D0D8670D3C716AC5C98398C8067943C7292F787F5451E8202F4C8BAEFA6CA787BC79B73A99CC4C85743EC7320E17195D560A380356A9D32AA81EF276A9DE8B9F6728647851AAD0090A458FB928BCE86884BD7CC7AC3CF226CE546E596135A948B820E1865D6A3395DF2BD5EB26FE5259B2B950CC61F887C0D5A81F77549D8F792D32552870358EC5B2B45552C35829D732CC1A08898FD2FFDFF5EBFE0BEE7D5702FCA240B377BFE7D2821E123F2A146725D01A5CF0A6C89FB7E73CA6F3B8640C44B0FA1A51B429BB3D4668495F20A25FB4185831C3B479C5041713C", "010001", "00BB5175E55C2F1BBAE52B0C1225F43385FF54B3BFEA88B42B21044328815B8742E303C843ABE76D147861AE92D563592EFD748BF2E5BE4D76793FB32FCF6B38F755D408D114C9DF89B3FAA77EDF0C9358AC3BC23C90CDAA8337927A3530DCF2AD6EFC023C96A7932F8A7935B9B3F5C84668B41FB39059A1B723A40D59A7B1BD03F56933D641409F2A49E614BBAA9F2573ED24899840585B73329A01071793332BA92A0C9033D7004B45FD01C3A850125FA2E4A40818F8E233B7B7595ABAB04B84AE88E4F7B516359EAB7C285F399A3EFF467113DDBDB17981F2F4F2DE405BA18863046570C1621AD9446CE8A3884893CEF50933CB60053B6862E2443CC8554121", 16);
//mont_go(res, "13", "05", "31",10);
- mpz_get_str(str, 16, res); // result is base 10!
+ //mpz_get_str(str, 16, res); // result is base 10!
- printf("%s\n",str);
+ //printf("%s\n",str);
//opencl_tests();
@@ -39,6 +39,10 @@ int main(int argc, char** argv)
// montgomery_test();
+ mont_rsa_tests();
+
+
+
return 0;
}
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -19,6 +19,8 @@
#include "ctype.h"
#include "time.h"
+//#include "gmp_GPU.h"
+
//
//#include "RSA-Montgomery.h"
//
@@ -674,6 +676,461 @@ int rsa_tests(void) {
return 0;
}
+
+// returns how many public keys were read – either 1 or n
+int mont_pairs_from_files(char *bases,
+ char *exponents,
+ char *moduli,
+ char *signatures,
+ unsigned int *n) {
+
+ FILE *pkfile;
+ FILE *msfile;
+
+ pkfile = fopen("lib-gpu-generate/publickey.txt", "r");
+ msfile = fopen("lib-gpu-generate/msgsig.txt", "r");
+
+ if (pkfile == NULL || msfile == NULL) {
+ printf("Auxiliary files not found.");
+ abort();
+ }
+
+
+ int i = 0;
+
+ while (1) {
+
+ char n_buf[2048];
+ char e_buf[2048];
+
+ if (fscanf(pkfile, "%s %s ", n_buf,e_buf) == -1)
+ break;
+
+ unsigned long n_buf_len = strlen(n_buf);
+ unsigned long e_buf_len = strlen(e_buf);
+
+ memcpy(moduli, n_buf, n_buf_len);
+ memcpy(exponents, e_buf, e_buf_len);
+
+ i++;
+
+ break; // testing with just one
+ }
+
+ int j = 0;
+
+ while (1) {
+
+ char m_buf[2048]; // temp storage, large enough
+ char s_buf[2048];
+
+ if (fscanf(msfile, "%s %s ", m_buf,s_buf) == -1)
+ break;
+
+ unsigned long m_buf_len = strlen(m_buf);
+ unsigned long s_buf_len = strlen(s_buf);
+
+ memcpy(bases, m_buf, m_buf_len);
+ memcpy(signatures, s_buf, s_buf_len);
+
+ j++;
+
+ break; // testing with just one
+
+ }
+
+ fclose(pkfile);
+ fclose(msfile);
+
+ *n = j;
+
+ return i;
+
+}
+
+
+int mont_verify_pairs_with_opencl(char *bases,
+ char *exponents,
+ char *moduli,
+ char *signatures,
+ const unsigned int n,
+ const unsigned int pks,
+ unsigned long *result) {
+
+ int err; // error code returned from api calls
+
+ size_t global; // global domain size for our calculation
+ size_t local; // local domain size for our calculation
+
+ cl_device_id device_id; // compute device id
+ cl_context context; // compute context
+ cl_command_queue commands; // compute command queue
+ cl_program program; // compute program
+ cl_kernel kernel; // compute kernel
+
+ cl_uint max_platforms = 4;
+ cl_platform_id platforms[max_platforms];
+ cl_uint num_platforms;
+ cl_int rplat;
+
+ rplat = clGetPlatformIDs (max_platforms, platforms, &num_platforms);
+ if (rplat != CL_SUCCESS)
+ {
+ printf("Error: Failed to lookup platforms! (%d)\n", rplat);
+ return EXIT_FAILURE;
+ }
+ printf("Found %u platforms\n", (unsigned int) num_platforms);
+ if (0 == num_platforms)
+ {
+ return EXIT_FAILURE;
+ }
+ for (unsigned int i=0;i<num_platforms;i++)
+ {
+ char buf[128];
+ size_t rbuf;
+ static struct {
+ cl_platform_info cpi;
+ const char *name;
+ } param[] = {
+ { CL_PLATFORM_PROFILE, "profile" },
+ { CL_PLATFORM_VENDOR, "vendor" },
+ { CL_PLATFORM_NAME, "name" },
+ { CL_PLATFORM_EXTENSIONS, "extensions" },
+ { 0, NULL }
+ };
+
+ for (unsigned int j=0;NULL != param[j].name;j++)
+ {
+ err = clGetPlatformInfo (platforms[i], param[j].cpi, sizeof (buf), buf, &rbuf);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to get platform info for %s! (%d)\n", param[j].name, err);
+ }
+ else
+ {
+ printf ("#%u %s %.*s\n", i, param[j].name, (int) rbuf, buf);
+ }
+ }
+ }
+
+ // Connect to a compute device
+ //
+ //int gpu = 1;
+ err = clGetDeviceIDs(platforms[0],
+ CL_DEVICE_TYPE_ALL, 1, &device_id, NULL);
+ // gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to create a device group! (%d)\n", err);
+ return EXIT_FAILURE;
+ }
+
+ size_t retSize_1 = 0;
+ clGetDeviceInfo(device_id, CL_DRIVER_VERSION, 0, NULL, &retSize_1);
+ char driver_version[retSize_1];
+ clGetDeviceInfo(device_id, CL_DRIVER_VERSION, retSize_1, &driver_version, &retSize_1);
+
+ //printf("driver version: %s\n", driver_version);
+
+
+ size_t retSize_2 = sizeof(cl_uint);
+ cl_uint address_bits = 0;
+ clGetDeviceInfo(device_id, CL_DEVICE_ADDRESS_BITS, 0, NULL, &retSize_2);
+ clGetDeviceInfo(device_id, CL_DEVICE_ADDRESS_BITS, retSize_2, &address_bits, &retSize_2);
+
+ if (address_bits == 32) {
+ printf("Kernel is only designed to run on 64-bit GPUs.");
+ abort();
+ }
+
+ // Create a compute context
+ //
+ context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
+ if (!context)
+ {
+ printf("Error: Failed to create a compute context!\n");
+ return EXIT_FAILURE;
+ }
+
+ // Create a command commands
+ //
+ commands = clCreateCommandQueue(context, device_id, 0, &err);
+ if (!commands)
+ {
+ printf("Error: Failed to create a command commands!\n");
+ return EXIT_FAILURE;
+ }
+
+ // get the kernel from a file instead of a constant
+
+ FILE *fp = fopen("./montgomery.cl", "r");
+ if (NULL == fp)
+ {
+ printf("Error: could not find 'montgomery.cl'\n");
+ return EXIT_FAILURE;
+ }
+ fseek(fp, 0L, SEEK_END);
+ size_t sz = ftell(fp);
+ rewind(fp);
+
+ char *kernelBuf = malloc(sz);
+ fread(kernelBuf, sizeof(char), sz, fp);
+ fclose(fp);
+
+ // Create the compute program from the source buffer
+ //
+ //program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
+ program = clCreateProgramWithSource(context, 1, (const char **) & kernelBuf, &sz, &err);
+ if (!program)
+ {
+ printf("Error: Failed to create compute program!\n");
+ return EXIT_FAILURE;
+ }
+
+ // Build the program executable
+ //
+ err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
+ if (err != CL_SUCCESS)
+ {
+ size_t len;
+ char buffer[352323];
+
+ printf("Error: Failed to build program executable!\n");
+ clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+ printf("%s\n", buffer);
+ exit(1);
+ }
+
+ size_t len;
+ char buffer[3523];
+
+ clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
+ printf("%s\n", buffer);
+
+
+
+ // Create the compute kernel in the program we wish to run
+ //
+ kernel = clCreateKernel(program, "montgomery", &err);
+ if (!kernel || err != CL_SUCCESS)
+ {
+ printf("Error: Failed to create compute kernel!\n");
+ exit(1);
+ }
+
+ // Create the input and output arrays in device memory for our calculation
+
+ 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 valid; // needs to be a buffer because it goes out
+
+ unsigned long signature_is_valid = 0;
+
+ size_t moduli_len = strlen(moduli) + 1;
+ size_t exponents_len = strlen(exponents) + 1;
+ size_t signatures_len = strlen(signatures) + 1;
+ size_t bases_len = strlen(bases) + 1;
+
+
+
+
+ if (pks == 1) {
+ mod_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, moduli_len, NULL, NULL);
+ exp_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, exponents_len, NULL, NULL);
+ } else {
+
+ }
+
+ sig_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, signatures_len , NULL, NULL);
+ comp_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, bases_len, NULL, NULL); // the base, to compare whether we get the same signature
+//
+// if (pks == 1) {
+// mod_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long), NULL, NULL);
+// exp_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long), NULL, NULL);
+// } else {
+//
+// }
+//
+// sig_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long) * n, NULL, NULL);
+// comp_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long) * n, NULL, NULL);
+
+ valid = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(unsigned long) ,NULL, NULL);
+
+ global = 65536; // MARK: ???
+
+ if (!sig_mem || !exp_mem || !mod_mem || !comp_mem || !valid )
+ {
+ printf("Error: Failed to allocate device memory!\n");
+ exit(1);
+ }
+
+ // Write our data set into the input array in device memory
+ //
+ err = clEnqueueWriteBuffer(commands, sig_mem, CL_TRUE, 0, signatures_len, signatures, 0, NULL, NULL);
+// err |= clEnqueueWriteBuffer(commands, sig_len, CL_TRUE, 0,sizeof(unsigned long) * n, s_len, 0, NULL, NULL);
+
+ if (pks == 1) {
+ err |= clEnqueueWriteBuffer(commands, exp_mem, CL_TRUE, 0, exponents_len, exponents, 0, NULL, NULL);
+// err |= clEnqueueWriteBuffer(commands, exp_len, CL_TRUE, 0,sizeof(unsigned long), e_len, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, mod_mem, CL_TRUE, 0, moduli_len, moduli, 0, NULL, NULL);
+// err |= clEnqueueWriteBuffer(commands, mod_len, CL_TRUE, 0,sizeof(unsigned long), m_len, 0, NULL, NULL);
+ } else {
+
+ }
+
+ err |= clEnqueueWriteBuffer(commands, comp_mem, CL_TRUE, 0, bases_len, bases, 0, NULL, NULL);
+// err |= clEnqueueWriteBuffer(commands, comp_len, CL_TRUE, 0,sizeof(unsigned long) * n, b_len, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, valid, CL_TRUE, 0, sizeof(unsigned long), &signature_is_valid, 0, NULL, NULL);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to write to source array!\n");
+ exit(1);
+ }
+
+ // Set the arguments to our compute kernel
+ //
+ err = 0;
+ err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &sig_mem);
+// err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &sig_len);
+ err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &exp_mem);
+// err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &exp_len);
+ err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &mod_mem);
+// err |= clSetKernelArg(kernel, 5, sizeof(cl_mem), &mod_len);
+ err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &comp_mem);
+// err |= clSetKernelArg(kernel, 7, sizeof(cl_mem), &comp_len);
+ err |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &valid);
+ //err |= clSetKernelArg(kernel, 5, sizeof(unsigned int), &n);
+ //err |= clSetKernelArg(kernel, 6, sizeof(unsigned int), &pks);
+
+ //err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);
+ if (err != CL_SUCCESS)
+ {
+ printf("RSA-Error: Failed to set kernel arguments! %d\n", err);
+ exit(1);
+ }
+
+ // Get the maximum work group size for executing the kernel on the device
+ //
+ err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to retrieve kernel work group info! %d\n", err);
+ exit(1);
+ }
+
+ // Execute the kernel over the entire range of our 1d input data set
+ // using the maximum number of work group items for this device
+ //
+
+ err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
+ if (err)
+ {
+ printf("Error: Failed to execute kernel!\n");
+ return EXIT_FAILURE;
+ }
+
+ printf("KERNEL IS EXECUTING...\n");
+
+ // Wait for the command commands to get serviced before reading back results
+ //
+ clFinish(commands);
+
+ // Read back the results from the device to verify the output
+ //
+ //err = clEnqueueReadBuffer( commands, res_mem, CL_TRUE, 0, res_len, res_buf, 0, NULL, NULL );
+ err = clEnqueueReadBuffer( commands, valid, CL_TRUE, 0, sizeof(unsigned long), &signature_is_valid, 0, NULL, NULL );
+ if (err != CL_SUCCESS)
+ {
+ printf("Error: Failed to read output array! %d\n", err);
+ exit(1);
+ }
+
+ *result = signature_is_valid;
+
+ // Shutdown and cleanup
+ //
+ clReleaseMemObject(comp_mem);
+ clReleaseMemObject(exp_mem);
+ clReleaseMemObject(mod_mem);
+ clReleaseMemObject(sig_mem);
+
+// clReleaseMemObject(comp_len);
+// clReleaseMemObject(exp_len);
+// clReleaseMemObject(mod_len);
+// clReleaseMemObject(sig_len);
+
+ clReleaseProgram(program);
+ clReleaseKernel(kernel);
+ clReleaseCommandQueue(commands);
+ clReleaseContext(context);
+
+ return 0;
+
+}
+
+
+int mont_rsa_tests(void) {
+
+ setup_gcry();
+
+ unsigned int pairs = 1;
+
+
+ int str_sz = (2048);
+
+
+ char *b = malloc(str_sz);
+ char *e = malloc(str_sz);
+ char *m = malloc(str_sz);
+ char *s = malloc(str_sz);
+
+ unsigned int pks = pairs;
+
+ pks = mont_pairs_from_files(b, e, m, s, &pairs);
+
+ printf("--");
+
+
+ unsigned long result = 0;
+
+ struct timespec t1, t2;
+
+ clock_gettime(CLOCK_REALTIME, &t1);
+
+
+ // montgomery(b, e, m, s, &result);
+
+
+ mont_verify_pairs_with_opencl(
+ b,e,m,s,
+ pairs, pks, &result);
+
+ clock_gettime(CLOCK_REALTIME, &t2);
+
+ printf("VERIFICATION RESULT: %lu\n",result);
+
+ printf("\nGPU verification took %ld ms\n", (t2.tv_nsec - t1.tv_nsec) / 1000000);
+
+
+
+ free(b);
+ free(e);
+ free(m);
+ free(s);
+
+
+}
+
+
/*
static void show_sexp(const char *prefix, gcry_sexp_t a) {
char *buf;
diff --git a/source/rsa-test.h b/source/rsa-test.h
@@ -14,7 +14,9 @@
int rsa_tests(void);
-static void show_sexp(const char *prefix, gcry_sexp_t a);
+int mont_rsa_tests(void);
+
+//static void show_sexp(const char *prefix, gcry_sexp_t a);
void bigNum_tests(unsigned char* n, unsigned char* e, unsigned char* d);
@@ -24,6 +26,6 @@ int verify(unsigned char* sign, unsigned char* ee, unsigned char* nn, unsigned c
//int verify_gmp(unsigned char* sign, unsigned char* ee, unsigned char* nn, unsigned char* mm);
-static void upper(unsigned char* str);
+//static void upper(unsigned char* str);
#endif /* rsa_test_h */
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
@@ -8,9 +8,8 @@
/* Begin PBXBuildFile section */
6A36F8892B0F938E00AB772D /* montgomery.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A36F8882B0F938E00AB772D /* montgomery.cl */; };
- 6A7914CF2B0CF320001EDCC1 /* gmp.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A7914CB2B0CF320001EDCC1 /* gmp.c */; };
- 6A7914D02B0CF320001EDCC1 /* montgomery.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A7914CD2B0CF320001EDCC1 /* montgomery.c */; };
6A8A795F2A89672700116D7D /* verify.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6A8A795E2A89672700116D7D /* verify.cl */; };
+ 6A99B0672B125F430004E4B7 /* gmp.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A7914CB2B0CF320001EDCC1 /* gmp.c */; };
6AA38E5B2B0A97FC00E85243 /* main.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AA38E5A2B0A97FC00E85243 /* main.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 */; };
@@ -228,11 +227,10 @@
files = (
6AD85E0C2AFA510C00662919 /* openssl-test.c in Sources */,
6AD85E072AF71AD900662919 /* big-int-test.c in Sources */,
+ 6A99B0672B125F430004E4B7 /* gmp.c in Sources */,
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */,
6A8A795F2A89672700116D7D /* verify.cl in Sources */,
- 6A7914CF2B0CF320001EDCC1 /* gmp.c in Sources */,
6AF748832ADADF4500D58E08 /* rsa-test.c in Sources */,
- 6A7914D02B0CF320001EDCC1 /* montgomery.c in Sources */,
6AF748862ADADFAD00D58E08 /* opencl-test.c in Sources */,
6A36F8892B0F938E00AB772D /* montgomery.cl 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
@@ -1262,8 +1262,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "71"
- endingLineNumber = "71"
+ startingLineNumber = "73"
+ endingLineNumber = "73"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1278,8 +1278,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "606"
- endingLineNumber = "606"
+ startingLineNumber = "608"
+ endingLineNumber = "608"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -1461,8 +1461,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "593"
- endingLineNumber = "593"
+ startingLineNumber = "595"
+ endingLineNumber = "595"
landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
landmarkType = "9">
<Locations>
@@ -1554,8 +1554,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "111"
- endingLineNumber = "111"
+ startingLineNumber = "113"
+ endingLineNumber = "113"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1570,8 +1570,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "138"
- endingLineNumber = "138"
+ startingLineNumber = "140"
+ endingLineNumber = "140"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1586,8 +1586,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "133"
- endingLineNumber = "133"
+ startingLineNumber = "135"
+ endingLineNumber = "135"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
<Locations>
@@ -1757,8 +1757,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "252"
- endingLineNumber = "252"
+ startingLineNumber = "254"
+ endingLineNumber = "254"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1773,8 +1773,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "210"
- endingLineNumber = "210"
+ startingLineNumber = "212"
+ endingLineNumber = "212"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
<Locations>
@@ -1866,8 +1866,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "206"
- endingLineNumber = "206"
+ startingLineNumber = "208"
+ endingLineNumber = "208"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1882,8 +1882,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "245"
- endingLineNumber = "245"
+ startingLineNumber = "247"
+ endingLineNumber = "247"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1898,8 +1898,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "197"
- endingLineNumber = "197"
+ startingLineNumber = "199"
+ endingLineNumber = "199"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1914,8 +1914,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "563"
- endingLineNumber = "563"
+ startingLineNumber = "565"
+ endingLineNumber = "565"
landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
landmarkType = "9">
</BreakpointContent>
@@ -1930,8 +1930,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "653"
- endingLineNumber = "653"
+ startingLineNumber = "655"
+ endingLineNumber = "655"
landmarkName = "rsa_tests()"
landmarkType = "9">
</BreakpointContent>
@@ -1946,8 +1946,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "572"
- endingLineNumber = "572"
+ startingLineNumber = "574"
+ endingLineNumber = "574"
landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
landmarkType = "9">
</BreakpointContent>
@@ -1962,8 +1962,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "638"
- endingLineNumber = "638"
+ startingLineNumber = "640"
+ endingLineNumber = "640"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -2010,8 +2010,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "306"
- endingLineNumber = "306"
+ startingLineNumber = "308"
+ endingLineNumber = "308"
landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
landmarkType = "9">
</BreakpointContent>
@@ -2026,8 +2026,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "239"
- endingLineNumber = "239"
+ startingLineNumber = "241"
+ endingLineNumber = "241"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -2042,8 +2042,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "247"
- endingLineNumber = "247"
+ startingLineNumber = "249"
+ endingLineNumber = "249"
landmarkName = "pairs_from_files(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -2872,8 +2872,8 @@
filePath = "../source/gmp.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "1001"
- endingLineNumber = "1001"
+ startingLineNumber = "1004"
+ endingLineNumber = "1004"
landmarkName = "mpn_div_qr_1_preinv(qp, np, nn, inv)"
landmarkType = "9">
</BreakpointContent>
@@ -2881,48 +2881,979 @@
<BreakpointProxy
BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
<BreakpointContent
- uuid = "4A0A8B0C-87AE-40A2-BD5A-D53B0C6B1F51"
+ uuid = "B7D2B136-D2CC-494F-B46C-D6FF5545EE39"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "4219"
+ endingLineNumber = "4219"
+ landmarkName = "mpz_sizeinbase(u, base)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "B8675182-6B29-4414-80E9-6A957CF4BF1A"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "438"
+ endingLineNumber = "438"
+ landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "A4D3CD3E-C63D-4684-A93A-EE22635022EE"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "701"
+ endingLineNumber = "701"
+ landmarkName = "mont_pairs_from_files(bases, exponents, moduli, signatures, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "4184BD65-645D-4022-94C1-79216BEF6823"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "896"
+ endingLineNumber = "896"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "6A825320-5EF5-46A3-88EB-CBD1ECAB327F"
shouldBeEnabled = "Yes"
ignoreCount = "0"
continueAfterRunningActions = "No"
- filePath = "montgomery.cl"
+ filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "1664"
- endingLineNumber = "1664"
- landmarkName = "mpz_sizeinbase()"
+ startingLineNumber = "1050"
+ endingLineNumber = "1050"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "6A825320-5EF5-46A3-88EB-CBD1ECAB327F - 2905f85b306cd1a2"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1035"
+ endingLineNumber = "1035"
+ offsetFromSymbolStart = "3230">
+ </Location>
+ <Location
+ uuid = "6A825320-5EF5-46A3-88EB-CBD1ECAB327F - 2905f85b306cd2ba"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1043"
+ endingLineNumber = "1043"
+ offsetFromSymbolStart = "3294">
+ </Location>
+ <Location
+ uuid = "6A825320-5EF5-46A3-88EB-CBD1ECAB327F - 2905f85b306cd3df"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1048"
+ endingLineNumber = "1048"
+ offsetFromSymbolStart = "3154">
+ </Location>
+ <Location
+ uuid = "6A825320-5EF5-46A3-88EB-CBD1ECAB327F - 2905f85b306cd391"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1050"
+ endingLineNumber = "1050"
+ offsetFromSymbolStart = "3154">
+ </Location>
+ </Locations>
</BreakpointContent>
</BreakpointProxy>
<BreakpointProxy
BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
<BreakpointContent
- uuid = "B7D2B136-D2CC-494F-B46C-D6FF5545EE39"
+ uuid = "A0206612-E01F-4D5F-BA1C-09E3D6085CD0"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1045"
+ endingLineNumber = "1045"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "A0206612-E01F-4D5F-BA1C-09E3D6085CD0 - 2905f85b306cd00d"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1030"
+ endingLineNumber = "1030"
+ offsetFromSymbolStart = "3218">
+ </Location>
+ <Location
+ uuid = "A0206612-E01F-4D5F-BA1C-09E3D6085CD0 - 2905f85b306cd105"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1038"
+ endingLineNumber = "1038"
+ offsetFromSymbolStart = "3282">
+ </Location>
+ <Location
+ uuid = "A0206612-E01F-4D5F-BA1C-09E3D6085CD0 - 2905f85b306cd2ba"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1043"
+ endingLineNumber = "1043"
+ offsetFromSymbolStart = "3142">
+ </Location>
+ <Location
+ uuid = "A0206612-E01F-4D5F-BA1C-09E3D6085CD0 - 2905f85b306cd27c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1045"
+ endingLineNumber = "1045"
+ offsetFromSymbolStart = "3142">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "8BB8F280-7EFB-48F8-A762-C497F0136785"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "908"
+ endingLineNumber = "908"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "F6E2F2AB-21A2-4DC2-8B9C-59FF6326A2DC"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "946"
+ endingLineNumber = "946"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "C8B7770D-BAC1-4203-AF94-A85245DE1081"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "941"
+ endingLineNumber = "941"
+ landmarkName = "mont_verify_pairs_with_opencl(bases, exponents, moduli, signatures, n, pks, result)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "C8B7770D-BAC1-4203-AF94-A85245DE1081 - 2905f85b306c25c2"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "939"
+ endingLineNumber = "939"
+ offsetFromSymbolStart = "1812">
+ </Location>
+ <Location
+ uuid = "C8B7770D-BAC1-4203-AF94-A85245DE1081 - 2905f85b306c2584"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_verify_pairs_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "941"
+ endingLineNumber = "941"
+ offsetFromSymbolStart = "1812">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "403AE82D-01F5-40B2-BBD9-A4420ED3D914"
shouldBeEnabled = "No"
ignoreCount = "0"
continueAfterRunningActions = "No"
filePath = "../source/gmp.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "4216"
- endingLineNumber = "4216"
- landmarkName = "mpz_sizeinbase(u, base)"
+ startingLineNumber = "1000"
+ endingLineNumber = "1000"
+ landmarkName = "mpn_div_qr_1_preinv(qp, np, nn, inv)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
<BreakpointProxy
BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
<BreakpointContent
- uuid = "B8675182-6B29-4414-80E9-6A957CF4BF1A"
+ uuid = "2123D5D5-C635-404E-91CB-F962716A9093"
shouldBeEnabled = "No"
ignoreCount = "0"
continueAfterRunningActions = "No"
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "436"
- endingLineNumber = "436"
- landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks, result)"
+ startingLineNumber = "1110"
+ endingLineNumber = "1110"
+ landmarkName = "mont_rsa_tests()"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "8ECD43E8-C440-47F2-92EE-5C0EFBCD0487"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "495"
+ endingLineNumber = "495"
+ landmarkName = "mpz_set(r, x)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "472223C6-07BA-494A-A0A7-058648BF776C"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "492"
+ endingLineNumber = "492"
+ landmarkName = "mpz_set(r, x)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "1E44FD05-54D7-4E86-9450-6A8BDE263CE7"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2802"
+ endingLineNumber = "2802"
+ landmarkName = "mont_mulmod(res, a, b, mod)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "D1044013-3029-4CCB-9700-A4ED809E7749"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2858"
+ endingLineNumber = "2858"
+ landmarkName = "montgomery(signature, exponent, modulus, base, valid)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "D1044013-3029-4CCB-9700-A4ED809E7749 - 941e65ef5d451165"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2852"
+ endingLineNumber = "2852"
+ offsetFromSymbolStart = "235">
+ </Location>
+ <Location
+ uuid = "D1044013-3029-4CCB-9700-A4ED809E7749 - 941e65ef5d4511a7"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2850"
+ endingLineNumber = "2850"
+ offsetFromSymbolStart = "235">
+ </Location>
+ <Location
+ uuid = "D1044013-3029-4CCB-9700-A4ED809E7749 - 941e65ef5d451104"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2853"
+ endingLineNumber = "2853"
+ offsetFromSymbolStart = "235">
+ </Location>
+ <Location
+ uuid = "D1044013-3029-4CCB-9700-A4ED809E7749 - 941e65ef5d4512fa"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2855"
+ endingLineNumber = "2855"
+ offsetFromSymbolStart = "235">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "78DCF329-4308-45D8-AC1F-CCBE9A16757A"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2903"
+ endingLineNumber = "2903"
+ landmarkName = "montgomery(signature, exponent, modulus, base, valid)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "4BE2D23D-FE66-40D8-AAC5-D21A82847405"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2901"
+ endingLineNumber = "2901"
+ landmarkName = "montgomery(signature, exponent, modulus, base, valid)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "D6011B18-999E-4186-A3BD-BCBFB4665557"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2669"
+ endingLineNumber = "2669"
+ landmarkName = "mont_prepare(b, e, m, r, r_1, ni, M, x)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "D6011B18-999E-4186-A3BD-BCBFB4665557 - 64518ca852e0645"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2659"
+ endingLineNumber = "2659"
+ offsetFromSymbolStart = "138">
+ </Location>
+ <Location
+ uuid = "D6011B18-999E-4186-A3BD-BCBFB4665557 - 64518ca852e07f9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2663"
+ endingLineNumber = "2663"
+ offsetFromSymbolStart = "180">
+ </Location>
+ <Location
+ uuid = "D6011B18-999E-4186-A3BD-BCBFB4665557 - 64518ca852e075c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mont_prepare"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2666"
+ endingLineNumber = "2666"
+ offsetFromSymbolStart = "180">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "48FEE708-B88E-43F9-AAF0-4D30F363C472"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "836"
+ endingLineNumber = "836"
+ landmarkName = "mpz_div_qr(q, r, n, d, mode)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "48FEE708-B88E-43F9-AAF0-4D30F363C472 - 50dbcca7228f7fc5"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_invert_3by2"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "827"
+ endingLineNumber = "827"
+ offsetFromSymbolStart = "931">
+ </Location>
+ <Location
+ uuid = "48FEE708-B88E-43F9-AAF0-4D30F363C472 - b538a7a206838259"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpz_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "834"
+ endingLineNumber = "834"
+ offsetFromSymbolStart = "66">
+ </Location>
+ <Location
+ uuid = "48FEE708-B88E-43F9-AAF0-4D30F363C472 - b538a7a20683823a"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpz_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "837"
+ endingLineNumber = "837"
+ offsetFromSymbolStart = "66">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "2EE70F0A-89B5-47E5-95F0-727F46783CEE"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "907"
+ endingLineNumber = "907"
+ landmarkName = "mpz_div_qr(q, r, n, d, mode)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "2EE70F0A-89B5-47E5-95F0-727F46783CEE - b538a7a206838a7e"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpz_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "897"
+ endingLineNumber = "897"
+ offsetFromSymbolStart = "766">
+ </Location>
+ <Location
+ uuid = "2EE70F0A-89B5-47E5-95F0-727F46783CEE - b538a7a2068389fa"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpz_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "901"
+ endingLineNumber = "901"
+ offsetFromSymbolStart = "800">
+ </Location>
+ <Location
+ uuid = "2EE70F0A-89B5-47E5-95F0-727F46783CEE - b538a7a206838967"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpz_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "904"
+ endingLineNumber = "904"
+ offsetFromSymbolStart = "800">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "98374BB9-2845-4DB0-825F-5112D3F5E50D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1033"
+ endingLineNumber = "1033"
+ landmarkName = "mpn_div_qr_invert(inv, dp, dn)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "98374BB9-2845-4DB0-825F-5112D3F5E50D - 392da6fc5c4f34b3"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr_2_invert"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1023"
+ endingLineNumber = "1023"
+ offsetFromSymbolStart = "282">
+ </Location>
+ <Location
+ uuid = "98374BB9-2845-4DB0-825F-5112D3F5E50D - 7554351b5843b57"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr_invert"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1030"
+ endingLineNumber = "1030"
+ offsetFromSymbolStart = "20">
+ </Location>
+ <Location
+ uuid = "98374BB9-2845-4DB0-825F-5112D3F5E50D - 7554351b5843af4"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr_invert"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "1033"
+ endingLineNumber = "1033"
+ offsetFromSymbolStart = "20">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "8DBF09D6-C8CB-4D1F-B149-A4BCF77FB839"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "948"
+ endingLineNumber = "948"
+ landmarkName = "mpn_div_qr(qp, np, nn, dp, dn)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "8DBF09D6-C8CB-4D1F-B149-A4BCF77FB839 - 7090678fa8dcffd6"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "949"
+ endingLineNumber = "949"
+ offsetFromSymbolStart = "60">
+ </Location>
+ <Location
+ uuid = "8DBF09D6-C8CB-4D1F-B149-A4BCF77FB839 - 7090678fa8dcffd6"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "949"
+ endingLineNumber = "949"
+ offsetFromSymbolStart = "66">
+ </Location>
+ <Location
+ uuid = "8DBF09D6-C8CB-4D1F-B149-A4BCF77FB839 - 7090678fa8dcff6b"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "952"
+ endingLineNumber = "952"
+ offsetFromSymbolStart = "66">
+ </Location>
+ <Location
+ uuid = "8DBF09D6-C8CB-4D1F-B149-A4BCF77FB839 - 7090678fa8dcff6b"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpn_div_qr"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "952"
+ endingLineNumber = "952"
+ offsetFromSymbolStart = "60">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "7B64655D-0DD5-4A2E-A066-2B84000FEDCF"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2840"
+ endingLineNumber = "2840"
+ landmarkName = "montgomery(signature, exponent, modulus, base, valid)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "7B64655D-0DD5-4A2E-A066-2B84000FEDCF - 941e65ef5d4517f1"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2832"
+ endingLineNumber = "2832"
+ offsetFromSymbolStart = "88">
+ </Location>
+ <Location
+ uuid = "7B64655D-0DD5-4A2E-A066-2B84000FEDCF - 941e65ef5d451756"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2835"
+ endingLineNumber = "2835"
+ offsetFromSymbolStart = "88">
+ </Location>
+ <Location
+ uuid = "7B64655D-0DD5-4A2E-A066-2B84000FEDCF - 941e65ef5d451714"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "montgomery"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2837"
+ endingLineNumber = "2837"
+ offsetFromSymbolStart = "88">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "0307EA98-8D0F-4EAD-8ECA-A057B11BE8EB"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2586"
+ endingLineNumber = "2586"
+ landmarkName = "mpz_set_str(r, sp, base)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "C92C231A-CD9B-47F3-91FF-937373714089"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "856"
+ endingLineNumber = "856"
+ landmarkName = "mpz_div_qr(q, r, n, d, mode)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "BBB53FD7-763F-4D9E-85C9-A827113A6971"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2213"
+ endingLineNumber = "2213"
+ landmarkName = "mpz_gcdext(g, s, t, u, v)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "A13F2AE0-69B2-403C-A34C-03F8E2D0358E"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2664"
+ endingLineNumber = "2664"
+ landmarkName = "mont_prepare(b, e, m, r, r_1, ni, M, x)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "F1B8A13D-2350-47BF-BCEB-CDEAD7827EA9"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp_GPU.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "2919"
+ endingLineNumber = "2919"
+ landmarkName = "montgomery(signature, exponent, modulus, base, valid)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "DF3913BB-0DF6-48AA-A6C9-ECFDE4C32A54"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/gmp.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "999"
+ endingLineNumber = "999"
+ landmarkName = "mpn_div_qr_1_preinv(qp, np, nn, inv)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
diff --git a/xcode/montgomery.cl b/xcode/montgomery.cl
@@ -1,8 +1,17 @@
+//
+// gmp_GPU.c
+// lib-gpu-verify
+//
+// Created by Cedric Zwahlen on 25.11.2023.
+//
#ifndef MINI_GMP_LIMB_TYPE
#define MINI_GMP_LIMB_TYPE long
#endif
+
+#define ULONG_MAX_gpu 0xFFFFFFFFUL
+
#define GMP_LIMB_BITS (sizeof(mp_limb_t) * CHAR_BIT)
#define GMP_LIMB_MAX ((mp_limb_t) ~ (mp_limb_t) 0)
@@ -22,9 +31,8 @@
#define GMP_CMP(a,b) (((a) > (b)) - ((a) < (b)))
-#define assert(x)
-
-#define NULL 0
+#define GMP_MPN_OVERLAP_P(xp, xsize, yp, ysize) \
+ ((xp) + (xsize) > (yp) && (yp) + (ysize) > (xp))
#define gmp_clz(count, x) do { \
@@ -190,6 +198,11 @@
(y) = __mp_bitcnt_t_swap__tmp; \
} while (0)
+
+#define assert(x){}
+
+#define NULL ((void*)0)
+
typedef unsigned MINI_GMP_LIMB_TYPE mp_limb_t;
typedef long mp_size_t;
typedef unsigned long mp_bitcnt_t;
@@ -226,38 +239,226 @@ struct gmp_div_inverse
mp_limb_t di;
};
+
+struct mpn_base_info
+{
+ /* bb is the largest power of the base which fits in one limb, and
+ exp is the corresponding exponent. */
+ unsigned exp;
+ mp_limb_t bb;
+};
+
+
enum mpz_div_round_mode { GMP_DIV_FLOOR, GMP_DIV_CEIL, GMP_DIV_TRUNC };
-void mpz_sub (mpz_t r, const mpz_t a, const mpz_t b);
-void mpz_add (mpz_t, const mpz_t, const mpz_t);
-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_mod (mpz_t, const mpz_t, const mpz_t);
+void mpz_init (mpz_t r);
+void mpn_copyi (mp_ptr d, mp_srcptr s, mp_size_t n);
+void mpz_set (mpz_t r, const mpz_t x);
+void
+mpz_set (mpz_t r, const mpz_t x);
+void
+mpz_set_ui (mpz_t r, unsigned long int x);
+void
+mpz_set_si (mpz_t r, signed long int x);
+void
+mpz_init_set_si (mpz_t r, signed long int x);
+void
+mpz_init_set (mpz_t r, const mpz_t x);
+void
+mpz_init2 (mpz_t r, mp_bitcnt_t bits);
+void
+mpz_init_set_ui (mpz_t r, unsigned long int x);
+void
+mpz_clear (mpz_t r);
+void
+gmp_die (const char *msg);
-int mpz_sgn (const mpz_t);
-void mpz_mul (mpz_t, const mpz_t, const mpz_t);
-void mpz_mul_2exp (mpz_t, const mpz_t, mp_bitcnt_t);
+mp_size_t mpn_normalized_size (mp_srcptr xp, mp_size_t n);
+void
+mpz_add_ui (mpz_t r, const mpz_t a, unsigned long b);
+void
+mpz_ui_sub (mpz_t r, unsigned long a, const mpz_t b);
+void
+mpz_sub_ui (mpz_t r, const mpz_t a, unsigned long b);
+int
+mpn_absfits_ulong_p (mp_srcptr up, mp_size_t un);
+unsigned long int
+mpz_get_ui (const mpz_t u);
+int
+mpz_cmpabs_ui (const mpz_t u, unsigned long v);
+mp_limb_t
+mpn_sub_1 (mp_ptr rp, mp_srcptr ap, mp_size_t n, mp_limb_t b);
+mp_limb_t
+mpn_sub_n (mp_ptr rp, mp_srcptr ap, mp_srcptr bp, mp_size_t n);
+mp_limb_t
+mpn_sub (mp_ptr rp, mp_srcptr ap, mp_size_t an, mp_srcptr bp, mp_size_t bn);
+mp_limb_t
+mpn_invert_3by2 (mp_limb_t u1, mp_limb_t u0);
+int
+mpz_div_qr (mpz_t q, mpz_t r,
+ const mpz_t n, const mpz_t d, enum mpz_div_round_mode mode);
+void
+mpz_mod (mpz_t r, const mpz_t n, const mpz_t d);
+void
+mpn_div_qr_1_invert (struct gmp_div_inverse *inv, mp_limb_t d);
-void mpz_gcdext (mpz_t, mpz_t, mpz_t, const mpz_t, const mpz_t);
-void mpz_powm (mpz_t, const mpz_t, const mpz_t, const mpz_t);
+void
+mpn_div_qr_2_invert (struct gmp_div_inverse *inv,
+ mp_limb_t d1, mp_limb_t d0);
-void mpz_addmul (mpz_t, const mpz_t, const mpz_t);
+void
+mpn_div_qr_invert (struct gmp_div_inverse *inv,
+ mp_srcptr dp, mp_size_t dn);
+int
+mpz_cmp_ui (const mpz_t u, unsigned long v);
+int
+mpn_cmp (mp_srcptr ap, mp_srcptr bp, mp_size_t n);
+mp_limb_t
+mpn_lshift (mp_ptr rp, mp_srcptr up, mp_size_t n, unsigned int cnt);
+mp_limb_t
+mpn_rshift (mp_ptr rp, mp_srcptr up, mp_size_t n, unsigned int cnt);
+int
+mpz_invert (mpz_t r, const mpz_t u, const mpz_t m);
+mp_limb_t
+mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
+ const struct gmp_div_inverse *inv);
+mp_limb_t
+mpn_add_n (mp_ptr rp, mp_srcptr ap, mp_srcptr bp, mp_size_t n);
+void
+mpn_div_qr_2_preinv (mp_ptr qp, mp_ptr np, mp_size_t nn,
+ const struct gmp_div_inverse *inv);
+mp_limb_t
+mpn_submul_1 (mp_ptr rp, mp_srcptr up, mp_size_t n, mp_limb_t vl);
+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);
+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);
+void
+mpz_powm (mpz_t r, const mpz_t b, const mpz_t e, const mpz_t m);
+int
+mpn_cmp4 (mp_srcptr ap, mp_size_t an, mp_srcptr bp, mp_size_t bn);
+mp_size_t
+mpz_abs_sub (mpz_t r, const mpz_t a, const mpz_t b);
+mp_limb_t
+mpn_add_1 (mp_ptr rp, mp_srcptr ap, mp_size_t n, mp_limb_t b);
+mp_limb_t
+mpn_add (mp_ptr rp, mp_srcptr ap, mp_size_t an, mp_srcptr bp, mp_size_t bn);
+mp_size_t
+mpz_abs_add (mpz_t r, const mpz_t a, const mpz_t b);
+void
+mpz_sub (mpz_t r, const mpz_t a, const mpz_t b);
+mp_limb_t
+mpn_addmul_1 (mp_ptr rp, mp_srcptr up, mp_size_t n, mp_limb_t vl);
+mp_limb_t
+mpn_mul_1 (mp_ptr rp, mp_srcptr up, mp_size_t n, mp_limb_t vl);
+mp_limb_t
+mpn_mul (mp_ptr rp, mp_srcptr up, mp_size_t un, mp_srcptr vp, mp_size_t vn);
+void
+mpz_mul (mpz_t r, const mpz_t u, const mpz_t v);
+void
+mpn_copyd (mp_ptr d, mp_srcptr s, mp_size_t n);
+void
+mpn_zero (mp_ptr rp, mp_size_t n);
+void
+mpz_mul_2exp (mpz_t r, const mpz_t u, mp_bitcnt_t bits);
+int
+mpn_zero_p(mp_srcptr rp, mp_size_t n);
+void
+mpz_div_q_2exp (mpz_t q, const mpz_t u, mp_bitcnt_t bit_index,
+ enum mpz_div_round_mode mode);
+void
+mpz_tdiv_q_2exp (mpz_t r, const mpz_t u, mp_bitcnt_t cnt);
+int
+mpz_cmp (const mpz_t a, const mpz_t b);
+void
+mpz_add (mpz_t r, const mpz_t a, const mpz_t b);
+int
+mpz_tstbit (const mpz_t d, mp_bitcnt_t bit_index);
+mp_bitcnt_t
+mpn_limb_size_in_base_2 (mp_limb_t u);
+size_t
+mpz_sizeinbase (const mpz_t u, int base);
+int
+mpz_sgn (const mpz_t u);
+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_bitcnt_t
+mpn_scan1 (mp_srcptr ptr, mp_bitcnt_t bit);
+mp_bitcnt_t
+mpz_scan1 (mpz_t u, mp_bitcnt_t starting_bit);
+mp_bitcnt_t
+mpz_make_odd (mpz_t r);
+void
+mpz_tdiv_qr (mpz_t q, mpz_t r, const mpz_t n, const mpz_t d);
+void
+mpz_abs_add_bit (mpz_t d, mp_bitcnt_t bit_index);
+void
+mpz_abs_sub_bit (mpz_t d, mp_bitcnt_t bit_index);
+void
+mpz_setbit (mpz_t d, mp_bitcnt_t bit_index);
+void
+mpz_divexact (mpz_t q, const mpz_t n, const mpz_t d);
+int
+mpz_cmpabs (const mpz_t u, const mpz_t v);
+void
+mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v);
+void
+mpz_addmul_ui (mpz_t r, const mpz_t u, unsigned long int v);
-int mpz_tstbit (const mpz_t, mp_bitcnt_t);
+unsigned
+mpn_base_power_of_two_p (unsigned b);
+void
+mpn_get_base_info (struct mpn_base_info *info, mp_limb_t b);
+int isspace_gpu(unsigned char c);
+int strlen_c(__constant char *c);
+mp_size_t mpn_set_str_bits (mp_ptr rp, const unsigned char *sp, size_t sn,
+ unsigned bits);
+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);
+int
+mpz_set_str (mpz_t r, __constant char *sp, int base);
+int
+mpz_init_set_str (mpz_t r, __constant char *sp, int base);
-int mpz_cmp_ui (const mpz_t u, unsigned long v);
+//void mpz_sub (mpz_t r, const mpz_t a, const mpz_t b);
+////void mpz_add (mpz_t, const mpz_t, const mpz_t);
+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_mod (mpz_t, const mpz_t, const mpz_t);
+//
+////int mpz_sgn (const mpz_t);
+//
+////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_gcdext (mpz_t, mpz_t, mpz_t, const mpz_t, const mpz_t);
+////void mpz_powm (mpz_t, const mpz_t, const mpz_t, const mpz_t);
+//
+void mpz_addmul (mpz_t, const mpz_t, const mpz_t);
+//
+//int mpz_tstbit (const mpz_t, mp_bitcnt_t);
+//
+//int mpz_cmp_ui (const mpz_t u, unsigned long v);
+//
void mpn_div_qr (mp_ptr qp, mp_ptr np, mp_size_t nn, mp_srcptr dp, mp_size_t dn);
+//
+//mp_limb_t mpn_invert_3by2 (mp_limb_t, mp_limb_t);
-mp_limb_t mpn_invert_3by2 (mp_limb_t, mp_limb_t);
#define mpn_invert_limb(x) mpn_invert_3by2 ((x), 0)
#define MPZ_REALLOC(z,n) (z)->_mp_d
-
-
void
mpz_init (mpz_t r)
{
@@ -265,7 +466,10 @@ mpz_init (mpz_t r)
r->_mp_alloc = 0;
r->_mp_size = 0;
-// r->_mp_d = (mp_ptr) &dummy_limb;
+
+ // memset(r->_mp_d, 0, 256);
+
+ // r->_mp_d = (mp_ptr) &dummy_limb;
}
void
@@ -286,7 +490,7 @@ mpz_set (mpz_t r, const mpz_t x)
mp_ptr rp;
n = GMP_ABS (x->_mp_size);
- //rp = MPZ_REALLOC (r, n);
+ rp = MPZ_REALLOC (r, n);
mpn_copyi (rp, x->_mp_d, n);
r->_mp_size = x->_mp_size;
@@ -300,14 +504,14 @@ mpz_set_ui (mpz_t r, unsigned long int x)
if (x > 0)
{
r->_mp_size = 1;
- //MPZ_REALLOC (r, 1)[0] = x;
+ 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;
+ MPZ_REALLOC (r, r->_mp_size)[r->_mp_size - 1] = x;
}
}
}
@@ -338,7 +542,7 @@ mpz_set_si (mpz_t r, signed long int x)
else
{
r->_mp_size = -1;
- //MPZ_REALLOC (r, 1)[0] = GMP_NEG_CAST (unsigned long int, x);
+ MPZ_REALLOC (r, 1)[0] = GMP_NEG_CAST (unsigned long int, x);
}
}
@@ -389,7 +593,7 @@ void
gmp_die (const char *msg)
{
//fprintf (stderr, "%s\n", msg);
- abort();
+ //abort();
}
mp_size_t mpn_normalized_size (mp_srcptr xp, mp_size_t n)
@@ -430,7 +634,7 @@ mpn_absfits_ulong_p (mp_srcptr up, mp_size_t un)
mp_limb_t ulongrem = 0;
if (GMP_ULONG_BITS % GMP_LIMB_BITS != 0)
- ulongrem = (mp_limb_t) (ULONG_MAX >> GMP_LIMB_BITS * ulongsize) + 1;
+ ulongrem = (mp_limb_t) (ULONG_MAX_gpu >> GMP_LIMB_BITS * ulongsize) + 1;
return un <= ulongsize || (up[ulongsize] < ulongrem && un == ulongsize + 1);
}
@@ -741,6 +945,54 @@ mpz_div_qr (mpz_t q, mpz_t r,
}
void
+mpn_div_qr (mp_ptr qp, mp_ptr np, mp_size_t nn, mp_srcptr dp, mp_size_t dn)
+{
+ struct gmp_div_inverse inv;
+ // mp_ptr tp = NULL;
+
+ mpz_t tp;
+
+ assert (dn > 0);
+ assert (nn >= dn);
+
+ mpn_div_qr_invert (&inv, dp, dn);
+ if (dn > 2 && inv.shift > 0)
+ {
+ //tp = gmp_alloc_limbs (dn);
+ gmp_assert_nocarry (mpn_lshift (tp->_mp_d, dp, dn, inv.shift));
+ dp = tp->_mp_d;
+ }
+ mpn_div_qr_preinv (qp, np, nn, dp, dn, &inv);
+ if (tp) {}
+ //gmp_free_limbs (tp, dn);
+}
+
+void
+mpz_addmul (mpz_t r, const mpz_t u, const mpz_t v)
+{
+ mpz_t t;
+ mpz_init (t);
+ mpz_mul (t, u, v);
+ mpz_add (r, r, t);
+ mpz_clear (t);
+}
+
+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;
+
+}
+
+void
mpz_mod (mpz_t r, const mpz_t n, const mpz_t d)
{
mpz_div_qr (NULL, r, n, d, d->_mp_size >= 0 ? GMP_DIV_FLOOR : GMP_DIV_CEIL);
@@ -937,7 +1189,8 @@ mpn_div_qr_1_preinv (mp_ptr qp, mp_srcptr np, mp_size_t nn,
tp = qp;
if (!tp)
{
- // tn = nn;
+ tn = nn;
+
// tp = gmp_alloc_limbs (tn);
}
r = mpn_lshift (tp, np, nn, inv->shift);
@@ -1764,7 +2017,7 @@ mpn_scan1 (mp_srcptr ptr, mp_bitcnt_t bit)
}
mp_bitcnt_t
-mpz_scan1 (const mpz_t u, mp_bitcnt_t starting_bit)
+mpz_scan1 (mpz_t u, mp_bitcnt_t starting_bit)
{
mp_ptr up;
mp_size_t us, un, i;
@@ -2061,7 +2314,7 @@ mpz_gcdext (mpz_t g, mpz_t s, mpz_t t, const mpz_t u, const mpz_t v)
mpz_sub (s0, s0, s1);
mpz_add (t0, t0, t1);
}
- assert (mpz_even_p (t0) && mpz_even_p (s0));
+ //assert (mpz_even_p (t0) && mpz_even_p (s0));
mpz_tdiv_q_2exp (s0, s0, 1);
mpz_tdiv_q_2exp (t0, t0, 1);
}
@@ -2123,13 +2376,7 @@ mpn_base_power_of_two_p (unsigned b)
}
}
-struct mpn_base_info
-{
- /* bb is the largest power of the base which fits in one limb, and
- exp is the corresponding exponent. */
- unsigned exp;
- mp_limb_t bb;
-};
+
void
mpn_get_base_info (struct mpn_base_info *info, mp_limb_t b)
@@ -2146,13 +2393,13 @@ mpn_get_base_info (struct mpn_base_info *info, mp_limb_t b)
info->bb = p;
}
-int isspace(unsigned char c) {
+int isspace_gpu(unsigned char c) {
if (c == '\n' || c == ' ' || c == '\t' || c == '\r' || c == '\f' || c == '\v')
return 1;
return 0;
}
-int strlen(__constant char *c) {
+int strlen_c(__constant char *c) {
// rather naive implementation – we assume a string is terminated, and is not 0 characters long.
@@ -2160,6 +2407,7 @@ int strlen(__constant char *c) {
while (1) {
if (c[i] == '\0')
return i;
+ i++;
}
return i;
}
@@ -2239,11 +2487,11 @@ mpz_set_str (mpz_t r, __constant char *sp, int base)
mp_ptr rp;
size_t dn, sn;
int sign;
- unsigned char dp[2048];
+ unsigned char dp[4096];
assert (base == 0 || (base >= 2 && base <= 62));
- while (isspace( (unsigned char) *sp))
+ while (isspace_gpu( (unsigned char) *sp))
sp++;
sign = (*sp == '-');
@@ -2275,7 +2523,7 @@ mpz_set_str (mpz_t r, __constant char *sp, int base)
r->_mp_size = 0;
return -1;
}
- sn = strlen(sp);
+ sn = strlen_c(sp);
//dp = (unsigned char *) gmp_alloc (sn);
@@ -2284,7 +2532,7 @@ mpz_set_str (mpz_t r, __constant char *sp, int base)
{
unsigned digit;
- if (isspace ((unsigned char) *sp))
+ if (isspace_gpu ((unsigned char) *sp))
continue;
else if (*sp >= '0' && *sp <= '9')
digit = *sp - '0';
@@ -2348,12 +2596,6 @@ mpz_init_set_str (mpz_t r, __constant char *sp, int base)
}
-
-
-
-
-
-
// Montgomery multiplication
void mont_prepare(mpz_t b, mpz_t e, mpz_t m,
@@ -2423,6 +2665,7 @@ void mont_prepare(mpz_t b, mpz_t e, mpz_t m,
mpz_set_si(one, 0);
+
mpz_gcdext(one, r_1, ni, r, m); // set r_1 and ni
int sgn = mpz_sgn(r_1);
@@ -2555,7 +2798,6 @@ void mont_product(mpz_t ret,
}
-
// not the fastest... but it does not increase the variable sizes
void mont_mulmod(mpz_t res, const mpz_t a, const mpz_t b, const mpz_t mod) {
@@ -2579,11 +2821,19 @@ void mont_mulmod(mpz_t res, const mpz_t a, const mpz_t b, const mpz_t mod) {
-__kernel void montgomery(__constant unsigned long *valid, __constant char *base, __constant char *exponent, __constant char *modulus, __constant char *signature) {
+__kernel void montgomery(__constant char *signature,
+ __constant char *exponent,
+ __constant char *modulus,
+ __constant char *base,
+ __global unsigned long *valid)
+{
+
int radix = 16;
- mpz_t b,e,m, res;
+ mpz_t b,e,m,res;
+
+
mpz_init(res);
@@ -2598,9 +2848,13 @@ __kernel void montgomery(__constant unsigned long *valid, __constant char *base,
mpz_init(M);
mpz_init(x);
+
mpz_t xx;
mpz_init(xx);
+
+
+
if (mpz_even_p(m)) {
mpz_t bb, x1, x2, q, powj;
@@ -2640,8 +2894,6 @@ __kernel void montgomery(__constant unsigned long *valid, __constant char *base,
mpz_set(res, x1);
- printf("--\n");
-
} else {
@@ -2650,10 +2902,27 @@ __kernel void montgomery(__constant unsigned long *valid, __constant char *base,
mont_modexp(xx, x, e, M, m, ni, r, r_1);
- mont_finish(res, xx, m, ni, r, r_1);
+ mont_finish(res, xx, m, ni, r, r_1);
+
}
+
+
+
+
+ mpz_t sig;
+ mpz_init_set_str(sig,signature,radix);
+
+ if (mpz_cmp(sig,res) == 0) {
+
+ *valid = 1;
+
+ } else {
+
+ }
+
+
}
