commit f65ca92b4717c3fec7083b7033bf0b352a64f890
parent 78ab2791f6a47cfbf4dfc3baa10382023c213dfe
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Fri, 17 Nov 2023 02:39:26 +0100
Fix memory issue
Diffstat:
4 files changed, 33 insertions(+), 11 deletions(-)
diff --git a/source/rsa-test.c b/source/rsa-test.c
@@ -388,7 +388,7 @@ int rsa_tests(void) {
setup_gcry();
- int gen_n_pairs = 4;
+ int gen_n_pairs = 16;
DIGIT_T *q = malloc(32768); // does not set memory to 0 on linux, keep that in mind
DIGIT_T *r = malloc(32768);
@@ -417,7 +417,7 @@ int rsa_tests(void) {
-
+ printf("____\n");
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
@@ -2003,7 +2003,7 @@
BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
<BreakpointContent
uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66"
- shouldBeEnabled = "Yes"
+ shouldBeEnabled = "No"
ignoreCount = "0"
continueAfterRunningActions = "No"
filePath = "../source/rsa-test.c"
@@ -2204,5 +2204,21 @@
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
+ <BreakpointProxy
+ BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint">
+ <BreakpointContent
+ uuid = "FB747E07-3A28-4AD5-9FA0-539B30DAC5A0"
+ shouldBeEnabled = "No"
+ ignoreCount = "0"
+ continueAfterRunningActions = "No"
+ filePath = "../source/rsa-test.c"
+ startingColumnNumber = "9223372036854775807"
+ endingColumnNumber = "9223372036854775807"
+ startingLineNumber = "420"
+ endingLineNumber = "420"
+ landmarkName = "rsa_tests()"
+ landmarkType = "9">
+ </BreakpointContent>
+ </BreakpointProxy>
</Breakpoints>
</Bucket>
diff --git a/xcode/verify.cl b/xcode/verify.cl
@@ -1205,31 +1205,37 @@ __kernel void several(__global DIGIT_T* x, __global const size_t *s_len,
) {
int index = get_global_id(0);
-
-
if (index < count) {
- int ndigits = max( max( n_len[index], s_len[index]), mm_len[index]);
- int edigits = e_len[index];
+ int ndigits = max( max( n_len[index] - (index == 0 ? 0 : n_len[index - 1]) , mm_len[index] - (index == 0 ? 0 : mm_len[index - 1]) ), s_len[index] - (index == 0 ? 0 : s_len[index - 1]) );
+ int edigits = e_len[index] - ( index == 0 ? 0 : e_len[index - 1] );
+
+ // int ndigits = 64;
+ // int edigits = 1;
// the result is copied in here, compare it to mm
DIGIT_T yout[MAX_ALLOC_SIZE *2];
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])];
-
+//
+// window_x = &x[0];
+// window_e = &e[0];
+// window_m = &m[0];
+// window_mm = &mm[0];
+//
// can probably be smaller
__private DIGIT_T t1[MAX_ALLOC_SIZE *2];
__private DIGIT_T t2[MAX_ALLOC_SIZE *2];
@@ -1280,7 +1286,7 @@ __kernel void several(__global DIGIT_T* x, __global const size_t *s_len,
// MARK: valid cannot be written to by several at once (the same unit anyway)
- if (mpCompare_lg(yout,window_mm,len) == 0 && index == 0) {
+ if (mpCompare_lg(yout,window_mm,len) == 0 && index == count - 1) {
*valid |= 0x1 << index;
}
