commit 78ab2791f6a47cfbf4dfc3baa10382023c213dfe
parent f7d71405a6d636b67b0d34a32c9afdd2edfc8a9b
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Fri, 17 Nov 2023 01:51:08 +0100
Batch preparations complete
But there seems to be an issue, either with memory mapping or the amount of data to be processed
Diffstat:
5 files changed, 722 insertions(+), 154 deletions(-)
diff --git a/source/big-int-test.c b/source/big-int-test.c
@@ -787,7 +787,8 @@ 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, size_t edigits)
+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 */
@@ -800,10 +801,14 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
DIGIT_T t2[nn];
DIGIT_T y[nn];
+ DIGIT_T const *window_x = &x[0];
+ DIGIT_T const *window_e = &e[0];
+ DIGIT_T *window_m = &m[0];
+
assert(ndigits <= MAX_FIXED_DIGITS);
assert(ndigits != 0);
- n = mpSizeof(e, edigits);
+ n = mpSizeof(window_e, edigits);
/* Catch e==0 => x^0=1 */
if (0 == n)
{
@@ -819,7 +824,7 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
mpNEXTBITMASK(mask, n);
/* Set y = x */
- mpSetEqual(y, x, ndigits);
+ mpSetEqual(y, window_x, ndigits);
// the number of bits in e
@@ -831,11 +836,11 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
for(xyz = bitlength_n; xyz > 0; xyz--)
{
/* Square y = y * y mod n */
- mpMODSQUARETEMP(y, m, ndigits, t1, t2);
+ mpMODSQUARETEMP(y, window_m, 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);
+ mpMODMULTTEMP(y, window_x, window_m, ndigits, t1, t2);
}
/* Move to next bit */
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -22,7 +22,6 @@
#define NEED_LIBGCRYPT_VERSION "1.10.2"
-#define DATA_SIZE (1024)
void setup_gcry(void) {
@@ -53,14 +52,14 @@ void generate_random_pairs(DIGIT_T *bases, size_t *b_len,
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++) {
+ char *template = "(genkey(rsa(nbits 4:2048)))";
+ gcry_sexp_t parms;
+
+ gcry_sexp_new(&parms, template, strlen(template), 1);
+
+
gcry_sexp_t key;
gcry_pk_genkey(&key,parms);
@@ -105,32 +104,32 @@ void generate_random_pairs(DIGIT_T *bases, size_t *b_len,
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];
+ DIGIT_T base [sz*2];
+ DIGIT_T exponent [sz*2];
+ DIGIT_T modulus [sz*2];
+ DIGIT_T signature [sz*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);
+ mpSetZero(base, sz*2);
+ mpSetZero(exponent, sz*2);
+ mpSetZero(modulus, sz*2);
+ mpSetZero(signature, sz*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) );
+ size_t max_len = max( max( mpSizeof(base, sz*2), mpSizeof(modulus, sz*2) ), mpSizeof(signature, sz*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;
+ b_len[i] = (i == 0 ? 0 : b_len[i - 1]) + max_len;
+ e_len[i] = (i == 0 ? 0 : e_len[i - 1]) + mpSizeof(exponent, sz*2);
+ m_len[i] = (i == 0 ? 0 : m_len[i - 1]) + max_len;
+ s_len[i] = (i == 0 ? 0 : s_len[i - 1]) + 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));
+ 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);
@@ -146,17 +145,255 @@ void generate_random_pairs(DIGIT_T *bases, size_t *b_len,
}
-int rsa_tests(void) {
+int verify_with_opencl(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 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
+
+
+ // Connect to a compute device
+ //
+ int gpu = 1;
+ err = clGetDeviceIDs(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!\n");
+ 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);
+
+ //printf("device address bits: %i\n", address_bits);
+
+
+ // 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("./verify.cl", "r");
+ 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, NULL, &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[2048];
+
+ 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);
+ }
+
+ // Create the compute kernel in the program we wish to run
+ //
+ kernel = clCreateKernel(program, "several", &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;
+
+ mod_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(DIGIT_T) * m_len[n-1], NULL, NULL);
+ exp_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(DIGIT_T) * e_len[n-1], NULL, NULL);
+ sig_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(DIGIT_T) * s_len[n-1], NULL, NULL);
+ comp_mem = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(DIGIT_T) * b_len[n-1], NULL, NULL); // the base, to compare whether we get the same signature
+
+ mod_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(size_t) * n, NULL, NULL);
+ exp_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(size_t) * n, NULL, NULL);
+ sig_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(size_t) * n, NULL, NULL);
+ comp_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(size_t) * 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, sizeof(DIGIT_T) * s_len[n-1], signatures, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, sig_len, CL_TRUE, 0,sizeof(size_t) * n, s_len, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, exp_mem, CL_TRUE, 0,sizeof(DIGIT_T) * e_len[n-1], exponents, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, exp_len, CL_TRUE, 0,sizeof(size_t) * n, e_len, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, mod_mem, CL_TRUE, 0, sizeof(DIGIT_T) * m_len[n-1], moduli, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, mod_len, CL_TRUE, 0,sizeof(size_t) * n, m_len, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, comp_mem, CL_TRUE, 0, sizeof(DIGIT_T) * b_len[n-1], bases, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(commands, comp_len, CL_TRUE, 0,sizeof(size_t) * 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, 2, sizeof(cl_mem), &exp_mem);
+ err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &exp_len);
+ err |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &mod_mem);
+ err |= clSetKernelArg(kernel, 5, sizeof(cl_mem), &mod_len);
+ err |= clSetKernelArg(kernel, 6, sizeof(cl_mem), &comp_mem);
+ err |= clSetKernelArg(kernel, 7, sizeof(cl_mem), &comp_len);
+ err |= clSetKernelArg(kernel, 8, sizeof(cl_mem), &valid);
+ err |= clSetKernelArg(kernel, 9, sizeof(unsigned int), &n);
+
+ //err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);
+ if (err != CL_SUCCESS)
+ {
+ printf("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;
+ }
+
+ // 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);
+ }
+
+ printf("VERIFICATION RESULT: %lu\n",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);
+
+}
+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
+ int gen_n_pairs = 4;
- DIGIT_T *q = malloc(2048);
- DIGIT_T *r = malloc(2048);
- DIGIT_T *s = malloc(2048);
- DIGIT_T *t = malloc(2048);
+ DIGIT_T *q = malloc(32768); // does not set memory to 0 on linux, keep that in mind
+ DIGIT_T *r = malloc(32768);
+ DIGIT_T *s = malloc(32768);
+ DIGIT_T *t = malloc(32768);
size_t *u = malloc(gen_n_pairs * sizeof(size_t));
size_t *v = malloc(gen_n_pairs * sizeof(size_t));
@@ -168,6 +405,23 @@ int rsa_tests(void) {
s, w,
t, x, gen_n_pairs);
+
+
+ verify_with_opencl(q, u,
+ r, v,
+ s, w,
+ t, x, gen_n_pairs);
+
+
+
+
+
+
+
+
+
+
+
DIGIT_T *y = malloc(2048);
DIGIT_T *s_window = &s[0];
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 = "478"
- endingLineNumber = "478"
+ startingLineNumber = "732"
+ endingLineNumber = "732"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -738,8 +738,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "492"
- endingLineNumber = "492"
+ startingLineNumber = "746"
+ endingLineNumber = "746"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -891,8 +891,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "675"
- endingLineNumber = "675"
+ startingLineNumber = "929"
+ endingLineNumber = "929"
landmarkName = "verify(sign, ee, nn, mm)"
landmarkType = "9">
</BreakpointContent>
@@ -939,8 +939,8 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "808"
- endingLineNumber = "808"
+ startingLineNumber = "813"
+ endingLineNumber = "813"
landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
</BreakpointContent>
@@ -955,8 +955,8 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "831"
- endingLineNumber = "831"
+ startingLineNumber = "836"
+ endingLineNumber = "836"
landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
<Locations>
@@ -1018,8 +1018,8 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "847"
- endingLineNumber = "847"
+ startingLineNumber = "852"
+ endingLineNumber = "852"
landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
<Locations>
@@ -1081,8 +1081,8 @@
filePath = "../source/big-int-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "832"
- endingLineNumber = "832"
+ startingLineNumber = "837"
+ endingLineNumber = "837"
landmarkName = "mpModExpO(yout, x, e, m, ndigits, edigits)"
landmarkType = "9">
</BreakpointContent>
@@ -1400,8 +1400,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "702"
- endingLineNumber = "702"
+ startingLineNumber = "956"
+ endingLineNumber = "956"
landmarkName = "unknown"
landmarkType = "0">
</BreakpointContent>
@@ -1496,24 +1496,8 @@
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"
+ startingLineNumber = "67"
+ endingLineNumber = "67"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1528,8 +1512,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "188"
- endingLineNumber = "188"
+ startingLineNumber = "442"
+ endingLineNumber = "442"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -1711,8 +1695,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "663"
- endingLineNumber = "663"
+ startingLineNumber = "917"
+ endingLineNumber = "917"
landmarkName = "verify(sign, ee, nn, mm)"
landmarkType = "9">
</BreakpointContent>
@@ -1727,8 +1711,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "177"
- endingLineNumber = "177"
+ startingLineNumber = "431"
+ endingLineNumber = "431"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -1792,6 +1776,21 @@
endingLineNumber = "177"
offsetFromSymbolStart = "371">
</Location>
+ <Location
+ uuid = "E76A4300-645A-48D3-AFAA-F40E9454639D - b0b9078e770c8f6b"
+ 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 = "427"
+ endingLineNumber = "427"
+ offsetFromSymbolStart = "402">
+ </Location>
</Locations>
</BreakpointContent>
</BreakpointProxy>
@@ -1805,8 +1804,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "108"
- endingLineNumber = "108"
+ startingLineNumber = "107"
+ endingLineNumber = "107"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1821,8 +1820,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "648"
- endingLineNumber = "648"
+ startingLineNumber = "902"
+ endingLineNumber = "902"
landmarkName = "verify(sign, ee, nn, mm)"
landmarkType = "9">
</BreakpointContent>
@@ -1837,8 +1836,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "135"
- endingLineNumber = "135"
+ startingLineNumber = "134"
+ endingLineNumber = "134"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
</BreakpointContent>
@@ -1853,8 +1852,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "130"
- endingLineNumber = "130"
+ startingLineNumber = "129"
+ endingLineNumber = "129"
landmarkName = "generate_random_pairs(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
landmarkType = "9">
<Locations>
@@ -1888,6 +1887,51 @@
endingLineNumber = "130"
offsetFromSymbolStart = "1123">
</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 = "1465">
+ </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 = "1668">
+ </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 = "1784">
+ </Location>
</Locations>
</BreakpointContent>
</BreakpointProxy>
@@ -1901,11 +1945,264 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "166"
- endingLineNumber = "166"
+ startingLineNumber = "403"
+ endingLineNumber = "403"
landmarkName = "rsa_tests()"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "310"
+ endingLineNumber = "310"
+ landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE - acfd4e71f26ef54b"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "verify_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "298"
+ endingLineNumber = "298"
+ offsetFromSymbolStart = "1959">
+ </Location>
+ <Location
+ uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE - acfd4e71f26ef37a"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "verify_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "313"
+ endingLineNumber = "313"
+ offsetFromSymbolStart = "2488">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "419"
+ endingLineNumber = "419"
+ landmarkName = "rsa_tests()"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8ee7"
+ 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 = "423"
+ endingLineNumber = "423"
+ offsetFromSymbolStart = "454">
+ </Location>
+ <Location
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e42"
+ 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 = "420"
+ endingLineNumber = "420"
+ offsetFromSymbolStart = "398">
+ </Location>
+ <Location
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e04"
+ 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 = "422"
+ endingLineNumber = "422"
+ offsetFromSymbolStart = "398">
+ </Location>
+ <Location
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e63"
+ 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 = "419"
+ endingLineNumber = "419"
+ offsetFromSymbolStart = "398">
+ </Location>
+ <Location
+ uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8ea9"
+ 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 = "425"
+ endingLineNumber = "425"
+ offsetFromSymbolStart = "410">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "470BAF83-0588-455C-AE68-F686E9954517"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "366"
+ endingLineNumber = "366"
+ landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "470BAF83-0588-455C-AE68-F686E9954517 - acfd4e71f26efc6c"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "verify_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "367"
+ endingLineNumber = "367"
+ offsetFromSymbolStart = "3307">
+ </Location>
+ <Location
+ uuid = "470BAF83-0588-455C-AE68-F686E9954517 - acfd4e71f26efcb2"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "verify_with_opencl"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "369"
+ endingLineNumber = "369"
+ offsetFromSymbolStart = "3307">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "9B0E2741-A817-4815-8AE4-26ED0DDEB4A6"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "410"
+ endingLineNumber = "410"
+ landmarkName = "rsa_tests()"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "9B0E2741-A817-4815-8AE4-26ED0DDEB4A6 - b0b9078e770c8c9b"
+ 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 = "411"
+ endingLineNumber = "411"
+ offsetFromSymbolStart = "310">
+ </Location>
+ <Location
+ uuid = "9B0E2741-A817-4815-8AE4-26ED0DDEB4A6 - b0b9078e770c8d5d"
+ 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 = "413"
+ endingLineNumber = "413"
+ offsetFromSymbolStart = "310">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "164290F2-14BC-4321-8C37-498198D7FC1A"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "333"
+ endingLineNumber = "333"
+ landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
</Breakpoints>
</Bucket>
diff --git a/xcode/verify.cl b/xcode/verify.cl
@@ -1,7 +1,6 @@
// macros
-#define mpDESTROY(b, n) do{if(b)mpSetZero(b,n);}while(0)
#define max(a,b) (((a) > (b)) ? (a) : (b))
// only for that string conversion
@@ -140,9 +139,6 @@ int mpModulo(__global DIGIT_T *r, DIGIT_T *u, size_t udigits, __global DIGIT_T *
/* Final r is only vdigits long */
mpSetEqual_gl(r, rr, vdigits);
- mpDESTROY(rr, udigits);
- mpDESTROY(qq, udigits);
-
return 0;
}
@@ -1200,82 +1196,98 @@ void assert(bool precondition) {
}
-__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
+__kernel void several(__global DIGIT_T* x, __global const size_t *s_len,
+ __global DIGIT_T* e, __global const size_t *e_len,
+ __global DIGIT_T* m, __global const size_t *n_len,
+ __global DIGIT_T *mm, __global const size_t *mm_len,
+ __global unsigned long* valid,
+ const int count
) {
int index = get_global_id(0);
+
- // the result is copied in here, compare it to mm
- DIGIT_T yout[MAX_FIXED_DIGITS * 2];
-
- DIGIT_T mask;
- 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(window_e, ndigits);
- /* Catch e==0 => x^0=1 */
- if (0 == n)
- {
- mpSetDigit(yout, 1, ndigits);
- goto done;
- }
- /* Find second-most significant bit in e */
- for (mask = HIBITMASK; mask > 0; mask >>= 1)
- {
- if (window_e[n-1] & mask)
- break;
- }
- mpNEXTBITMASK(mask, n);
-
- /* Set y = x */
- 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, window_mm, ndigits, t1, t2);
+ if (index < count) {
+
+ int ndigits = max( max( n_len[index], s_len[index]), mm_len[index]);
+ int edigits = e_len[index];
+ // the result is copied in here, compare it to mm
+ DIGIT_T yout[MAX_ALLOC_SIZE *2];
- if (e[n-1] & mask)
- { /* if e(j) == 1 then multiply
- y = y * x mod n */
- mpMODMULTTEMP(y, window_x, window_m, ndigits, t1, t2);
+ DIGIT_T mask;
+ size_t n;
+ size_t nn = ndigits * 2;
+
+ __global DIGIT_T * __private window_x;
+ __global DIGIT_T * __private window_e;
+ __global DIGIT_T * __private window_m;
+ __global DIGIT_T * __private window_mm;
+ //
+ window_x = &x[index == 0 ? 0 : (s_len[index - 1])];
+ window_e = &e[index == 0 ? 0 : (e_len[index - 1])];
+ window_m = &m[index == 0 ? 0 : (n_len[index - 1])];
+ window_mm = &mm[index == 0 ? 0 : (mm_len[index - 1])];
+
+ // can probably be smaller
+ __private DIGIT_T t1[MAX_ALLOC_SIZE *2];
+ __private DIGIT_T t2[MAX_ALLOC_SIZE *2];
+ __private DIGIT_T y[MAX_ALLOC_SIZE *2];
+
+ assert(ndigits <= MAX_FIXED_DIGITS);
+ assert(ndigits != 0);
+
+ n = mpSizeof_g(window_e, edigits);
+ /* Catch e==0 => x^0=1 */
+ if (0 == n)
+ {
+ mpSetDigit(yout, 1, ndigits);
}
-
- /* Move to next bit */
+ /* Find second-most significant bit in e */
+ for (mask = HIBITMASK; mask > 0; mask >>= 1)
+ {
+ if (window_e[n-1] & mask)
+ break;
+ }
mpNEXTBITMASK(mask, n);
+
+ /* Set y = x */
+ 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, window_m, ndigits, t1, t2);
+
+
+ if (e[n-1] & mask)
+ { /* if e(j) == 1 then multiply
+ y = y * x mod n */
+ mpMODMULTTEMP(y, window_x, window_m, ndigits, t1, t2);
+
+ }
+
+ /* Move to next bit */
+ mpNEXTBITMASK(mask, n);
+ }
+
+ mpSetEqual(yout, y, ndigits);
+
+ int len = ( mm_len[index] - (index == 0 ? 0 : mm_len[index - 1]) );
+
+
+ // MARK: valid cannot be written to by several at once (the same unit anyway)
+ if (mpCompare_lg(yout,window_mm,len) == 0 && index == 0) {
+ *valid |= 0x1 << index;
+ }
+
+
+
}
- mpSetEqual(yout, y, ndigits);
-
- int len = ( mm_len[index] - (index == 0 ? 0 : mm_len[index]) );
-
- // equal
- *valid = mpCompare_lg(yout,window_mm,len) == 0 ? 1 : 0;
+ // if (index == 8) { *valid = 0xBA; }
-done:
- mpDESTROY(t1, nn);
- mpDESTROY(t2, nn);
- mpDESTROY(y, ndigits);
}
