commit e1f873f49be90a485f4bf9510feb91400c966c82
parent 94c4edf4a5b936503b6e7afb84c59e9b33376761
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Tue, 7 Nov 2023 11:46:58 +0100
some minor changes to the kernel
Diffstat:
8 files changed, 223 insertions(+), 68 deletions(-)
diff --git a/.DS_Store b/.DS_Store
Binary files differ.
diff --git a/source/big-int-test.c b/source/big-int-test.c
@@ -545,13 +545,17 @@ int spMultiply(uint32_t p[2], uint32_t x, uint32_t y)
uint32_t spDivide(uint32_t *pq, uint32_t *pr, const uint32_t u[2], uint32_t v)
{
- uint64_t uu, q;
+
+ uint64_t uu, q;
uu = (uint64_t)u[1] << 32 | (uint64_t)u[0];
q = uu / (uint64_t)v;
//r = uu % (uint64_t)v;
*pr = (uint32_t)(uu - q * v);
*pq = (uint32_t)(q & 0xFFFFFFFF);
return (uint32_t)(q >> 32);
+
+
+
}
int mpCompare(const DIGIT_T a[], const DIGIT_T b[], size_t ndigits)
@@ -823,9 +827,15 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
/* Set y = x */
mpSetEqual(y, x, ndigits);
+
+
+ // the number of bits in e
+ size_t bitlength_n = log2(mask) + sizeof(DIGIT_T) * (n - 1) + 1;
+
+ //size_t ctr = 0;
- /* For bit j = k-2 downto 0 */
- while (n)
+ size_t xyz = 0;
+ for(xyz = bitlength_n; xyz > 0; xyz--)
{
/* Square y = y * y mod n */
mpMODSQUARETEMP(y, m, ndigits, t1, t2);
@@ -836,7 +846,10 @@ int mpModExpO(DIGIT_T *yout, const DIGIT_T *x, const DIGIT_T *e, DIGIT_T *m, siz
}
/* Move to next bit */
- mpNEXTBITMASK(mask, n);
+ mpNEXTBITMASK(mask, n);
+
+ // ctr++;
+
}
/* Return y */
diff --git a/source/big-int-test.h b/source/big-int-test.h
@@ -18,6 +18,8 @@
#include <string.h> // only used for the convert from hex function
#include <assert.h>
+#include <math.h>
+
// MARK: definitions
typedef uint32_t DIGIT_T; // for gpu might need to be half? is that half?
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -21,6 +21,8 @@
int rsa_tests(void) {
+
+
// MARK: UNSAFE init
// consider disabling optimizations, since they dont make for a fair comparison
@@ -359,9 +361,7 @@ int rsa_tests(void) {
}
- struct timespec t3, t4;
- clock_gettime(CLOCK_REALTIME, &t3);
// Execute the kernel over the entire range of our 1d input data set
// using the maximum number of work group items for this device
@@ -373,12 +373,20 @@ int rsa_tests(void) {
printf("Error: Failed to execute kernel!\n");
return EXIT_FAILURE;
}
+
+ struct timespec t3, t4;
+
+ clock_gettime(CLOCK_REALTIME, &t3);
// Wait for the command commands to get serviced before reading back results
//
clFinish(commands);
+ clock_gettime(CLOCK_REALTIME, &t4);
+
+ float seconds_2 = (t3.tv_nsec - t4.tv_nsec) / 1000;
+ printf("\nGPU verification: %f micro seconds\n", seconds_2);
// Read back the results from the device to verify the output
@@ -390,11 +398,7 @@ int rsa_tests(void) {
exit(1);
}
- clock_gettime(CLOCK_REALTIME, &t4);
-
- float seconds_2 = (t3.tv_nsec - t4.tv_nsec) / 1000;
- printf("\nGPU verification: %f micro seconds\n", seconds_2);
size_t sz_res = mpSizeof(res_buf, MAX_ALLOC_SIZE*2);
diff --git a/xcode/.DS_Store b/xcode/.DS_Store
Binary files differ.
diff --git a/xcode/lib-gpu-verify.xcodeproj/project.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 = "408"
- endingLineNumber = "408"
+ startingLineNumber = "412"
+ endingLineNumber = "412"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -783,8 +783,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "312"
- endingLineNumber = "312"
+ startingLineNumber = "314"
+ endingLineNumber = "314"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -831,8 +831,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "325"
- endingLineNumber = "325"
+ startingLineNumber = "327"
+ endingLineNumber = "327"
landmarkName = "rsa_tests()"
landmarkType = "9">
<Locations>
@@ -984,8 +984,8 @@
filePath = "../source/rsa-test.c"
startingColumnNumber = "9223372036854775807"
endingColumnNumber = "9223372036854775807"
- startingLineNumber = "515"
- endingLineNumber = "515"
+ startingLineNumber = "519"
+ endingLineNumber = "519"
landmarkName = "verify(sign, ee, nn, mm)"
landmarkType = "9">
</BreakpointContent>
@@ -1022,5 +1022,179 @@
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "1755B2D7-6C20-42C4-80E3-1B33107D8061"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "815"
+ endingLineNumber = "815"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "11706F8A-090F-445E-8F26-D709842A002D"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "838"
+ endingLineNumber = "838"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "11706F8A-090F-445E-8F26-D709842A002D - 8650bac5cd4fd621"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "828"
+ endingLineNumber = "828"
+ offsetFromSymbolStart = "475">
+ </Location>
+ <Location
+ uuid = "11706F8A-090F-445E-8F26-D709842A002D - 8650bac5cd4fd667"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "830"
+ endingLineNumber = "830"
+ offsetFromSymbolStart = "594">
+ </Location>
+ <Location
+ uuid = "11706F8A-090F-445E-8F26-D709842A002D - 8650bac5cd4fd6eb"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "834"
+ endingLineNumber = "834"
+ offsetFromSymbolStart = "617">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "B79E85F3-5E9B-4CE0-B21A-3C70696435B8"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "856"
+ endingLineNumber = "856"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ landmarkType = "9">
+ <Locations>
+ <Location
+ uuid = "B79E85F3-5E9B-4CE0-B21A-3C70696435B8 - 8650bac5cd4fd8da"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "849"
+ endingLineNumber = "849"
+ offsetFromSymbolStart = "925">
+ </Location>
+ <Location
+ uuid = "B79E85F3-5E9B-4CE0-B21A-3C70696435B8 - 8650bac5cd4fd8fb"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "850"
+ endingLineNumber = "850"
+ offsetFromSymbolStart = "969">
+ </Location>
+ <Location
+ uuid = "B79E85F3-5E9B-4CE0-B21A-3C70696435B8 - 8650bac5cd4fdb18"
+ shouldBeEnabled = "Yes"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ symbolName = "mpModExpO"
+ moduleName = "lib-gpu-verify"
+ usesParentBreakpointCondition = "Yes"
+ urlString = "file:///Users/cedriczwahlen/libgpuverify/source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "851"
+ endingLineNumber = "851"
+ offsetFromSymbolStart = "980">
+ </Location>
+ </Locations>
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "E91582E0-58FD-4DDB-9944-DD7D8E1EDA8A"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "839"
+ endingLineNumber = "839"
+ landmarkName = "mpModExpO(yout, x, e, m, ndigits)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "7A159638-F11E-4DFA-9DD6-9431F4B807BF"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/big-int-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "548"
+ endingLineNumber = "548"
+ landmarkName = "spDivide(pq, pr, u, v)"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
</Breakpoints>
</Bucket>
diff --git a/xcode/verify.cl b/xcode/verify.cl
@@ -131,20 +131,9 @@ int mpModulo(__global DIGIT_T *r, DIGIT_T *u, size_t udigits, __global DIGIT_T *
Use remainder from mpDivide function.
*/
- size_t nn = max(udigits, vdigits);
-/* Allocate temp storage */
-//#ifdef NO_ALLOCS
- // [v2.6] increased to two times
DIGIT_T qq[MAX_FIXED_DIGITS*2];
DIGIT_T rr[MAX_FIXED_DIGITS*2];
- // assert(nn <= (MAX_FIXED_DIGITS*2));
-/*#else
- DIGIT_T *qq, *rr;
- qq = mpAlloc(udigits);
- rr = mpAlloc(nn);
-#endif
-*/
-
+
/* rr[nn] = u mod v */
mpDivide(qq, rr, u, udigits, v, vdigits);
@@ -157,25 +146,6 @@ int mpModulo(__global DIGIT_T *r, DIGIT_T *u, size_t udigits, __global DIGIT_T *
return 0;
}
-//int mpModMult(__global DIGIT_T *a, __global DIGIT_T *x, const DIGIT_T *y, DIGIT_T *m, size_t ndigits)
-//{ /* Computes a = (x * y) mod m */
-//
-///* Double-length temp variable p */
-//
-// DIGIT_T p[MAX_FIXED_DIGITS * 2];
-//// assert(ndigits <= MAX_FIXED_DIGITS);
-//
-// //Calc p[2n] = x * y
-// mpMultiply(p, x, y, ndigits);
-//
-// /* Then modulo (NOTE: a is OK at only ndigits long) */
-// mpModulo(a, p, ndigits * 2, m, ndigits);
-//
-// mpDESTROY(p, ndigits * 2);
-//
-// return 0;
-//}
-
int mpMultiply( DIGIT_T *w, __global DIGIT_T *u, const DIGIT_T *v, size_t ndigits)
{
/* Computes product w = u * v
@@ -895,11 +865,15 @@ DIGIT_T mpShiftRight_gg(__global DIGIT_T *a, __global const DIGIT_T *b, size_t s
int spMultiply(uint p[2], uint x, uint y)
{
+
+
+
+
/* Use a 64-bit temp for product */
- ulong t = (ulong)x * (ulong)y;
+ //ulong t = (ulong)x * (ulong)y;
/* then split into two parts */
- p[1] = (uint)(t >> 32);
- p[0] = (uint)(t & 0xFFFFFFFF);
+ p[1] = mul_hi(x,y);
+ p[0] = x * y;
return 0;
}
@@ -1240,17 +1214,7 @@ __kernel void single(__global DIGIT_T* x, const unsigned int s_len,
//const unsigned int count
)
{
-
- // memory(res);
-
-
- // __global DIGIT_T * __local ptr_x;
-
- // ptr_x = x;
-
- // ptr_x[3] = 4;
-
- // mpModExpO(res,s,e,n,max_len);
+
DIGIT_T mask;
size_t n;
@@ -1283,14 +1247,13 @@ __kernel void single(__global DIGIT_T* x, const unsigned int s_len,
mpSetEqual_lg(y, x, ndigits);
/* For bit j = k-2 downto 0 */
- while (n)
+ while (n) // I think it just goes the bit length of e
{
/* Square y = y * y mod n */
mpMODSQUARETEMP(y, m, ndigits, t1, t2);
//mpModSquareTemp(*y, *m, ndigits, t1, t2);
- //mpSquare(t1,y,n);
- // mpDivide(t2,y,t1,n*2,m,n);
+
if (e[n-1] & mask)
{ /* if e(j) == 1 then multiply
@@ -1298,8 +1261,7 @@ __kernel void single(__global DIGIT_T* x, const unsigned int s_len,
mpMODMULTTEMP(y, x, m, ndigits, t1, t2);
//mpModMultTemp(*y, *x, *m, ndigits, t1, t2);
- // mpMultiply(t1,x,y,n);
- // mpDivide(t2,y,t1,n*2,m,n);
+
}
