libgpuverify

Signature verification on GPUs (WiP)
Log | Files | Refs | README | LICENSE

commit 6fb6e8a74ab8fa368234ab28da2089968fe3263d
parent 194e79eaa0b5b534a2cad45c35102e11c17ea6a4
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date:   Mon, 20 Nov 2023 18:02:24 +0100

Keys and Signatues can be read from files

Also, the kernel can handle either one key for all signatures or an individual key for each signature

Diffstat:
Msource/rsa-test.c | 209++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++-----------
Mxcode/.DS_Store | 0
Mxcode/lib-gpu-generate/main.c | 2+-
Mxcode/lib-gpu-generate/msgsig.txt | 48++++++++++++++++++++++++++++++++----------------
Mxcode/lib-gpu-generate/publickey.txt | 2+-
Mxcode/lib-gpu-verify.xcodeproj/project.xcworkspace/xcuserdata/cedriczwahlen.xcuserdatad/UserInterfaceState.xcuserstate | 0
Mxcode/lib-gpu-verify.xcodeproj/xcuserdata/cedriczwahlen.xcuserdatad/xcdebugger/Breakpoints_v2.xcbkptlist | 432++++++++++++++++++++++++++++++++++++++++++++++++-------------------------------
Mxcode/verify.cl | 28++++++++++++++++++++++------
8 files changed, 500 insertions(+), 221 deletions(-)

diff --git a/source/rsa-test.c b/source/rsa-test.c @@ -149,11 +149,120 @@ void generate_random_pairs(DIGIT_T *bases, unsigned long *b_len, } -int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, +// returns how many public keys were read – either 1 or n + +int pairs_from_buffers(DIGIT_T *bases, unsigned long *b_len, DIGIT_T *exponents, unsigned long *e_len, DIGIT_T *moduli, unsigned long *m_len, DIGIT_T *signatures, unsigned long *s_len, - const unsigned int n) { + 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 sz = 2048 / sizeof(DIGIT_T); + + 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); + + // printf("%s: %lu\n", n_buf, n_buf_len); + // printf("%s: %lu\n", e_buf, e_buf_len); + + DIGIT_T exponent [sz*2]; + DIGIT_T modulus [sz*2]; + + mpSetZero(exponent, sz*2); + mpSetZero(modulus, sz*2); + + mpConvFromHex(exponent, e_buf_len, e_buf); + mpConvFromHex(modulus, n_buf_len, n_buf); + + unsigned long max_len = 64; + + 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; + + memcpy(&moduli[i == 0 ? 0 : m_len[i - 1]], modulus, ( m_len[i] - (i == 0 ? 0 : m_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)); + + i++; + } + + int j = 0; + + while (1) { + + char m_buf[2048]; + 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); + + // printf("%s: %lu\n", m_buf, m_buf_len); + // printf("%s: %lu\n", s_buf, s_buf_len); + + DIGIT_T base [sz*2]; + DIGIT_T signature [sz*2]; + + mpSetZero(base, sz*2); + mpSetZero(signature, sz*2); + + mpConvFromHex(base, m_buf_len, m_buf); + mpConvFromHex(signature, s_buf_len, s_buf); + + unsigned long max_len = 64; + + b_len[j] = (j == 0 ? 0 : b_len[j - 1]) + max_len; + s_len[j] = (j == 0 ? 0 : s_len[j - 1]) + max_len; + + + + memcpy(&bases[j == 0 ? 0 : b_len[j - 1]], base, ( b_len[j] - (j == 0 ? 0 : b_len[j - 1]) ) * sizeof(DIGIT_T)); + memcpy(&signatures[j == 0 ? 0 : s_len[j - 1]], signature, ( s_len[j] - (j == 0 ? 0 : s_len[j - 1]) ) * sizeof(DIGIT_T)); + + j++; + + } + + fclose(pkfile); + fclose(msfile); + + *n = j; + + return i; + +} + +int verify_pairs_with_opencl(DIGIT_T *bases, unsigned long *b_len, + DIGIT_T *exponents, unsigned long *e_len, + DIGIT_T *moduli, unsigned long *m_len, + DIGIT_T *signatures, unsigned long *s_len, + const unsigned int n, + const unsigned int pks) { int err; // error code returned from api calls @@ -235,9 +344,11 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, 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); - + + if (address_bits == 32) { + printf("Kernel is only designed to run on 64-bit GPUs."); + abort(); + } // Create a compute context // @@ -259,7 +370,6 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, // get the kernel from a file instead of a constant - FILE *fp = fopen("./verify.cl", "r"); if (NULL == fp) { @@ -324,15 +434,30 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, 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); + + + + if (pks == 1) { + mod_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(DIGIT_T) * m_len[0], NULL, NULL); + exp_mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(DIGIT_T) * e_len[0], NULL, NULL); + } else { + 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); + + 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 { + mod_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long) * n, NULL, NULL); + exp_len = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned long) * n, NULL, NULL); + } + + 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); @@ -344,17 +469,25 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, 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, 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,sizeof(DIGIT_T) * e_len[0], 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, sizeof(DIGIT_T) * m_len[0], moduli, 0, NULL, NULL); + err |= clEnqueueWriteBuffer(commands, mod_len, CL_TRUE, 0,sizeof(unsigned long), m_len, 0, NULL, NULL); + } else { + 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(unsigned long) * 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(unsigned long) * 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, 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) { @@ -375,6 +508,7 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, 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, 10, sizeof(unsigned int), &pks); //err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count); if (err != CL_SUCCESS) @@ -403,6 +537,8 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, return EXIT_FAILURE; } + printf("KERNEL IS EXECUTING..."); + // Wait for the command commands to get serviced before reading back results // clFinish(commands); @@ -441,35 +577,52 @@ int verify_with_opencl(DIGIT_T *bases, unsigned long *b_len, } + + int rsa_tests(void) { setup_gcry(); + + - int gen_n_pairs = 16; + unsigned 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); - DIGIT_T *s = malloc(32768); - DIGIT_T *t = malloc(32768); + // MARK: must be 0ed first + DIGIT_T *q = malloc(516 * gen_n_pairs); // does not set memory to 0 on linux, keep that in mind + DIGIT_T *r = malloc(516 * gen_n_pairs); + DIGIT_T *s = malloc(516 * gen_n_pairs); + DIGIT_T *t = malloc(516 * gen_n_pairs); + size_t *u = malloc(gen_n_pairs * sizeof(unsigned long)); size_t *v = malloc(gen_n_pairs * sizeof(unsigned long)); size_t *w = malloc(gen_n_pairs * sizeof(unsigned long)); size_t *x = malloc(gen_n_pairs * sizeof(unsigned long)); + + unsigned int pks = gen_n_pairs; + + + pks = pairs_from_buffers(q, u, + r, v, + s, w, + t, x, &gen_n_pairs); + + printf("--"); +/* generate_random_pairs(q, u, r, v, s, w, t, x, gen_n_pairs); - +*/ struct timespec t1, t2; clock_gettime(CLOCK_REALTIME, &t1); - verify_with_opencl(q, u, + verify_pairs_with_opencl(q, u, r, v, s, w, - t, x, gen_n_pairs); + t, x, gen_n_pairs, pks); clock_gettime(CLOCK_REALTIME, &t2); diff --git a/xcode/.DS_Store b/xcode/.DS_Store Binary files differ. diff --git a/xcode/lib-gpu-generate/main.c b/xcode/lib-gpu-generate/main.c @@ -134,7 +134,7 @@ int main(int argc, const char * argv[]) { gcry_mpi_print(GCRYMPI_FMT_HEX,mm,2048,&m_len,m_mpi); gcry_mpi_print(GCRYMPI_FMT_HEX,ss,2048,&s_len,s_mpi); - fprintf(msfile, "%s:",mm); + fprintf(msfile, "%s\n",mm); fprintf(msfile, "%s\n",ss); gcry_mpi_release(s_mpi); diff --git a/xcode/lib-gpu-generate/msgsig.txt b/xcode/lib-gpu-generate/msgsig.txt @@ -1,16 +1,32 @@ -436FB24308EF665D:7AE29746A8D2E8F7BD0A46CE41A3B6EC41FEBBACD631222E3FBF8BD0581AB405F193425670EF88CFD55159812F328F10712788A4112F5142F07E58732005E33ABED94B52B26E196D430E32AE250A283BD9ECA07CE111B872C25EA9763A8FB6F5CBC8C1B035A2762CDC86FCC57EE335CD23B5D32D78175C0C897743E501F487001E086AE3C9B04FCB2C88356D24067C609A999A08FAE659D801A3366ED27D84C2824533378552CD85225C94B1325818A4ABF983D97A3684C2440BFAFDE209100B48589C3467B45DC282984774E2747D96DE5298623B6596B9030BF35CE3096A74FA725DD4053EA341381F9D0EBB5D74D0830D70F72C277069606E83724A725B41 -00FF2397B8AA99F05D:00B7B531089B254E648B1D22FF42DE5405FCFA93C54790D4B8D6D3EECAFA101A7F8AB72FED1C782A6B51B13DBB489976F26350897ACE99D94B0C671E506ACED56C9D39D21DFDAB9ECD5B65082353536FA167935ED4ACF434E98C4A7C49DA34C518BA74C41CF75FACC4C21435B4A40507FA22968628EB579E7593828FB55279DC0046A56AC27CD768DDE481D6E79C865F13E6D6425012F361964C718E357B28911613633624D15493999C25F0D47860007BAC89403D1251008216DB03AEABFC84A844329F7E14D2BFB835A3324C2A622E4CA08AA7C1E69C17EA7F4931572A0AF1506CA255A5B02BAC03F6A64EC94FFB0B2F9AF3B3A81B349501B1474B2A4D8A0352 -5F18CF0FE3D498B6:6CE98C13C69D3084A4A5C157FD7E55E7DFC5295B95F46E8DD58B02470D6CF190F6D54298EFDC9CCDCC35C1A2F515E5342E38248E03E9967AC1F4AAAE8E5365C58BB0A1ECAE857295D80C870305A2920A4BC32A61AC614AFCBC14F6CBCAB6F4722F1D41F0FE0EDEDE819964F4B7368EDC8472EF8C70EE2B2361A1C2F20226FDC4172EC0253E165587952BD277D873855C826A0FE4D7ECA03539FD0765269DE20DF99715267F51935510247BE9820E31ABB56B316AF77D6636EBFC13CDD9BA9219EB50973B04A12CDF87120EC77C41C02FD6CAFB20FB8571670BB4744059979530E28786B7A6E769C082E0D7A8E5FA4B3C7242627F91714EF40A129EF0A757EE37 -7E5332A3B84EF33E:00B9A48FA8058A2391BACB0413BBDEA70021476DFE960AA11E76BE0415CAE055D4E25DB547E1FA96814D3F1995609F1AA176D71708067533500169FED35E834E417E082E91DB24FBFE7C06CA12382DB1205DA0A525AB7DC0638B69EE93A59C51FBF38A2C89ADE05B166956E2F7799390E8B2DE1C0574622619C57C8D88EC8B0B91C40126356A21990E92C7732104E9CB13E70B1ED28221BED853596B7DC62D11D981D634C6272B27D15998F1260DB9ADB9F3D3AC7DFCE9675E421A0196A889144B8FC066CF17EA408ABC644E037EE47FD15D423A62BAA7638D3B8B951CBA42CCE2B8CD1807A25A64233BB92AC26EE6D606A2A74D2C98475F33C45902E2A6C7B80E -37EF192BE2750051:77541F192DB5B755115ADEF2C8C14F8D2712D3087D790FB0F5D134C7DA5BA6DC661156F9971586610F2D73C9083D6DB6E8978410CFFC0454CB7BA219D3B928E5A4779ACAD4AF72CC1F8557D7D649745E2F6BAAD06957AA34C7CBA164ECA9907CFD3BC18E9B082056424E831E684262A65B482AC53DFE65AE189F1C3C3B42673C1C22322C2F9F97B164EA31026BE711BE9AF128D586BCC3BFDD28591AC09CBF885A46583D13390EBE02AB0CE4A827BAD8C5A3AD9A5A2AB975CA6B2AD0B58F35A38604A598728F626E68236796833E72015B098B5D1F029011E3E19AFA7FFBEF6BB4D5C0C96D869BD6DE96EE18CDC1F0C13A30D6C6367F90D37AC88BF09BF1945D -00C4C187289937B8E3:27696442A4E507A30DB8234029F9ABA46124376051B5DE6B67D1B89F00FA79BB2D57D1A05C540351521B8ECA6A7FC74B1C7689A3706E08A40CC1F1EAD6886FB657DB5C7A28FBD331D524176007F88B2FE59F87CD5B98CE9FAA843BE415AD082BEC45720BEC97ABC9A46CAC5387F03121E8F582A06CD55F534B056E271679C1102EF52867393A46A0C4F6579945C4D338EF123B85A8CAC6B06578F32EE7B66E38240A005B95E3F7C637B076B955EDF9967D152795188959A733A6F225B4245173AC629EE3F71334E0CF5C98238EADFBF9E4F78DB9FB593EA5546378FBCF4902EB1A6113AEDB2E3E913B7941EB8C6759F32DF5FD0DDE50491D40D90753050641BB -0E98251C2BFEA0F1:6D480050E47EA73C1B9D8C625BD149C53EF1A6CF7ADA24CC177DC29712364EE28D3A80DAEC25BD4F3A6854AB5ECCE56F0F66532FEE9D9A728B69E9A60DD9AFD7B864504D1A2E41414EED9BB5170ECC68E094FF0D3F52067A7F7C37221B71E24D94D113383BEAB79BAB01E5BD89473BE3D8610B429B48D8867679DCF35E42D1FEB1D0CB4E2970FC200F6DCF1CEEDF2A46D71FA298E2C752C65AF8516F7EE03DFB4C386D3E6BB96EC8848BBCC8003777BA0E1AEBA2AEFFE5A2C1D8D2FB8E18D4D45F9A55834C7B0DED72DC8732D7B54C72E391E8E4E94017E522B6D38A99EAA1362A95539077356B304D67556D376FBFC06CE5E2D165BD63B0EBEF748E11C4B7E8 -682E0844E409AEDC:009700455FFA6C25BEB00A31DD1613CC94AFE28D339EDF703C235463158B5D40B26A3AB96A20872ACA98D6A1746C2CF4C69BBCCF93C18EC4DD451B80210B854F85722C0E8476C8460DA02076BB097BD27B77A63CBDAE69597FD7596A235512A1BD72B616CAA464EF872EC0812C2A35A298D5C18EDA76D3349BDB76B18CC9EBC2510D30998E203108F395060D1D1DED1537A2B9AEEF8AE4F8D077F6DE9A82CFFBA8CD1171B5DBD66E5CBF39320CC04F54D51E0AB2819C74C69F749153C93F9307FB274B79C66982CD57D41DBDC9A1FCA2E15EE0F6557B25681DE223A9BF596739D13C94F3893E66657BFFFD452840558C9857BB7725D873C093F1E72F6C34925F46 -76728A4999496B9C:04B793AF801AE18965DD0E37F21100950384B5D8059FC526F1143BD6BC182940F5E42392CC9534E8DAC1190E966F10DB619FC4F45B2AFFBB415F733E7CB971A5DDED0737A9B81646AE41C6742BFD6DD97CB1AEC7BAF923FF237C329FD4E307906B6ACB92F084626AC1E20FFB2BD780DCDFCD7460092E1DAF5AAC6CA56BC1D265B6BD02D1C45E868A289E294A2AEF784600D24EBD108686CF6AFA82D39A69DF508E0F28C7D31D2FBB3BF6E2D4226CCF197E92B02DBDBEEC04B03CD0D899997387D8CCE04AB41C02C4967FFB0CA19E8864360BD12A9DE5703F38B394C2EA8B7F6294823464911D479E7700CC6B68E07F23E9D69A9405E93928258FF771DDF7657E -00F00C6DB96830BD0A:00930DA0FB9918FE343B8CC5E77258FCC75860D15E6A6B72D3B955FB39C62162EE3638E14EAD1473D81D500F89FEF9F98CC0039271290775689F8728E87D93E9654E599E052088D35654DCCBB4E199EFB4EE88B2681CE79AB8D0E75EC07C25865173DFFB8865610FCE8B9450FF494A8CEC24BE3CC2EFD22FAE17F357663A3FF7C279175A9BDFD523C267EA6B24210C434A7682D25A38FB3F25BAB62E44D792AC4D0848DF462A68A3953380B6F67235F65FA68091893D7331D4C7DD93F79DA572405013ADD7D8FF28ECA4BF4FA97DF37E0D16C24F6CA06E952C722BED51D980B6B88D2B812CF88EC9F0BB717E974BD62076F177B4ED9EDDDFB479C590DCE2AA3F87 -046C3C377F612702:647D3FA8F730455CB553CA225891FE8A9C05BBB5D5EEF8A93536ED579D79C11C8B71CC2780E28718553FA2FB7EE7749624D98421D4F7125D941CFEA3A2D2BD5ACBE68970BE222079EF3EB7939C0E5A2EFC921E2261E35C93CEF2544EF8AF811D4D3292B488C6067B1B4F8C29DA303DE590A2264B072FCDCDB333EBAD51B8792883068436CAD86A1B3396FE371E928C7E854966FB8C021D4B0086EF718E612E2A926C321459EAA9E645F378E761BC467CF090750C385D56C2E3287BCD6646EE1F07DD9472980A72D22F82FAD5E8FB0E0B3E087E295A6EAB1BA22927B1E2EC6E213174C729DC00F2956CDF9E30BC2E17AC037B4FDF705C52F405473B6D61957D3A -00EF8540C5BC629D9B:00B423B022FAF9319FA8D7079BA64EEF3370AF01603D1122E742385E60D6FE6924A65412EF46A867076859D5DC32647F995933505F3DE83F7D8F2C7108AFE6D00DA0107E474729297BEC813F2D529A501D7ABE360345F12094BE558369B3CEA9BD5B61C990F5C6E0A7B5F5B64896F5DB75D4D0CE934C3E84E098092470833D518F6CB3779C199F7CDB2052561FC5B6F86B44299BE6C5420C199DA4F46FBEC220186CA086FBBC4E629AC7A2A0B637A440A993C96BDFC9B32EE00954C92D9194F079DBEEFE8B4BA636384EAC2A166408EB944DE41FB304837F10A5037C78F983E9B5DB7B44AB1D619825EC444815E169A904D77F7E9F74BF356097C566447BC3FD9A -59A5EE5E908DECF9:568344411E8EEE5E4B86EF8B468FDEFAC49EDBFE2D00F3A148FEC5D18B709F13E90D6936488F966805933C349DC5E61C12AC2DCE4B573F46D9F7FD85BBB4A9AC5996A83356D659DD75892D4E32272968B4F2B05944DE54C81C4476FC88C2C2A817F616D671B78EFD940DAB906B76D2695E8E582E3359031176BAD71C69831FEE68927A37742BCB707F6EC24B768B1BB325CF5B57FF239E263FBA35DFA65EE729E8D6FA0615A7D003F12476069569AFFF50BE05AC737EDDF8086D52E2801B7FEE9927EC4D60EA351E08EFF284FDBC66AD5AD0CA48B47EE187F41ED45290945D960436C3D581A9E05950F840F978854BBBF45492D408D6535D42E495D3D2D262D3 -33657640176515E9:663EF3A5401E0ED033094AE4D1F45E852742F7426273A03BB43FC6765EB722A8BF9802A042906570528613527B58B5B603851D5F5A73009FA16EC3BA35B7E61D9EE2E40653405151CF26B335C20E9DEF0EBFE5768DB4155D84B380E58E059CCAA526E3F2EBD01DD197863B715B9E6855D3A5442E78244F50049100DB9DF785C2418B0C6814B75DEE956E67E79D966641DCF81FF76609C79BB7432509E2777F74F52ACC239892654D1B7EA711133C28C23FF9167E13B096B41C30673B5049F759B389A460D9FF617CC539FC63EDBD1E3FE6D41C66D628F3FED99FBC159306CD879A0281C6B2B7DD0FC70AFC2ED7C68E1BA6C74B4FC8BB2280E6544556E31D564C -32C660AF720F1228:00ABB2A2B45B48D39E86F6FCC72974772EA337427AD64A8BDCFA625C938B2B162162AA9295527DFF6D1CD254966FD80FE014183E52F9A3DDBDBFED250900BD979E37AB158507CEEE849934D8C64DE6F064F300C7571335E3280BA05F4D4640141DBA54765B5643CF01FC3D99D103A87F03EC006D3499DCCD0CF4D553D004AA17AB8A17883B17E305A63540708A54507021B5CDBD8689AB2D2F273AD8BD5BD3A5E45BABD46A052153C3280AEA28297153EB079247D78B37C3E6251DDC80FD866FEB07A6ACA4098A34C25956D32F8947CCD94FB58085F1CC46B63ECD98797ACA16C8CB3B6C6D6953B91B0AEC958FCBC09201F4073C815AD539579B003D1F1C8303C2 -49F8BE8FCBD9AD8B:008EDA31256FE91B1C76B50413E3C0D934E4EBC83B0352EDF37CB279B835F57754FE91244D62F07A16C2B56C6ADEDAAAD8C07CB111846FADCDFF7EAE626F81933C4EB1C42026FB220B175DFACA9170492DB4221E4864FAA2D7F247D5052CC3E0EBE2B298FFFAEF84509822A1CD3B3B24A7F569BE7AB18A036F78A926605FF51F95F3DAEA40856FA1E3F6C0F9CC0B6B245545566CCBE0E6EBB8B74F8319AC3B9EA06FE31DAC0DD4E1289DDF1F5111F97D51C5A6E9652974F56D0D253FB1C7D8F67DCB52091C62D59458ADBA524FA23306BB816EB863AA7F4401B3DBD81197FC1C0D021EA04128EF02FB1BD65CC604BDC49F97B3794D660C2CE0E8F1A17E6BBBB51B +0090D8325E43E77A1B +009A2E7A3ABAC99E3531FE4DD97E8823EA04FAE925BE05F8DE5D87836FEFFD7D28FECFA41C573BBAE99FA26333E66B7EC7B13EA893395B831221EEE1BAD36F6D75F403FC5E9F64A440A39D34E756D33302B9BA915FAA3A6D819A40B0AEF9A5AEF34C024C45B68139BCB2D4970D538E2F1416482B406487CD49A15CE7F73C7F3E49457D0F27FCF44F896C340AE6FB18E378C54B0E99C1BB0D5B86F88C9F5739CEC50B76474EE4E8D803D1BEC51BB4CDD00ADA20DC267312BB305EF8DB7D488D32D08940853FA3CEA4C5005195465FFE411F3C08D13EE3B84B0E3F436A2162EB5E373B47CE0C4C10CA9EA19003412420E8668659D3012BAE5F373720ADCC11EA1497 +0D939385724C504B +78B5660DEE2FC8104BFD5E52EDD024EFF5C0692455BEC165BA758238BC9FA05E597AE246366825A736673AB3F0922BF8E040343512E548BEA1F4AA54F493FDC95198CC082271D986F1EE86F7B8FE38C6BE91D4E772F583374A4AB6CB60A518D0EC08A6A0C7BC43CF908D9E97D4D1DD578E369505574D9BCAEAB64030F7F52984EEBEF3F71B5D51CFA10989CB1C1E2250B6E9554A5E5F402FB845C2271662FED518FC8356B618B3BC299B0614E806DAE5BB84958186E9A1CF94E434D6F9135088037BA0FD1F81576B00F4F26648E15AA113E8A500437DEDFFE7D68719546A12A4BE94072D746CFEC4D320099992FB831B5E57699E71CD439BB06F7FA5263154D4 +329BB45CE74502D2 +1E6B8D60EEBE802A3603E0EF51622A78E7A0D0CBFAF932B953A0EC8CB0ECE830495D8A6618777B5439AA558FA2B9AB1F38638197591CADC929B892F64E129F5D584F37515BAFB0728D57B77F0136AE07F332964E6F53DA576859BCA785349291A62EDEB708FBA386E302B07C06A8DE9D83C717B14FEDA82CDB749B2E9E622EDBA67501BDF1B05D511D5EE4B95B9B764F584032ABA5685ED3680F43095436951596D4190FA93A345FCBE29599BD073990E412902B98D36122CF70BABF0D96FE35C818AD175E667CA6D0C194ADF7BCE6ACD6B948C10D215E3C9E901511D5318997D8E592B2627DCCB7D897CA54F42A1418EDE4A2C510BE4B83246A5BC02E414A5F +4690F9B10DC84BA0 +3833F1D92E17CEE8C71B29CDCD9611120CC20DAC979536075756F4B454D36E931A74DAF91B9F09DCEEB8DF1704F78122AF1BFA4E512DE3FF6D2245CD332A7713FFEB68A4150775CC42229A0EB0E24DAFB141D9698B397C0A7879C75975D2743A0EDE9F26CBBED21FFA121B179460DB81CA73EC85FA3BFAB5C0C5E058A2F7A1D202C242C6CAB84ACFEB6FC8C20F3B65BA309741127FA3D03C285FEEC3A20854F8B70A2FDEE253A5FDB2EFF61082CA0F630312B7959AC4E38AFA0BA82847450A0F7595521FAF14EA163C6BB526851922FC9959111111ECEA1C7EC13F86AB49D3521A061D69AEB40E3E5C6ED9572DFAE4E9A08135D23B4293AE11D940116A47FFC8 +17507073E32A3562 +121B51C2C8FDAA14BB2157C62F57B137B3777C7034F58B180437CD481F6183FF8BC76B631C0495C26240BBE532B6BC6A68C263B5593317AF35E45A03DAFE8CB154DD14EED8DCB6AB558BA001E5B4751383BB123578FD6926449AF09C9D69C376EFFEA35AB16D866D0A39DE16DD400859AC642BB5BADEFB9895FFA778DA0885C15FC0F1B3D6747ECE8EAE00F395CD42F96C4A7C0E757C523F3E0122FE2D1B1991B8894E16125662C51C9A75F64D88658BA6B52F25795753897555BFDAC950995E26218C884474E1AE74C9B269CA8C5BFC5E4B3FF23041626948125316AB6F0BBBD14C26D939F2395342FE1D4A3AF4857D39B9267C94E5C50F5424E17D882EB5C3 +00ACFCF8A7DC12E7FD +009149648B0B020681C7BB7966D9A5D95F9B96EDF74A9748ACA85D4D3463CAFD0B956EA115B0188C8AFCBF3A777041026C3AAF8312C65EAB1563C31F777D4A370C1E3A3690B94DD1DF44DF64E0D4957C8ADF5C8708DEE96DF601476DFE751ED209452AF612FE065749CCA16A9876B96B596CF24826C619706AE7FC31C06E7B726A3C2CA9F6CD278C68209EE9BFB40A6599DC2137099A9BBDF8E434617F1A8048FE82CC6990710B7E347C3FDEC7766134F5F2B4D1B1BF8CEFBFE8A902624A50EF90FC4915D8F443BD6094DBD367F81C3BB6043949406B783A4D23FD2732568F50490F4878AE325C60891BE2F18190B37E62BEA9C88001A3A1E66D133FE6D9E59EA4 +0098C4901E4E475AB9 +494D83F5AB6DDE76517E2A610F7F7412FB0FC44B2474F3BB404102813CDF7675CB16F4BD193B3B99EB2CAD063D381A7556C389A9147A4338185DB0A25F80B8F4C8A2247F1224AEB08CFDC3FDF47A32702FAD243719466DB3B290B4250CD98C7F855337BA8570A095556BBAF0AA9AF04F7B70ADE40A8C785C7243288A584154601417E6262276A291FAF715AD89ED01A9E102EC81C77A5770B14D731D32F796CF41ABB9D138AE44B09289AEF28B126B4D61B4BE58DAFF139859183856B6CD7741B31364CC62906023DAAD6AEE39CC310FE3CD7769E8835D015B82C9D8DB5B9294F52C88DFE0371BA3EF1A5C75574678AA66BD55A4EAB3867872751833785853C4 +00DD4ADC302D41DE2A +7671B379F7E8216D159598B28F78055D05A1313E42446F7B818E0C0E98B9B5163127AEDF2A35EEFD27FEAADA7A5816AD1468C113EB69E0F8E2C2690999C6709174A70F56D8E10307F5DBCD1A6745B057C29B9F5502FE1E5C8E0CE82EEAF9FD53A144450213BD94AEC27139D4D0D8057EA2266B8CC472E8FBE5818660905D8259609EDE3E0709243867177894C5B99208C1C9640DD77771668631446373F0F56F47C6994B27E9E9AD77EBE14FD328656F1103D915FD5A8C35CC955EE2AEE8DBB8D7C2A93440FC02AF8F908D9803EC0FF0B59A2A37E264B5C439AE34E0BE26691CD19F8E508B345843440332441327D32B9A8F984CAC48B83139144ED183C8D537 +11BFF9FB7C1681B0 +1D65893019C63343A22038380E1553698AEE82E895545AA24B8E6EB3F6FBD8ED2FC940C3EA33FCA6F990F84C14A0145D3DDD42E9E30FAC93D3F711D32AD6DFFE3DAB594E3666E3891DB9007569E998210CC7B2363338D8233004AB8D69EADC5DC68B60D398F803F6BBEC24F1C6458479DEF087E6D0CE5C659E2F94DBAD85DC6895D7ACF595B3D3F90B1186672DAB0CFAEFAFDDFA39EC649DA18830CD2EF2FC49A3B5667E462B03599D3DAA7B06BA663985C2BBEC53E9B71E546EAE199D806CA41E939B1F3C149799DB2131B5B88884F86A760429485A40C9F6C29A1ED5CB28BFA8C48E66E69A10925FEFACEF3F5A3B244DB9E24CD1D88DAA20E76A3E60DD00FA +00915C5BCB718411C6 +0087B8189FB165EE2E9C1CA6ED4A7C0D98F33A4CB6724BE52E60C53345C546B22205F462FB803DEAA817146E3060935F015509A3EEA30491A92CE6B970980A11E07E5F7ADE2FFE89F39E239F7EDABE36D6DD3D52ADB1D7595962E074B0B08E8E097D34E862D4BEFA090EF94B6ED9FB1A7815399CB2DE82EF9B18B14D07DB4AE96B3ACF29A98B477B29E267F8249F998AD07C1AA59FC67423B82F497DA2486BA02223228BDF663F9274951F7974C3AA8A3285EE5EF2D07D9A34E716A62CF111C683E8C889BDCBC5BCE685AFFEE161E36B974EA7F2D6A24B30077665962D11F23B6313EE53352384559094E60FE7D12F083E84A8C1400FD08F03DC290E095B3160BB +592A6A94F317309B +7A7BCE2335533AF60929795A3DE8A56BE1CD0615CA11005625B216014908130C9F8130202975BD5930B5CC1346ADF21522579060BDAE4FE6765FA278B469C0365A0C514772D4E2BAB4A8BF92716840FEEEB8DBDC00A90D10DE8105D1E6A5FBEEF84706295B02ED0D96D35FD80EDF6F0AE35AC58D988335D1AB229E93E502B4A9DE723661D9FFA923E48282F2B74E6115F02E6B5174B7AB79FA4B3532CBC27CC3677833ADD638E32B7CF2C6EBE104A174BA94E508E9741285DB9BE9EA70076838C3AFDCD71C81F171718055D51A714A10FBF0AC16B8A7F2D0114A9FC1DF65CA873D472779057EEFE92C192AF2728A7004910C756C9BD93028C20BB179DDF215F1 +0E9D0414A05989D9 +5A2BDC6DD3013DCE4E15FE50424CFC275BE3518F56A592D4290FDAA3058252B51010CDCF6B1455F45561CC51B916A569BBCC7EF84B973B6C1C4713E3AF0B2EC18A3A6EE6067C16131B73DD279B452F0FACB256261D9C1464E20A372048F51E881F1F1B9B828C31D9B1EBE05EE26E59DF68CF8149326A0D6F70B8A3A7BA6D993440DA3A6C74C745B59F4DB6B4542C8E51113817D7A92EE8F1E7CFF67607F97F6B9DA318247EF9EFC484BE7B8EE58B3460889BC490B0153306F69C6809792656E500311953DF90EA44DAC1E1C4E8189E83308D9443B8EE89009B0BB214B4F80E9817F217CF7AA071B875CC784B88866E040FD04A450BCF5BB0125DF185D579AAD5 +0096AA789DD034DB92 +387AA57B2027DB04CB8C7A3714C55228F209026EC3442AD6FBFBFDC953C6E1F3782108E1AB6FB92200992FFF13F5B41B5DCDE57B8E0C4D1168C65369BE641231244690CB5ED6B121CBEF3D08D7C9376938211A67B886AC35D35DF45F4300B3B090BEC933005DC7DD7B63FC2558B135033DB44DB7191791EE811F04F4C8830956DF10F9C23BB8E867FC09D62953D9804ACD95746ED37166897AF0B616B4E5F68F8A2C8D3721FD1EE2087D6DA2C294D7F76D9AE275300385C92D03B58E67DDBAFDDCC4982A2E3EFA2491BFBF110CBCB849234EBB3B8BB51BCCD2E9A9F8285BECDBB43B23FE58EA6BFDF95E4BB570A4FEDB184FC70F78A74DFE257E857C9FCE2772 +0098FF0D98581E7B2F +26F5D1CA1A9106EB39ACD225B74DCB51FCBD200237B8082D3ACBC15C5BA83021806941DCB4BD97BA58D1956046B52EFD7D8023A187F30CC46F0AF56ABDA373CC55DE92C70896A4B7EEFA97A4FC4FE56088BA4E73450616A2F234BDB3039FB5D7C0475EC6558ED9732D2CEA9C7DA5664E2218DE210CFB403AB4A74459ED2ADA3273362632BCCD5D6D3A36837510F28CD0B14C396CF3FDD1ADB0FCF6071B5C71052F19594FBDE997BA166E5C445FE2285EBEE4773B623FC300EE0C370D60E595F638813DF105481BA2009D386E59502016768916A9EB2982E50639F636D0DE7FDB5EF6613711796F62ACE42D12F0900C445280B55A81A0B61B572EA34D1C5CAFA0 +3012F2FB2F58E6EC +71BF267C57B5847244A217B02CDF8B5595574EECE917C44F0A9B6723A2945AD138644A25CFD89F9700F337AA42C56AF10CF968976AECC5D5FAD438F0B86417208511D43AB73BC45C8622AF5409B49184346A667E9303FB872F6D20C2C2A72BAD110CB3422E33782278D452A48B7985FC67CDEF2EEFC59BC25C3F4614E66E8CC82207574BFB431B7BCD17CC9D75D47DC86DA6775DAAD045B9DC6C5B80003C23FB766FF14393EC38A4E0751E0369C5B3D416E5648FCD715C2E1850974A218267410B7454AAD27E9B9090361EA3DC679A2CFA93EAC15DA4B218C6AE04D79EA8BDFD9EAFC9B08BC706F7F9F0C19E2FAE7A771F02195FECB61A678259081666CF066F +3C2BDE5A818B5439 +08191FAB1992DB18592519FD2F36BB86034953AA143816C93699A34BE405A7B0183691205636E69DFA4A872E1000878226F4E1A0CF030023E09BEF303240174EDF7CCAABE052CA0EF397CF0AABCBFD2574C93C39B21659EA0CE4BA9AC6BAFDEF0B7C3DB4AB5927F9F0C2CE55699B4E636AB484CC68BA5F4892012CC36992D82CAC33135265133F532C7330871A32DEA89E3363D5E1E65CE047CF296A5CA840C93B86074111E1CE1E273FF23C5F5372AAF75C1D448193E14DB99BB8BBFB08FED72DF19AB78B1CD030C3679E581F46DCBBF4C193BE00911C974BE03A51F2705E59E4484C7BBA118988FC5478E75FF2DBB247D49F4F9A706D5F0FD73C0AB68C03D8 diff --git a/xcode/lib-gpu-generate/publickey.txt b/xcode/lib-gpu-generate/publickey.txt @@ -1,2 +1,2 @@ -00BAF4220C41EBD7F0CB74BD9914CEF0D0F3C3C743B8EFA194CA13E25C387EFCD55DD8D4A18F076420D24151D72A0F1584FF1B58A7F6A39C7B0E862422A6BEE695C84C1769323331749609020D12E163D7CF890CE88E1EFFDFC8DFF2B870D8661FFE451DC3BF945383C8CBA9EF46202CC66CE3CAA76D1BB56842101957421E1D6F4023272E05ABD108F38C63A0F7BA13E1357E83C8BD72F8CD00755ED51FCDC745C3A3CCCB224FC2FEFED9E88E8650FD50B284AD3779DC5EB046826B21F03F4F5B41E97AB6C797DA6C4C0A7EE444779A53C0CE64B97AD4093A79779D879389C36E55A5E0B34F5EF031A20201D78FC702C54CC18AD85C2844190791F58970D79C5B +00C1D957F4DFBBA9A0D24C766F54CBC6C46E868CF9317E398690014454218004A68EB8CB69D6C819B60CA2097094BE7B52CAA39B1F7267940151C94208953E9AE244E58F4689994D09D780EF426718495597BBE5B8FAFDDFCA158FD3594B03C0B8FAF8F9D69E2AA598AABC36FDACE52EEEDD05ACFF0A368B68302EC41F264D3C49CB0ED468D10683EA43D80A19AC113753B4561323338A6F9C7945BE88D1864CCDFC6A13B0F398F83EBE7F60D32426FC1B80076ED89B2166215DA50543A64D21E6881B18F1C2E7C54D7DDF034293D96C755B47AAE86A4A4291DBB506F25BABCDA02BCA79FF31BF4F1142D43D08BC94CF6C6572CA2818221A326070D93F1F54CCA3 010001 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 @@ -1279,8 +1279,8 @@ filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "448" - endingLineNumber = "448" + startingLineNumber = "588" + endingLineNumber = "588" landmarkName = "rsa_tests()" landmarkType = "9"> <Locations> @@ -1462,9 +1462,9 @@ filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "435" - endingLineNumber = "435" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "571" + endingLineNumber = "571" + landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks)" landmarkType = "9"> <Locations> <Location @@ -1673,63 +1673,77 @@ <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "AE6F2C60-36E2-4F5D-94EA-115E01CF5285" - shouldBeEnabled = "No" - ignoreCount = "0" - continueAfterRunningActions = "No" - filePath = "../source/rsa-test.c" - startingColumnNumber = "9223372036854775807" - endingColumnNumber = "9223372036854775807" - startingLineNumber = "407" - endingLineNumber = "407" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" - landmarkType = "9"> - </BreakpointContent> - </BreakpointProxy> - <BreakpointProxy - BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> - <BreakpointContent - uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE" + uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF" shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" - filePath = "../source/rsa-test.c" + filePath = "lib-gpu-generate/main.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "314" - endingLineNumber = "314" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "123" + endingLineNumber = "123" + landmarkName = "main(argc, argv)" landmarkType = "9"> <Locations> <Location - uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE - acfd4e71f26ef54b" + uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e836" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "verify_with_opencl" - moduleName = "lib-gpu-verify" + symbolName = "main" + moduleName = "lib-gpu-generate" usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c" + urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "298" - endingLineNumber = "298" - offsetFromSymbolStart = "1959"> + startingLineNumber = "129" + endingLineNumber = "129" + offsetFromSymbolStart = "791"> </Location> <Location - uuid = "A3962F92-3CCB-485D-A314-5608CDB551EE - acfd4e71f26ef37a" + uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e9f4" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "verify_with_opencl" - moduleName = "lib-gpu-verify" + symbolName = "main" + moduleName = "lib-gpu-generate" usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c" + urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "127" + endingLineNumber = "127" + offsetFromSymbolStart = "831"> + </Location> + <Location + uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e9f4" + shouldBeEnabled = "Yes" + ignoreCount = "0" + continueAfterRunningActions = "No" + symbolName = "main" + moduleName = "lib-gpu-generate" + usesParentBreakpointCondition = "Yes" + urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "127" + endingLineNumber = "127" + offsetFromSymbolStart = "827"> + </Location> + <Location + uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e970" + shouldBeEnabled = "Yes" + ignoreCount = "0" + continueAfterRunningActions = "No" + symbolName = "main" + moduleName = "lib-gpu-generate" + usesParentBreakpointCondition = "Yes" + urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "313" - endingLineNumber = "313" - offsetFromSymbolStart = "2488"> + startingLineNumber = "123" + endingLineNumber = "123" + offsetFromSymbolStart = "727"> </Location> </Locations> </BreakpointContent> @@ -1737,92 +1751,108 @@ <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66" + uuid = "100A7C3C-BFD8-4C13-98ED-E1BA3A91E6D1" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "256" + endingLineNumber = "256" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A" shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "423" - endingLineNumber = "423" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "214" + endingLineNumber = "214" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" landmarkType = "9"> <Locations> <Location - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8ee7" + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A - e8b83ddc77e63c85" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "214" + endingLineNumber = "214" + offsetFromSymbolStart = "2727"> </Location> <Location - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e42" + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A - e8b83ddc77e63c85" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "214" + endingLineNumber = "214" + offsetFromSymbolStart = "2553"> </Location> <Location - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e04" + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A - e8b83ddc77e63c85" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "214" + endingLineNumber = "214" + offsetFromSymbolStart = "2537"> </Location> <Location - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8e63" + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A - e8b83ddc77e63cc7" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "216" + endingLineNumber = "216" + offsetFromSymbolStart = "2584"> </Location> <Location - uuid = "6C61A97F-A37D-44D7-9731-EADED9F5AA66 - b0b9078e770c8ea9" + uuid = "E1B7F08A-27CE-46BC-9DE5-F321A723594A - e8b83ddc77e63c85" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "214" + endingLineNumber = "214" + offsetFromSymbolStart = "2539"> </Location> </Locations> </BreakpointContent> @@ -1830,95 +1860,93 @@ <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "470BAF83-0588-455C-AE68-F686E9954517" + uuid = "154F320D-BB50-456C-98E0-9EB7D9A6FD14" shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "370" - endingLineNumber = "370" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "210" + endingLineNumber = "210" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "C1D44B56-001C-4AD8-A710-1D82ADEC89C9" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "178" + endingLineNumber = "178" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" landmarkType = "9"> <Locations> <Location - uuid = "470BAF83-0588-455C-AE68-F686E9954517 - acfd4e71f26efc6c" + uuid = "C1D44B56-001C-4AD8-A710-1D82ADEC89C9 - e8b83ddc77e639e1" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "verify_with_opencl" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "178" + endingLineNumber = "178" + offsetFromSymbolStart = "1400"> </Location> <Location - uuid = "470BAF83-0588-455C-AE68-F686E9954517 - acfd4e71f26efcb2" + uuid = "C1D44B56-001C-4AD8-A710-1D82ADEC89C9 - e8b83ddc77e639e1" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "verify_with_opencl" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "178" + endingLineNumber = "178" + offsetFromSymbolStart = "1392"> </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 = "414" - endingLineNumber = "414" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" - landmarkType = "9"> - <Locations> <Location - uuid = "9B0E2741-A817-4815-8AE4-26ED0DDEB4A6 - b0b9078e770c8c9b" + uuid = "C1D44B56-001C-4AD8-A710-1D82ADEC89C9 - e8b83ddc77e639e1" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "178" + endingLineNumber = "178" + offsetFromSymbolStart = "1431"> </Location> <Location - uuid = "9B0E2741-A817-4815-8AE4-26ED0DDEB4A6 - b0b9078e770c8d5d" + uuid = "C1D44B56-001C-4AD8-A710-1D82ADEC89C9 - e8b83ddc77e639e1" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "rsa_tests" + symbolName = "pairs_from_buffers" 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"> + startingLineNumber = "178" + endingLineNumber = "178" + offsetFromSymbolStart = "1393"> </Location> </Locations> </BreakpointContent> @@ -1926,112 +1954,178 @@ <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "164290F2-14BC-4321-8C37-498198D7FC1A" + uuid = "F9997565-9649-43A6-92BE-1BBC7B14FA12" shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "337" - endingLineNumber = "337" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "249" + endingLineNumber = "249" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" landmarkType = "9"> </BreakpointContent> </BreakpointProxy> <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "FB747E07-3A28-4AD5-9FA0-539B30DAC5A0" + uuid = "54AE6D41-A6DF-42E1-AAFE-2167C9737016" shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "424" - endingLineNumber = "424" - landmarkName = "verify_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + startingLineNumber = "201" + endingLineNumber = "201" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" landmarkType = "9"> </BreakpointContent> </BreakpointProxy> <BreakpointProxy BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> <BreakpointContent - uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF" - shouldBeEnabled = "Yes" + uuid = "6AEAD907-DAF2-4B97-AA78-487F7920EDCE" + shouldBeEnabled = "No" ignoreCount = "0" continueAfterRunningActions = "No" - filePath = "lib-gpu-generate/main.c" + filePath = "../source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "123" - endingLineNumber = "123" - landmarkName = "main(argc, argv)" + startingLineNumber = "540" + endingLineNumber = "540" + landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "DB4CB114-B80C-4F47-BEBB-CB3665D7983F" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "622" + endingLineNumber = "622" + landmarkName = "rsa_tests()" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "AB02A779-8551-47F6-8552-C9F9B2C4C62C" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "549" + endingLineNumber = "549" + landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "996DA6FD-606B-4A8B-A251-42B5EC250727" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "610" + endingLineNumber = "610" + landmarkName = "rsa_tests()" landmarkType = "9"> <Locations> <Location - uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e836" + uuid = "996DA6FD-606B-4A8B-A251-42B5EC250727 - b0b9078e770cf6a3" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "main" - moduleName = "lib-gpu-generate" - usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" - startingColumnNumber = "9223372036854775807" - endingColumnNumber = "9223372036854775807" - startingLineNumber = "129" - endingLineNumber = "129" - offsetFromSymbolStart = "791"> - </Location> - <Location - uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e9f4" - shouldBeEnabled = "Yes" - ignoreCount = "0" - continueAfterRunningActions = "No" - symbolName = "main" - moduleName = "lib-gpu-generate" - usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" - startingColumnNumber = "9223372036854775807" - endingColumnNumber = "9223372036854775807" - startingLineNumber = "127" - endingLineNumber = "127" - offsetFromSymbolStart = "831"> - </Location> - <Location - uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e9f4" - shouldBeEnabled = "Yes" - ignoreCount = "0" - continueAfterRunningActions = "No" - symbolName = "main" - moduleName = "lib-gpu-generate" + symbolName = "rsa_tests" + moduleName = "lib-gpu-verify" usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" + urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "127" - endingLineNumber = "127" - offsetFromSymbolStart = "827"> + startingLineNumber = "611" + endingLineNumber = "611" + offsetFromSymbolStart = "240"> </Location> <Location - uuid = "FB9B9AD4-F59B-4392-927C-1207B6B04BBF - e0abeefb3692e970" + uuid = "996DA6FD-606B-4A8B-A251-42B5EC250727 - b0b9078e770cf6c0" shouldBeEnabled = "Yes" ignoreCount = "0" continueAfterRunningActions = "No" - symbolName = "main" - moduleName = "lib-gpu-generate" + symbolName = "rsa_tests" + moduleName = "lib-gpu-verify" usesParentBreakpointCondition = "Yes" - urlString = "file:///Users/cedriczwahlen/libgpuverify/xcode/lib-gpu-generate/main.c" + urlString = "file:///Users/cedriczwahlen/libgpuverify/source/rsa-test.c" startingColumnNumber = "9223372036854775807" endingColumnNumber = "9223372036854775807" - startingLineNumber = "123" - endingLineNumber = "123" - offsetFromSymbolStart = "727"> + startingLineNumber = "610" + endingLineNumber = "610" + offsetFromSymbolStart = "240"> </Location> </Locations> </BreakpointContent> </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "06B1A996-AC77-49EE-BEAC-7216310A73B5" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "283" + endingLineNumber = "283" + landmarkName = "verify_pairs_with_opencl(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n, pks)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "C52F82BA-9FA2-4A77-AAB1-DAA52E57E040" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "244" + endingLineNumber = "244" + landmarkName = "pairs_from_buffers(bases, b_len, exponents, e_len, moduli, m_len, signatures, s_len, n)" + landmarkType = "9"> + </BreakpointContent> + </BreakpointProxy> + <BreakpointProxy + BreakpointExtensionID = "Xcode.Breakpoint.FileBreakpoint"> + <BreakpointContent + uuid = "CA94DAC1-5B12-4685-87CA-E53E27704AB0" + shouldBeEnabled = "No" + ignoreCount = "0" + continueAfterRunningActions = "No" + filePath = "../source/rsa-test.c" + startingColumnNumber = "9223372036854775807" + endingColumnNumber = "9223372036854775807" + startingLineNumber = "251" + endingLineNumber = "251" + landmarkName = "pairs_from_buffers(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 @@ -1210,15 +1210,26 @@ __kernel void several(__global DIGIT_T* x, __global const unsigned long *s_len, __global DIGIT_T* m, __global const unsigned long *n_len, __global DIGIT_T *mm, __global const unsigned long *mm_len, __global unsigned long* valid, - const int count + const unsigned int count, + const unsigned int pks ) { int index = get_global_id(0); if (index < count) { - 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; + int edigits; + + if (pks == 1) { + ndigits = max( max( n_len[0], mm_len[index] - (index == 0 ? 0 : mm_len[index - 1]) ), s_len[index] - (index == 0 ? 0 : s_len[index - 1]) ); + edigits = e_len[0]; + } else { + 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]) ); + edigits = e_len[index] - ( index == 0 ? 0 : e_len[index - 1] ); + } + + // int ndigits = 64; // int edigits = 1; @@ -1234,10 +1245,15 @@ __kernel void several(__global DIGIT_T* x, __global const unsigned long *s_len, __global DIGIT_T * __private window_m; __global DIGIT_T * __private window_mm; - + if (pks == 1) { + window_e = e; + window_m = m; + } else { + window_e = &e[index == 0 ? 0 : (e_len[index - 1])]; + window_m = &m[index == 0 ? 0 : (n_len[index - 1])]; + } + 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];