commit 4197370d2fa86fd8e7a6652f571758256c080c1b
parent 246b31b827ad8f42ff15ba55703edbfaec4a3ec1
Author: Cedric <cedric.zwahlen@students.bfh.ch>
Date: Fri, 12 Jan 2024 14:34:21 +0100
Improvements
Diffstat:
16 files changed, 329 insertions(+), 938 deletions(-)
diff --git a/source/Makefile.am b/source/Makefile.am
@@ -11,8 +11,7 @@ libgpuverify_la_LDFLAGS = \
-no-undefined
libgpuverify_la_SOURCES = \
gpuv.c \
- gpuv-ref.c \
- gpuv-montg.c \
+ universal.c \
util.c
check_PROGRAMS = \
@@ -26,7 +25,6 @@ test_gpu_verify_SOURCES = \
# FIXME: on OSx, use '-framework OpenCL' instead -lOpenCL
test_gpu_verify_LDFLAGS = \
libgpuverify.la \
- -pthread \
-lgmp \
-lgcrypt \
-lOpenCL \
diff --git a/source/gpuv-montg.c b/source/gpuv-montg.c
@@ -12,14 +12,11 @@
#define ORDER -1 // I think we need to do this, because we want to write it in the 'wrong' way
#define END 0
-#define GPUV_BIT_LENGTH 2048
-
#define BITS 64
-//typedef uint32_t DIGIT_T;
-
// sizes are always the same 32 units for all, except exp_buf
-void gpuvt_read_files_montg(void *n_buf,
+void gpuvt_read_files(enum GPUV_VARIANT variant,
+ void *n_buf,
void *msg_buf,
void *exp_buf,
void *s_buf,
@@ -60,12 +57,21 @@ void gpuvt_read_files_montg(void *n_buf,
}
fclose (pk);
- gpu_register *n_buf_t = n_buf;
- gpu_register *msg_buf_t = msg_buf;
- gpu_register *s_buf_t = s_buf;
- unsigned long *exp_buf_t = exp_buf;
+ char *n_buf_t = n_buf;
+ char *msg_buf_t = msg_buf;
+ char *s_buf_t = s_buf;
+ char *exp_buf_t = exp_buf;
- int len = (GPUV_BIT_LENGTH / 8) / sizeof(gpu_register);
+ int len;
+
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ len = (GPUV_BIT_LENGTH_2048 / 8) / sizeof(DIGIT_T);
+ break;
+ default:
+ len = (GPUV_BIT_LENGTH_2048 / 8) / sizeof(gpu_register);
+ break;
+ }
char* message = strtok_r(ms_ptr, "\n", &ms_ptr_rest);
char* signature = strtok_r(0, "\n", &ms_ptr_rest);
@@ -96,8 +102,18 @@ void gpuvt_read_files_montg(void *n_buf,
exponent = strtok_r(0, "\n", &pk_ptr_rest);
offs = strtok_r(0, "\n", &pk_ptr_rest);
- mpz_export(&n_buf_t[len * j], NULL, ORDER, sizeof(gpu_register), END, 0, mod);
- mpz_export(&exp_buf_t[j], NULL, ORDER, sizeof(unsigned long), END, 0, e);
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ mpz_export(&n_buf_t[len * j * sizeof(DIGIT_T)], NULL, ORDER, sizeof(DIGIT_T), END, 0, mod);
+ mpz_export(&exp_buf_t[j * sizeof(uint32_t)], NULL, ORDER, sizeof(uint32_t), END, 0, e);
+ break;
+ default:
+ mpz_export(&n_buf_t[len * j* sizeof(gpu_register)], NULL, ORDER, sizeof(gpu_register), END, 0, mod);
+ mpz_export(&exp_buf_t[j * sizeof(unsigned long)], NULL, ORDER, sizeof(unsigned long), END, 0, e);
+ break;
+ }
+
+
j++;
@@ -110,8 +126,16 @@ void gpuvt_read_files_montg(void *n_buf,
message = strtok_r(0, "\n",&ms_ptr_rest);
signature = strtok_r(0, "\n",&ms_ptr_rest);
- mpz_export(&msg_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, msg);
- mpz_export(&s_buf_t[len * i], NULL, ORDER, sizeof(gpu_register), END, 0, s);
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ mpz_export(&msg_buf_t[len * i * sizeof(DIGIT_T)], NULL, ORDER, sizeof(DIGIT_T), END, 0, msg);
+ mpz_export(&s_buf_t[len * i * sizeof(DIGIT_T)], NULL, ORDER, sizeof(DIGIT_T), END, 0, s);
+ break;
+ default:
+ mpz_export(&msg_buf_t[len * i * sizeof(gpu_register)], NULL, ORDER, sizeof(gpu_register), END, 0, msg);
+ mpz_export(&s_buf_t[len * i * sizeof(gpu_register)], NULL, ORDER, sizeof(gpu_register), END, 0, s);
+ break;
+ }
i++;
}
@@ -145,7 +169,7 @@ void gpuvt_read_files_montg(void *n_buf,
static volatile bool finished = 0;
-void finally(void *cls, int valid, struct timespec time, unsigned long len, u_int32_t *res) {
+void finally(void *cls, int valid, struct timespec time, unsigned long len, uint32_t *res) {
if (valid) {
printf("VERIFICATION RESULT: OK\n\n");
@@ -153,24 +177,26 @@ void finally(void *cls, int valid, struct timespec time, unsigned long len, u_in
printf("VERIFICATION RESULT: NOT OK!\n");
}
- printf("GPU verification (MONTGOMERY) took %ld.%06ld s\n",time.tv_sec, time.tv_nsec);
+ printf("Verification took %ld.%06ld s\n",time.tv_sec, time.tv_nsec);
finished = 1;
}
-int gpuv_test_montg(void) {
+int gpuv_test(enum GPUV_VARIANT variant) {
+
+ finished = 0;
unsigned long pairs = gpuvt_estimate_pairs(); // returns an estimation of pairs
- unsigned long digit_sz = (GPUV_BIT_LENGTH / 8) * pairs;
+ unsigned long digit_sz = (GPUV_BIT_LENGTH_2048 / 8) * pairs;
unsigned long arr_sz = pairs * sizeof(unsigned long);
- gpu_register *n_buf = malloc(digit_sz);
- gpu_register *msg_buf = malloc(digit_sz);
- gpu_register *s_buf = malloc(digit_sz);
- gpu_register *exp_buf = malloc(pairs * sizeof(unsigned long));
+ char *n_buf = malloc(digit_sz);
+ char *msg_buf = malloc(digit_sz);
+ char *s_buf = malloc(digit_sz);
+ char *exp_buf = malloc(pairs * sizeof(unsigned long));
memset(n_buf, 0, digit_sz);
memset(msg_buf, 0, digit_sz);
@@ -178,12 +204,14 @@ int gpuv_test_montg(void) {
memset(exp_buf, 0, pairs * sizeof(unsigned long));
unsigned long *pks = malloc(arr_sz);
-
memset(pks, 0, arr_sz);
printf("READING KEYS...\n");
- gpuvt_read_files_montg(n_buf,
+ //enum GPUV_VARIANT variant = GPUV_CPU;
+
+ gpuvt_read_files(variant,
+ n_buf,
msg_buf,
exp_buf,
s_buf,
@@ -202,14 +230,18 @@ int gpuv_test_montg(void) {
pubks += 1;
-
-
-
-
-
struct gpuv_batch * batch = gpuv_prepare_batch();
- int len = (GPUV_BIT_LENGTH / 8) / sizeof(gpu_register);
+ int len;
+
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ len = (GPUV_BIT_LENGTH_2048 / 8) / sizeof(DIGIT_T);
+ break;
+ default:
+ len = (GPUV_BIT_LENGTH_2048 / 8) / sizeof(gpu_register);
+ break;
+ }
int x = 0;
@@ -218,21 +250,44 @@ int gpuv_test_montg(void) {
unsigned long range = (pks[i] + 1) - (i == 0 ? 0 : (pks[i - 1] + 1));
// inc & dec ref
+ size_t i_t;
+
+ struct gpuv_public_key * pub_key;
+
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ i_t = i * sizeof(DIGIT_T);
+ pub_key = gpuv_prepare_pubkey(((uint32_t *)exp_buf)[i], len, &n_buf[len * i_t]);
+ break;
+ default:
+ i_t = i * sizeof(gpu_register);
+ pub_key = gpuv_prepare_pubkey(((unsigned long *)exp_buf)[i], len, &n_buf[len * i_t]);
+ break;
+ }
+
- struct gpuv_public_key * pub_key = gpuv_prepare_pubkey(exp_buf[i], len, &n_buf[len * i]);
for (int j = 0; j < range; j++) {
- struct gpuv_signature_message * sig_msg = gpuv_prepare_sig_msg(pub_key);
+ size_t o;
- gpuv_add_message(sig_msg, len, &msg_buf[len * (j + x)]);
- gpuv_add_signature(sig_msg, len, &s_buf[len * (j + x)]);
+ switch (variant) {
+ case GPUV_GPU_REGULAR:
+ o = sizeof(DIGIT_T);
+ break;
+ default:
+ o = sizeof(gpu_register);
+ break;
+ }
- gpuv_add_to_batch(batch, sig_msg);
+ struct gpuv_signature_message * sig_msg = gpuv_prepare_sig_msg(pub_key);
+ gpuv_add_message(sig_msg, len, &msg_buf[len * (j + x) * o]);
+ gpuv_add_signature(sig_msg, len, &s_buf[len * (j + x) * o]);
+ gpuv_add_to_batch(batch, sig_msg);
- // destroy signature?
+ free(sig_msg); // sig_msg may be destroyed after adding it to the batch – public keys will be destroyed by freeing the state object
}
x+= range;
@@ -242,7 +297,7 @@ int gpuv_test_montg(void) {
struct gpuv_info *info;
// sets everything up...
- info = gpuv_init(MONTGOMERY);
+ info = gpuv_init(variant, GPUV_BIT_LENGTH_2048);
// every object represents a batch of data that should be processed on the GPU.
struct gpuv_state *state;
@@ -250,7 +305,7 @@ int gpuv_test_montg(void) {
state = gpuv_prepare(info, batch);
// A kernel always runs to completion – it can't be cancelled. But it is non-blocking, so the program can be terminated at any time.
- gpuv_start(state, &finally, NULL);
+ gpuv_start(state, &finally, NULL, batch);
fprintf(stderr, "KERNEL RUNNING...\n");
@@ -260,7 +315,7 @@ int gpuv_test_montg(void) {
- gpuv_free_batch(batch);
+ gpuv_free_batch(batch); // destroy public keys here
gpuv_free_state(state);
diff --git a/source/gpuv-montg.h b/source/gpuv-montg.h
@@ -13,7 +13,7 @@
#include "util.h"
-int gpuv_test_montg(void);
+int gpuv_test(enum GPUV_VARIANT variant);
diff --git a/source/gpuv-ref.c b/source/gpuv-ref.c
@@ -1,257 +0,0 @@
-//
-// gpuv-ref.c
-// lib-gpu-verify
-//
-// Created by Cedric Zwahlen on 06.12.2023.
-//
-
-#include "gpuv-ref.h"
-
-void ref_pairs_from_files(char *bases, unsigned long *b_off,
- char *exponents, unsigned long *e_off,
- char *moduli, unsigned long *m_off,
- char *signatures, unsigned long *s_off,
- unsigned long *pks,
- unsigned long *n) {
-
- FILE *pk;
- FILE *ms;
-
- pk = fopen("lib-gpu-generate/publickey.txt", "r");
- ms = fopen("lib-gpu-generate/msgsig.txt", "r");
-
- if (pk == NULL || ms == NULL) {
- printf("Auxiliary files not found.");
- abort();
- }
-
- fseek (ms, 0, SEEK_END);
- long ms_l = ftell(ms);
- fseek (ms, 0, SEEK_SET);
- char *ms_ptr = malloc(ms_l);
- char *ms_ptr_rest = malloc(ms_l);
- if (ms_ptr || ms_ptr_rest)
- {
- fread (ms_ptr, 1, ms_l, ms);
- memcpy(ms_ptr_rest, ms_ptr, ms_l);
- }
- fclose (ms);
-
- fseek (pk, 0, SEEK_END);
- long pk_l = ftell(pk);
- fseek (pk, 0, SEEK_SET);
- char *pk_ptr = malloc(pk_l);
- char *pk_ptr_rest = malloc(pk_l);
- if (pk_ptr && pk_ptr_rest)
- {
- fread (pk_ptr, 1, pk_l, pk);
- memcpy(pk_ptr_rest, pk_ptr, pk_l);
- }
- fclose (pk);
-
- char* message = strtok_r(ms_ptr, "\n", &ms_ptr_rest);
- char* signature = strtok_r(0, "\n", &ms_ptr_rest);
- char* modulus = strtok_r(pk_ptr, "\n", &pk_ptr_rest);
- char* exponent = strtok_r(0, "\n", &pk_ptr_rest);
- char* offs = strtok_r(0, "\n", &pk_ptr_rest);
-
- unsigned long b_offset = 0;
- unsigned long e_offset = 0;
- unsigned long m_offset = 0;
- unsigned long s_offset = 0;
-
- int i = 0;
- int j = 0;
-
- while (modulus != NULL && exponent != NULL && offs != NULL) {
-
- unsigned long n_buf_len = strlen(modulus);
- unsigned long e_buf_len = strlen(exponent);
-
- memcpy(&moduli[m_offset], modulus, n_buf_len);
- memcpy(&exponents[e_offset], exponent, e_buf_len);
-
- m_off[i] = m_offset;
- e_off[i] = e_offset;
-
- m_offset += n_buf_len + 1;
- e_offset += e_buf_len + 1;
-
- pks[i] = atoi(offs);
-
- modulus = strtok_r(0, "\n", &pk_ptr_rest);
- exponent = strtok_r(0, "\n", &pk_ptr_rest);
- offs = strtok_r(0, "\n", &pk_ptr_rest);
-
- i++;
- }
-
- while (message != NULL && signature != NULL) {
-
- unsigned long m_buf_len = strlen(message);
- unsigned long s_buf_len = strlen(signature);
-
- memcpy(&bases[b_offset], message, m_buf_len);
- memcpy(&signatures[s_offset], signature, s_buf_len);
-
- b_off[j] = b_offset;
- s_off[j] = s_offset;
-
- b_offset += m_buf_len + 1;
- s_offset += s_buf_len + 1;
-
- message = strtok_r(0, "\n",&ms_ptr_rest);
- signature = strtok_r(0, "\n",&ms_ptr_rest);
-
- j++;
-
- }
-
- *n = j;
-}
-
-gcry_sexp_t sexp_from_string(char* str, const char *format) {
-
- gcry_sexp_t sexp;
-
- gcry_mpi_t mpi = gcry_mpi_new((int)strlen(str) * 8);
- //size_t scanned = 0;
- gcry_mpi_scan(&mpi, GCRYMPI_FMT_HEX, str, 0, NULL);
-
- size_t errOff = 0;
- gcry_sexp_build(&sexp,&errOff,format,mpi);
-
- return sexp;
-}
-
-gcry_sexp_t sexp_from_string_key(char* str_1, char* str_2, const char *format) {
-
- gcry_sexp_t sexp;
-
- gcry_mpi_t mpi_1 = gcry_mpi_new((int)strlen(str_1) * 8);
- //size_t scanned = 0;
- gcry_mpi_scan(&mpi_1, GCRYMPI_FMT_HEX, str_1, 0, NULL);
-
- gcry_mpi_t mpi_2 = gcry_mpi_new((int)strlen(str_2) * 8);
- //size_t scanned = 0;
- gcry_mpi_scan(&mpi_2, GCRYMPI_FMT_HEX, str_2, 0, NULL);
-
- size_t errOff = 0;
- gcry_sexp_build(&sexp,&errOff,format,mpi_1,mpi_2);
-
- return sexp;
-}
-
-int gpuv_test_ref(void) {
-
- unsigned long pairs = gpuvt_estimate_pairs();
-
- unsigned long str_sz = (2048) * pairs;
-
-
- char *b = malloc(str_sz);
- char *e = malloc(str_sz);
- char *m = malloc(str_sz);
- char *s = malloc(str_sz);
-
- unsigned long *b_off = malloc(str_sz);
- unsigned long *e_off = malloc(str_sz);
- unsigned long *m_off = malloc(str_sz);
- unsigned long *s_off = malloc(str_sz);
-
- memset(b, 0, str_sz);
- memset(e, 0, str_sz);
- memset(m, 0, str_sz);
- memset(s, 0, str_sz);
-
- memset(b_off, 0, str_sz);
- memset(e_off, 0, str_sz);
- memset(m_off, 0, str_sz);
- memset(s_off, 0, str_sz);
-
- unsigned long *pks = malloc(str_sz);
- memset(pks, 0, str_sz);
-
- ref_pairs_from_files(b, b_off, e, e_off, m, m_off, s, s_off, pks,
- &pairs);
-
- unsigned long pk = 0;
-
- while (1) {
- if (pks[pk] + 1 == pairs)
- break;
- pk++;
- }
-
-
-
- gcry_sexp_t *m_sexps = malloc(pairs * sizeof(gcry_sexp_t));
- gcry_sexp_t *s_sexps = malloc(pairs * sizeof(gcry_sexp_t));
- gcry_sexp_t *key_sexps = malloc((pk + 1) * sizeof(gcry_sexp_t));
-
- for (int i = 0; i < pairs; i++) {
-
- m_sexps[i] = sexp_from_string(&b[b_off[i]], "(data (flags raw) (value %m))"); // message format (for comparison)
-
- s_sexps[i] = sexp_from_string(&s[s_off[i]], "(sig-val (rsa (s %m)))"); // signature format
- }
-
-
- for (int i = 0; i <= pk; i++) {
-
- key_sexps[i] = sexp_from_string_key(&m[m_off[i]], &e[e_off[i]], "(public-key (rsa (n %m) (e %m)))" ); // pub key data
-
- }
-
- unsigned long result = 0;
-
- struct timespec t1, t2;
-
- printf("VERIFYING %lu SIGNATURES...\n", pairs);
-
- clock_gettime(CLOCK_REALTIME, &t1);
-
- pk = 0; // reuse pk
-
- for (int i = 0; i < pairs; i++) {
-
- while (1) {
- if (pks[pk] >= i)
- break;
- pk++;
- }
-
- if ( gcry_pk_verify(s_sexps[i], m_sexps[i], key_sexps[pk]) == 0 )
- result += 1;
-
- }
-
- clock_gettime(CLOCK_REALTIME, &t2);
-
- printf("CPU (Reference) verification took %ld.%06ld s\n", ( t2.tv_nsec < t1.tv_nsec ? t2.tv_sec - (t1.tv_sec + 1) : t2.tv_sec - t1.tv_sec ), ( t2.tv_nsec < t1.tv_nsec ? ((999999999 - t1.tv_nsec) + t2.tv_nsec) : (t2.tv_nsec - t1.tv_nsec) ) / 1000);
-
- if (result == pairs) {
- printf("VERIFICATION RESULT: %lu - OK\n\n",result);
- } else {
- printf("VERIFICATION RESULT: %lu - NOT OK\n\n",result);
- }
-
-
- free(b);
- free(e);
- free(m);
- free(s);
-
- free(b_off);
- free(e_off);
- free(m_off);
- free(s_off);
-
- free(pks);
-
- free(m_sexps);
- free(s_sexps);
- free(key_sexps);
-
- return result == pairs ? 1 : 0;
-}
diff --git a/source/gpuv-ref.h b/source/gpuv-ref.h
@@ -1,16 +0,0 @@
-//
-// gpuv-ref.h
-// lib-gpu-verify
-//
-// Created by Cedric Zwahlen on 06.12.2023.
-//
-
-
-#ifndef gpuv_ref_h
-#define gpuv_ref_h
-
-#include "util.h"
-
-int gpuv_test_ref(void);
-
-#endif /* gpuv-ref_h */
diff --git a/source/gpuv.c b/source/gpuv.c
@@ -1,439 +0,0 @@
-/*
- * lib-gpu-verify
- *
- * This software contains code derived from or inspired by the BigDigit library,
- * <http://www.di-mgt.com.au/bigdigits.html>
- * which is distributed under the Mozilla Public License, version 2.0.
- *
- * The original code and modifications made to it are subject to the terms and
- * conditions of the Mozilla Public License, version 2.0. A copy of the
- * MPL license can be obtained at
- * https://www.mozilla.org/en-US/MPL/2.0/.
- *
- * Changes and additions to the original code are as follows:
- * - Copied some functions of the BigDigit library into this file, to convert strings read from files to BigDigit type numbers.
- *
- * Contributors:
- * - Cedric Zwahlen cedric.zwahlen@bfh.ch
- *
- * Please note that this software is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the Mozilla Public License, version 2.0, for the specific language
- * governing permissions and limitations under the License.
- */
-
-#include "gpuv.h"
-#include "math.h"
-
-#define BITS_PER_DIGIT 32
-
-#define MAX_ALLOC_SIZE 256
-
-#define BYTES_PER_DIGIT (BITS_PER_DIGIT / 8)
-
-
-
-#define ORDER -1 // I think we need to do this, because we want to write it in the 'wrong' way
-#define END 0
-
-#define GPUV_BIT_LENGTH 2048
-
-#define WORD_COUNT ((GPUV_BIT_LENGTH / 8) / sizeof(DIGIT_T))
-
-
-typedef uint32_t DIGIT_T;
-
-// MARK: OPENCL CODE
-void gpuvt_read_files_std(void *n_buf,
- void *msg_buf,
- void *exp_buf,
- void *s_buf,
- unsigned long *pks, unsigned long *n) {
-
- FILE * pk;
- FILE * ms;
-
- pk = fopen("lib-gpu-generate/publickey.txt", "r");
- ms = fopen("lib-gpu-generate/msgsig.txt", "r");
-
- if (pk == NULL || ms == NULL) {
- printf("Auxiliary files not found.");
- abort();
- }
-
- fseek (ms, 0, SEEK_END);
- long ms_l = ftell(ms);
- fseek (ms, 0, SEEK_SET);
- char *ms_ptr = malloc(ms_l);
- char *ms_ptr_rest = ms_ptr;
- if (ms_ptr || ms_ptr_rest)
- {
- fread (ms_ptr, 1, ms_l, ms);
- memcpy(ms_ptr_rest, ms_ptr, ms_l);
- }
- fclose (ms);
-
- fseek (pk, 0, SEEK_END);
- long pk_l = ftell(pk);
- fseek (pk, 0, SEEK_SET);
- char *pk_ptr = malloc(pk_l);
- char *pk_ptr_rest = pk_ptr;
- if (pk_ptr && pk_ptr_rest)
- {
- fread (pk_ptr, 1, pk_l, pk);
- memcpy(pk_ptr_rest, pk_ptr, pk_l);
- }
- fclose (pk);
-
- DIGIT_T *n_buf_t = n_buf;
- DIGIT_T *msg_buf_t = msg_buf;
- DIGIT_T *s_buf_t = s_buf;
- DIGIT_T *exp_buf_t = exp_buf;
-
- int len = WORD_COUNT;
-
- char* message = strtok_r(ms_ptr, "\n", &ms_ptr_rest);
- char* signature = strtok_r(0, "\n", &ms_ptr_rest);
- char* modulus = strtok_r(pk_ptr, "\n", &pk_ptr_rest);
- char* exponent = strtok_r(0, "\n", &pk_ptr_rest);
- char* offs = strtok_r(0, "\n", &pk_ptr_rest);
-
- int i = 0;
- int j = 0;
-
- mpz_t e,mod,msg,s;
-
- mpz_init(e);
- mpz_init(mod);
- mpz_init(msg);
- mpz_init(s);
-
- while (message != NULL && signature != NULL) {
-
- if (i == 0 || pks[j - 1] < i) {
-
- mpz_set_str(mod,modulus,16);
- mpz_set_str(e,exponent,16);
-
- pks[j] = atoi(offs);
-
- modulus = strtok_r(0, "\n", &pk_ptr_rest);
- exponent = strtok_r(0, "\n", &pk_ptr_rest);
- offs = strtok_r(0, "\n", &pk_ptr_rest);
-
- mpz_export(&n_buf_t[len * j], NULL, ORDER, sizeof(DIGIT_T), END, 0, mod);
- mpz_export(&exp_buf_t[j], NULL, ORDER, sizeof(DIGIT_T), END, 0, e);
-
-
- j++;
-
- }
-
- mpz_set_str(msg,message,16);
- mpz_set_str(s,signature,16);
-
- message = strtok_r(0, "\n",&ms_ptr_rest);
- signature = strtok_r(0, "\n",&ms_ptr_rest);
-
- mpz_export(&msg_buf_t[len * i], NULL, ORDER, sizeof(DIGIT_T), END, 0, msg);
- mpz_export(&s_buf_t[len * i], NULL, ORDER, sizeof(DIGIT_T), END, 0, s);
-
- i++;
- }
-
- mpz_clear(e);
- mpz_clear(mod);
- mpz_clear(msg);
- mpz_clear(s);
-
- free(ms_ptr);
- free(pk_ptr);
-
-
- *n = i;
-
-}
-
-//
-//
-//struct gpuv_public_key * gpuv_get_pubkey_std(DIGIT_T e, unsigned long len_n, DIGIT_T *n) {
-//
-// struct gpuv_public_key * pk = malloc(sizeof(struct gpuv_public_key));
-// memset(pk, 0, sizeof(struct gpuv_public_key));
-//
-// pk->e = e;
-// pk->n = (char *)n;
-// pk->len_n = len_n;
-//
-// return pk;
-//}
-//
-//struct gpuv_signature_message * gpuv_get_sigmsg_std(struct gpuv_public_key *pk,
-// unsigned long len_single_m, DIGIT_T *m,
-// unsigned long len_single_s, DIGIT_T *s,
-// unsigned long count) {
-//
-// struct gpuv_signature_message * sigmsg = malloc(sizeof(struct gpuv_signature_message));
-// memset(sigmsg, 0, sizeof(struct gpuv_signature_message));
-//
-// sigmsg->pubkey = pk;
-// sigmsg->m = (char *)m;
-// sigmsg->len_m = len_single_m * count;
-// sigmsg->s = (char *)s;
-// sigmsg->len_s = len_single_s * count;
-//
-// return sigmsg;
-//}
-
-
-///
-/// prepares a state object that contains buffers etc for the gpu
-//struct gpuv_state * gpuv_prepare_std(struct gpuv_info *info, struct gpuv_batch * batch) {
-//
-// struct timespec p1, p2;
-//
-// clock_gettime(CLOCK_REALTIME, &p1);
-//
-// struct gpuv_state *state = malloc(sizeof(struct gpuv_state));
-// memset(state, 0, sizeof(struct gpuv_state));
-//
-// state->info = info;
-// state->queue = create_command_queue (info->device_id, info->context);
-//
-// state->event_kernel = malloc(sizeof(cl_event));
-// state->event_results = malloc(sizeof(cl_event));
-//
-// unsigned long sig_pairs = 0;
-// for (int i = 0; i < batch->pairs_len; i++) { sig_pairs += batch->pairs[i].count; }
-//
-// state->sig_count = sig_pairs;
-// state->pubkey_count = batch->pairs_len;
-//
-// unsigned long res_len = ceil((double)sig_pairs / (double)(sizeof(u_int32_t) * 8)); // how many uint32
-// unsigned long res_len_bytes = res_len * sizeof(u_int32_t); // how many bytes needed
-//
-// u_int32_t *results_buf = malloc(res_len_bytes);
-// memset(results_buf, 0, res_len_bytes);
-//
-// // set up the results, which can be accessed by the user once the kernel has run
-// state->results = results_buf;
-// state->results_len = res_len;
-//
-// unsigned long len = (GPUV_BIT_LENGTH / 8) * sig_pairs;
-//
-// switch (info->variant) {
-// case MONTGOMERY:
-//
-// state->res_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, res_len_bytes ,NULL, NULL);
-// state->x_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
-// state->m_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
-// state->n_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
-// state->ni_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
-// state->exp_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sig_pairs * sizeof(DIGIT_T), NULL, NULL);
-// state->msg_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
-// // plus 1, because we need an additional element in the array, containing it's length
-// state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(unsigned long) * (batch->pairs_len + 1), NULL, NULL);
-//
-// if (!state->res_mem || !state->msg_mem || !state->pks_indices ||
-// !state->x_mem || !state->m_mem || !state->n_mem || !state->ni_mem || !state->exp_mem)
-// {
-// printf("Error: Failed to allocate device memory!\n");
-// exit(1);
-// }
-//
-// int l = sizeof(DIGIT_T);
-//
-// int err = 0;
-//
-// unsigned long off_x = 0;
-// unsigned long off_M = 0;
-// unsigned long off_n = 0;
-// unsigned long off_ni = 0;
-// unsigned long off_e = 0;
-// unsigned long off_m = 0;
-//
-// unsigned long off_pks = 0;
-//
-// err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, 0, sizeof(unsigned long), &state->pubkey_count, 0, NULL, NULL);
-//
-// for (int i = 0; i < batch->pairs_len; i++) {
-//
-// struct gpuv_signature_message * b = &batch->pairs[i];
-//
-// off_pks += b->count;
-//
-// unsigned long opks = off_pks - 1;
-//
-// // Write our data set into the input array in device memory
-// //
-//
-// err |= clEnqueueWriteBuffer(state->queue, state->x_mem, CL_TRUE, off_x, b->len_x * l, b->x, 0, NULL, NULL);
-// err |= clEnqueueWriteBuffer(state->queue, state->m_mem, CL_TRUE, off_M, b->len_M * l, b->M, 0, NULL, NULL);
-// err |= clEnqueueWriteBuffer(state->queue, state->n_mem, CL_TRUE, off_n, b->pubkey->len_n * l, b->pubkey->n, 0, NULL, NULL);
-// err |= clEnqueueWriteBuffer(state->queue, state->ni_mem, CL_TRUE, off_ni, b->pubkey->len_ni * l, b->pubkey->ni, 0, NULL, NULL);
-// err |= clEnqueueWriteBuffer(state->queue, state->exp_mem, CL_TRUE, off_e, l, &b->pubkey->e, 0, NULL, NULL);
-// err |= clEnqueueWriteBuffer(state->queue, state->msg_mem, CL_TRUE, off_m, b->len_m * l, b->m, 0, NULL, NULL);
-// // the kernel wants to know how many elements are in the array, so we write it at the beginning
-// err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, (i + 1) * sizeof(unsigned long), sizeof(unsigned long), &opks, 0, NULL, NULL);
-//
-//
-//
-// if (err != CL_SUCCESS)
-// {
-// printf("Error: Failed to write to source array!\n");
-// exit(1);
-// }
-//
-// off_x += b->len_x * l;
-// off_M += b->len_M * l;
-// off_n += b->pubkey->len_n * l;
-// off_ni += b->pubkey->len_ni * l;
-// off_e += l;
-// off_m += b->len_m * l;
-//
-// }
-//
-// err |= clEnqueueWriteBuffer(state->queue, state->res_mem, CL_TRUE, 0, state->results_len * sizeof(u_int32_t), state->results, 0, NULL, NULL);
-//
-// break;
-//
-// default:
-// break;
-// }
-//
-//
-// clock_gettime(CLOCK_REALTIME, &p2);
-//
-// state->p.tv_sec = ( p2.tv_nsec < p1.tv_nsec ? p2.tv_sec - (p1.tv_sec + 1) : p2.tv_sec - p1.tv_sec );
-// state->p.tv_nsec = ( p2.tv_nsec < p1.tv_nsec ? ((999999999 - p1.tv_nsec) + p2.tv_nsec) : (p2.tv_nsec - p1.tv_nsec) ) / 1000;
-//
-// state->ready = 1;
-//
-// return state;
-//
-//}
-
-void finally_std(void *cls, int valid, struct timespec time, unsigned long len, u_int32_t *res) {
-
- if (valid) {
- printf("VERIFICATION RESULT: OK\n\n");
- } else {
- printf("VERIFICATION RESULT: NOT OK!\n");
- }
-
- printf("GPU verification took %ld.%06ld s\n",time.tv_sec, time.tv_nsec);
-
-}
-
-int gpuv_test(void) {
-
- unsigned long pairs = gpuvt_estimate_pairs(); // returns an estimation of pairs
-
- unsigned long digit_sz = (GPUV_BIT_LENGTH / 8) * pairs;
-
- unsigned long arr_sz = pairs * sizeof(unsigned long);
-
- DIGIT_T *n_buf = malloc(digit_sz);
- DIGIT_T *msg_buf = malloc(digit_sz);
- DIGIT_T *s_buf = malloc(digit_sz);
- DIGIT_T *exp_buf = malloc(pairs * sizeof(DIGIT_T));
-
- memset(n_buf, 0, digit_sz);
- memset(msg_buf, 0, digit_sz);
- memset(s_buf, 0, digit_sz);
- memset(exp_buf, 0, pairs * sizeof(DIGIT_T));
-
- unsigned long *pks = malloc(arr_sz);
-
- memset(pks, 0, arr_sz);
-
- printf("READING KEYS...\n");
-
- gpuvt_read_files_std(n_buf,
- msg_buf,
- exp_buf,
- s_buf,
- pks, &pairs);
-
- printf("VERIFYING %lu SIGNATURES...\n", pairs);
-
- unsigned long pubks = 0;
-
- while (1) {
- if (pks[pubks] + 1 >= pairs)
- break;
- pubks++;
- }
-
- pubks += 1;
-
-
- struct gpuv_batch * batch;
- batch = gpuv_prepare_batch();
-
- int len = (GPUV_BIT_LENGTH / 8) / sizeof(DIGIT_T);
-
- int x = 0;
-
- for (int i = 0; i < pubks; i++) {
-
- unsigned long range = (pks[i] + 1) - (i == 0 ? 0 : (pks[i - 1] + 1));
-
- // inc & dec ref
-
- struct gpuv_public_key * pub_key = gpuv_prepare_pubkey(exp_buf[i], len, &n_buf[len * i]);
-
- for (int j = 0; j < range; j++) {
-
- struct gpuv_signature_message * sig_msg = gpuv_prepare_sig_msg(pub_key);
-
- gpuv_add_message(sig_msg, len, &msg_buf[len * (j + x)]);
- gpuv_add_signature(sig_msg, len, &s_buf[len * (j + x)]);
-
- gpuv_add_to_batch(batch, sig_msg);
-
-
-
- // destroy signature?
- }
-
- x+= range;
-
- }
-
-
-
-
- struct gpuv_info *info;
-
- // sets everything up...
- info = gpuv_init(REGULAR);
-
- // every object represents a batch of data that should be processed on the GPU.
- struct gpuv_state *state;
-
- state = gpuv_prepare(info, batch);
-
- gpuv_start(state, &finally_std, NULL);
-
-
- // wait for the job to finish
- sleep(3);
-
- gpuv_finish(info);
-
- gpuv_free_state(state);
- gpuv_free_batch(batch);
-
-// free(info);
-
- free(n_buf);
- free(msg_buf);
- free(s_buf);
- free(exp_buf);
-
- free(pks);
-
- return 0;
-}
-
diff --git a/source/gpuv.h b/source/gpuv.h
@@ -1,19 +0,0 @@
-//
-// gpuv.h
-// lib-gpu-verify
-//
-// Created by Cedric Zwahlen on 28.09.2023.
-//
-
-#ifndef gpuv_h
-#define gpuv_h
-
-#include "util.h"
-#include <stdint.h>
-#include <gmp.h>
-
-int gpuv_test(void);
-
-
-
-#endif /* gpuv_h */
diff --git a/source/lib-gpu-verify.c b/source/lib-gpu-verify.c
@@ -5,10 +5,6 @@
// Created by Cedric Zwahlen on 28.09.2023.
//
-
-#include "gpuv.h"
-
-#include "gpuv-ref.h"
#include "gpuv-montg.h"
@@ -19,11 +15,11 @@ int main(int argc, char** argv)
gpuv_prepare_gcry();
- gpuv_test_montg();
-
-// gpuv_test();
+ gpuv_test(GPUV_GPU_MONTGOMERY);
+
+// gpuv_test(GPUV_CPU);
- //gpuv_test_ref();
+// gpuv_test(GPUV_GPU_REGULAR);
diff --git a/source/universal.c b/source/universal.c
@@ -9,7 +9,7 @@
#include "util.h"
#include <math.h>
-struct gpuv_info * gpuv_init(enum Variant variant) {
+struct gpuv_info * gpuv_init(enum GPUV_VARIANT variant, enum GPUV_BIT_LENGTH bit_length) {
struct gpuv_info *info = malloc(sizeof(struct gpuv_info));
memset(info, 0, sizeof(struct gpuv_info));
@@ -21,15 +21,19 @@ struct gpuv_info * gpuv_init(enum Variant variant) {
info->context = create_compute_context (info->device_id);
switch (variant) {
- case MONTGOMERY:
+ case GPUV_GPU_MONTGOMERY:
info->program = compile_program (info->device_id, info->context, "gpuv-montg.cl");
info->kernel = create_kernel (info->program, "mont");
break;
- case REGULAR:
+ case GPUV_GPU_REGULAR:
info->program = compile_program (info->device_id, info->context, "gpuv.cl");
info->kernel = create_kernel (info->program, "several");
break;
+ case GPUV_CPU:
+
+ break;
+
default:
return NULL;
break;
@@ -46,11 +50,11 @@ struct gpuv_batch * gpuv_prepare_batch(void) {
struct gpuv_signature_message * ms = malloc(sizeof(struct gpuv_signature_message) * 64);
memset(ms, 0, sizeof(struct gpuv_signature_message) * 64);
- u_int32_t * pk_indices = malloc(sizeof(u_int32_t) * 64);
- memset(pk_indices, 0, sizeof(u_int32_t) * 64);
+ uint32_t * pk_indices = malloc(sizeof(uint32_t) * 64);
+ memset(pk_indices, 0, sizeof(uint32_t) * 64);
- u_int32_t * pk_list = malloc(sizeof(u_int32_t) * 64);
- memset(pk_list, 0, sizeof(u_int32_t) * 64);
+ uint32_t * pk_list = malloc(sizeof(uint32_t) * 64);
+ memset(pk_list, 0, sizeof(uint32_t) * 64);
batch->pairs = ms;
batch->max_count = 64;
@@ -106,10 +110,10 @@ int gpuv_add_to_batch(struct gpuv_batch * batch, struct gpuv_signature_message *
if (batch->current >= batch->max_count) {
unsigned long pl = batch->max_count * 2;
struct gpuv_signature_message * p = realloc(batch->pairs, pl * sizeof(struct gpuv_signature_message));
- u_int32_t * p_list = realloc(batch->pk_list, pl * sizeof(u_int32_t));
+ uint32_t * p_list = realloc(batch->pk_list, pl * sizeof(uint32_t));
if (p == NULL || p_list == NULL) { return 1; }
memset(&p[batch->max_count], 0, batch->max_count * sizeof(struct gpuv_signature_message));
- memset(&p_list[batch->max_count], 0, batch->max_count * sizeof(u_int32_t));
+ memset(&p_list[batch->max_count], 0, batch->max_count * sizeof(uint32_t));
batch->pairs = p;
batch->pk_list = p_list;
batch->max_count = pl;
@@ -121,8 +125,8 @@ int gpuv_add_to_batch(struct gpuv_batch * batch, struct gpuv_signature_message *
if (batch->pk_current >= batch->pk_max_count) {
unsigned long pl = batch->pk_max_count * 2;
- u_int32_t * p = realloc(batch->pk_indices, pl * sizeof(u_int32_t));
- memset(&p[batch->pk_max_count], 0, batch->pk_max_count * sizeof(u_int32_t));
+ uint32_t * p = realloc(batch->pk_indices, pl * sizeof(uint32_t));
+ memset(&p[batch->pk_max_count], 0, batch->pk_max_count * sizeof(uint32_t));
if (p == NULL) { return 1; }
batch->pk_indices = p;
batch->pk_max_count = pl;
@@ -146,9 +150,9 @@ int gpuv_add_to_batch(struct gpuv_batch * batch, struct gpuv_signature_message *
} else {
- batch->pk_indices[batch->pk_current] = (u_int32_t)batch->current;
+ batch->pk_indices[batch->pk_current] = (uint32_t)batch->current;
- batch->pk_list[batch->current] = (u_int32_t)batch->pk_current;
+ batch->pk_list[batch->current] = (uint32_t)batch->pk_current;
batch->pk_current++;
}
@@ -161,12 +165,16 @@ int gpuv_add_to_batch(struct gpuv_batch * batch, struct gpuv_signature_message *
void gpuv_free_batch(struct gpuv_batch * batch) {
- for(int i = 0; i < batch->max_count; i++) {
+ for(int i = 0; i < batch->pk_current; i++) {
-// free(batch->pairs[i].pubkey->ni);
-// free(batch->pairs[i].pubkey->r_1);
+ free(batch->pairs[ batch->pk_indices[i] ].pubkey->ni);
+ free(batch->pairs[ batch->pk_indices[i] ].pubkey->r_1);
+
+ free(batch->pairs[ batch->pk_indices[i] ].pubkey);
- //free(batch->pairs[i].pubkey);
+ }
+
+ for(int i = 0; i < batch->current; i++) {
free(batch->pairs[i].M);
free(batch->pairs[i].x);
@@ -174,6 +182,9 @@ void gpuv_free_batch(struct gpuv_batch * batch) {
}
free(batch->pairs);
+ free(batch->pk_indices);
+ free(batch->pk_list);
+
free(batch);
}
@@ -201,26 +212,26 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
state->info = info;
state->queue = create_command_queue (info->device_id, info->context);
- state->event_kernel = malloc(sizeof(cl_event));
- state->event_results = malloc(sizeof(cl_event));
+// state->event_kernel = malloc(sizeof(cl_event));
+// state->event_results = malloc(sizeof(cl_event));
unsigned long sig_pairs = batch->current;
state->sig_count = sig_pairs;
state->pubkey_count = batch->pk_current;
- unsigned long res_len = ceil((double)sig_pairs / (double)(sizeof(u_int32_t) * 8)); // how many uint32
- unsigned long res_len_bytes = res_len * sizeof(u_int32_t); // how many bytes needed
+ unsigned long res_len = ceil((double)sig_pairs / (double)(sizeof(uint32_t) * 8)); // how many uint32
+ unsigned long res_len_bytes = res_len * sizeof(uint32_t); // how many bytes needed
- u_int32_t *results_buf = malloc(res_len_bytes);
+ uint32_t *results_buf = malloc(res_len_bytes);
memset(results_buf, 0, res_len_bytes);
// set up the results, which can be accessed by the user once the kernel has run
state->results = results_buf;
state->results_len = res_len;
- unsigned long len = (GPUV_BIT_LENGTH / 8) * state->sig_count;
- unsigned long pk_len = (GPUV_BIT_LENGTH / 8) * state->pubkey_count;
+ unsigned long len = (GPUV_BIT_LENGTH_2048 / 8) * state->sig_count;
+ unsigned long pk_len = (GPUV_BIT_LENGTH_2048 / 8) * state->pubkey_count;
int err = 0;
@@ -233,7 +244,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
unsigned long off_s = 0;
switch (info->variant) {
- case MONTGOMERY:
+ case GPUV_GPU_MONTGOMERY:
state->res_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, res_len_bytes ,NULL, NULL);
state->x_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
@@ -242,7 +253,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
state->ni_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, pk_len, NULL, NULL);
state->exp_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, batch->pk_current * sizeof(unsigned long), NULL, NULL);
state->msg_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL);
- state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(u_int32_t) * batch->current, NULL, NULL);
+ state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(uint32_t) * batch->current, NULL, NULL);
if (!state->res_mem || !state->msg_mem || !state->pks_indices ||
!state->x_mem || !state->m_mem || !state->n_mem || !state->ni_mem || !state->exp_mem)
@@ -287,7 +298,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
}
- err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, 0, sizeof(u_int32_t) * batch->current, batch->pk_list, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, 0, sizeof(uint32_t) * batch->current, batch->pk_list, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(state->queue, state->res_mem, CL_TRUE, 0, res_len_bytes, state->results, 0, NULL, NULL);
if (err != CL_SUCCESS)
@@ -299,7 +310,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
break;
- case REGULAR:
+ case GPUV_GPU_REGULAR:
state->res_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, res_len_bytes ,NULL, NULL);
@@ -308,7 +319,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
state->sig_mem = clCreateBuffer(info->context, CL_MEM_READ_WRITE, len, NULL, NULL);
state->msg_mem = clCreateBuffer(info->context, CL_MEM_READ_ONLY, len, NULL, NULL); // the base, to compare whether we get the same signature
- state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(u_int32_t) * batch->current, NULL, NULL);
+ state->pks_indices = clCreateBuffer(info->context, CL_MEM_READ_ONLY, sizeof(uint32_t) * batch->current, NULL, NULL);
if (!state->sig_mem || !state->exp_mem || !state->n_mem || !state->msg_mem || !state->pks_indices)
@@ -344,7 +355,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
}
- err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, 0, sizeof(u_int32_t) * batch->current, batch->pk_list, 0, NULL, NULL);
+ err |= clEnqueueWriteBuffer(state->queue, state->pks_indices, CL_TRUE, 0, sizeof(uint32_t) * batch->current, batch->pk_list, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(state->queue, state->res_mem, CL_TRUE, 0, res_len_bytes, state->results, 0, NULL, NULL);
if (err != CL_SUCCESS)
@@ -381,7 +392,7 @@ struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * bat
Returns immediately
*/
-int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespec, unsigned long, u_int32_t *), void * arg) {
+int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespec, unsigned long, uint32_t *), void * arg, struct gpuv_batch *batch) {
if (state->stale) {
printf("State submitted twice.\n");
@@ -407,7 +418,7 @@ int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespe
int err = 0;
switch (state->info->variant) {
- case MONTGOMERY:
+ case GPUV_GPU_MONTGOMERY:
err |= clSetKernelArg(state->info->kernel, 0, sizeof(cl_mem), &state->x_mem);
err |= clSetKernelArg(state->info->kernel, 1, sizeof(cl_mem), &state->m_mem);
@@ -420,9 +431,9 @@ int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespe
break;
- case REGULAR:
+ case GPUV_GPU_REGULAR:
- err = clSetKernelArg(state->info->kernel, 0, sizeof(cl_mem), &state->sig_mem);
+ err |= clSetKernelArg(state->info->kernel, 0, sizeof(cl_mem), &state->sig_mem);
err |= clSetKernelArg(state->info->kernel, 1, sizeof(cl_mem), &state->exp_mem);
err |= clSetKernelArg(state->info->kernel, 2, sizeof(cl_mem), &state->n_mem);
err |= clSetKernelArg(state->info->kernel, 3, sizeof(cl_mem), &state->msg_mem);
@@ -432,6 +443,15 @@ int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespe
break;
+ case GPUV_CPU:
+
+
+ cpu_verify(batch, state);
+
+ return 0;
+
+ break;
+
default:
break;
}
@@ -447,16 +467,17 @@ int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespe
clock_gettime(CLOCK_REALTIME, &state->t);
- state->event_kernel = clCreateUserEvent(state->info->context, NULL);
+ cl_event e = clCreateUserEvent(state->info->context, NULL);
- err = clEnqueueNDRangeKernel(state->queue, state->info->kernel, 1, NULL, &state->sig_count, NULL, 0, NULL, &state->event_kernel);
+ err = clEnqueueNDRangeKernel(state->queue, state->info->kernel, 1, NULL, &state->sig_count, NULL, 0, NULL, &e);
if (err)
{
printf("Error: Failed to execute kernel!\n");
return 1;
}
- clSetEventCallback(state->event_kernel, CL_COMPLETE, callback_kernel, state);
+ clSetEventCallback(e, CL_COMPLETE, callback_kernel, state);
+ clRetainEvent(e);
return 0;
}
diff --git a/source/universal.h b/source/universal.h
@@ -10,8 +10,9 @@
#include <stdio.h>
#include <time.h>
+#include <stdint.h>
-#define GPUV_BIT_LENGTH 2048 // put in enum
+//#define GPUV_BIT_LENGTH 2048 // put in enum
@@ -29,15 +30,21 @@ struct gpuv_info;
struct gpuv_state;
-enum Variant {
+enum GPUV_VARIANT {
- REGULAR = 0,
- MONTGOMERY = 1,
- CPU = 2
+ GPUV_GPU_REGULAR = 0,
+ GPUV_GPU_MONTGOMERY = 1,
+ GPUV_CPU = 2
};
-struct gpuv_info * gpuv_init(enum Variant variant);
+enum GPUV_BIT_LENGTH {
+
+ GPUV_BIT_LENGTH_2048 = 2048,
+
+};
+
+struct gpuv_info * gpuv_init(enum GPUV_VARIANT variant, enum GPUV_BIT_LENGTH bit_length);
struct gpuv_batch * gpuv_prepare_batch(void);
int gpuv_add_to_batch(struct gpuv_batch * batch, struct gpuv_signature_message * sigmem);
@@ -51,7 +58,7 @@ void gpuv_free_batch(struct gpuv_batch * batch);
void gpuv_free_state(struct gpuv_state * state);
struct gpuv_state * gpuv_prepare(struct gpuv_info *info, struct gpuv_batch * batch);
-int gpuv_start(struct gpuv_state *state, void (*cls)(void*, int, struct timespec, unsigned long, u_int32_t *), void * arg);
+int gpuv_start(struct gpuv_state *state, void (*cls)(void *, int, struct timespec, unsigned long, uint32_t *), void * arg, struct gpuv_batch *batch);
void gpuv_finish(struct gpuv_info * info);
#endif /* universal_h */
diff --git a/source/util.c b/source/util.c
@@ -24,7 +24,7 @@ unsigned long gpuvt_estimate_pairs(void) {
fstat(msfile, &ss);
unsigned long len_f = ss.st_size;
- unsigned long len_sig = (2048 / 8) * 2 + 1; // this is the size of a 2048 bit signature in the file
+ unsigned long len_sig = (GPUV_BIT_LENGTH_2048 / 8) * 2 + 1; // this is the size of a 2048 bit signature in the file
unsigned long n_min = len_f / (len_sig + 3); // if each message was only one character, then this would be the maximum amount of signatures that could be in the file – use this estimate to allocate storage for the signatures
@@ -359,7 +359,7 @@ create_kernel (cl_program program,
*/
void CL_CALLBACK callback_result(cl_event event, cl_int event_command_status, void *user_data) {
- clReleaseEvent(event);
+ //clReleaseEvent(event);
struct gpuv_state *state = (struct gpuv_state *)user_data;
@@ -388,14 +388,14 @@ void CL_CALLBACK callback_result(cl_event event, cl_int event_command_status, vo
int ret = 1;
- unsigned long partial = state->sig_count / (sizeof(u_int32_t) * 8);
+ unsigned long partial = state->sig_count / (sizeof(uint32_t) * 8);
for(int i = 0; i < state->results_len; i++) {
uint32_t mask = 0;
if (i >= partial) {
- int remaining = state->sig_count % (sizeof(u_int32_t) * 8);
+ int remaining = state->sig_count % (sizeof(uint32_t) * 8);
for (int x = 0; x < remaining; x++) { mask |= 1 << x; }
@@ -421,28 +421,29 @@ void CL_CALLBACK callback_result(cl_event event, cl_int event_command_status, vo
*/
void CL_CALLBACK callback_kernel(cl_event event, cl_int event_command_status, void *user_data) {
- clReleaseEvent(event); // only call here if we wait for completion
+ //clReleaseEvent(event); // only call here if we wait for completion
struct gpuv_state *state = (struct gpuv_state *)user_data;
//MARK: no expensive operations here
unsigned long res_len = state->results_len;
- unsigned long res_len_bytes = res_len * sizeof(u_int32_t);
+ unsigned long res_len_bytes = res_len * sizeof(uint32_t);
int err = 0;
- state->event_results = clCreateUserEvent(state->info->context, NULL);
+ cl_event e = clCreateUserEvent(state->info->context, NULL);
// Read back the results from the device to verify the output
- err = clEnqueueReadBuffer(state->queue, state->res_mem, CL_TRUE, 0, res_len_bytes, state->results, 0, NULL, &state->event_results);
+ err = clEnqueueReadBuffer(state->queue, state->res_mem, CL_TRUE, 0, res_len_bytes, state->results, 0, NULL, &e);
if (err != CL_SUCCESS)
{
printf("Error: Failed to read output array! %d\n", err);
exit(1);
}
- clSetEventCallback(state->event_results, CL_COMPLETE, callback_result, state);
+ clSetEventCallback(e, CL_COMPLETE, callback_result, state);
+ clReleaseEvent(e);
}
@@ -450,7 +451,7 @@ void CL_CALLBACK callback_kernel(cl_event event, cl_int event_command_status, vo
#define ORDER -1 // I think we need to do this, because we want to write it in the 'wrong' way
#define END 0
-#define GPUV_BIT_LENGTH 2048
+//#define GPUV_BIT_LENGTH 2048
#define BITS 64
@@ -474,7 +475,7 @@ void pk_to_mont(struct gpuv_public_key * pk) {
mpz_init_set_si(one,1);
mpz_set_si(one, 1);
- mpz_mul_2exp(r,one,GPUV_BIT_LENGTH); // r
+ mpz_mul_2exp(r,one,GPUV_BIT_LENGTH_2048); // r
mpz_import(mod, pk->len_n, ORDER, sizeof(gpu_register), END, 0, pk->n);
@@ -536,7 +537,7 @@ void sig_msg_to_mont(struct gpuv_signature_message * sig_msg) {
mpz_init_set_si(one,1);
mpz_set_si(one, 1);
- mpz_mul_2exp(r,one,GPUV_BIT_LENGTH); // r
+ mpz_mul_2exp(r,one,GPUV_BIT_LENGTH_2048); // r
mpz_import(mod, sig_msg->pubkey->len_n, ORDER, sizeof(gpu_register), END, 0, sig_msg->pubkey->n);
@@ -566,3 +567,128 @@ void sig_msg_to_mont(struct gpuv_signature_message * sig_msg) {
sig_msg->prepared = 1;
}
+
+gcry_sexp_t *sexp_from_string(char* str, const char *format) {
+
+ gcry_sexp_t *sexp = malloc(sizeof(gcry_sexp_t));
+
+ gcry_mpi_t mpi = gcry_mpi_new((int)strlen(str) * 8);
+
+ char * str_buf = malloc((GPUV_BIT_LENGTH_2048 / 8) * 2 + 8);
+
+ mpz_t m;
+ mpz_init(m);
+
+ mpz_import(m, 32, ORDER, sizeof(gpu_register), END, 0, str);
+
+ mpz_get_str(str_buf, 16, m);
+
+
+ gcry_mpi_scan(&mpi, GCRYMPI_FMT_HEX, str_buf, 0, NULL);
+
+ size_t errOff = 0;
+ gcry_sexp_build(sexp,&errOff,format,mpi);
+
+ free(str_buf);
+
+ return sexp;
+}
+
+gcry_sexp_t *sexp_from_string_key(char* str_1, char* str_1_buf, unsigned long str_2, const char *format) {
+
+ gcry_sexp_t *sexp = malloc(sizeof(gcry_sexp_t));
+
+ gcry_mpi_t mpi_1 = gcry_mpi_new((int)strlen(str_1) * 8);
+
+
+ mpz_t m;
+ mpz_init(m);
+
+ mpz_import(m, 32, ORDER, sizeof(gpu_register), END, 0, str_1);
+
+ mpz_get_str(str_1_buf, 16, m);
+
+ gcry_mpi_scan(&mpi_1, GCRYMPI_FMT_HEX, str_1_buf, 0, NULL);
+
+ gcry_mpi_t mpi_2 = gcry_mpi_set_ui(NULL, str_2);
+
+ size_t errOff = 0;
+ gcry_sexp_build(sexp,&errOff,format,mpi_1,mpi_2);
+
+
+
+ return sexp;
+}
+
+void cpu_verify(struct gpuv_batch *batch, struct gpuv_state *state) {
+
+ struct timespec p1, p2;
+
+ clock_gettime(CLOCK_REALTIME, &p1);
+
+ for (int j = 0; j < batch->current; j++) {
+
+ struct gpuv_public_key * p = batch->pairs[ j ].pubkey;
+
+ char * str_key_buf = malloc((GPUV_BIT_LENGTH_2048 / 8) * 2 + 8);
+
+ gcry_sexp_t * key_sexp = sexp_from_string_key(p->n, str_key_buf, p->e, "(public-key (rsa (n %m) (e %m)))" ); // pub key data
+
+ gcry_sexp_t *m_sexp = sexp_from_string(batch->pairs[j].m, "(data (flags raw) (value %m))"); // message format (for comparison)
+
+ gcry_sexp_t *s_sexp = sexp_from_string(batch->pairs[j].s, "(sig-val (rsa (s %m)))"); // signature format
+
+ if ( gcry_pk_verify(*s_sexp, *m_sexp, *key_sexp) == 0 ) {
+
+ uint32_t out_offset = j / (sizeof(uint32_t) * 8); // 32 bit
+
+ uint32_t mv = 1 << j;
+
+ state->results[out_offset] |= mv;
+
+ }
+
+ gcry_sexp_release(*m_sexp);
+ gcry_sexp_release(*s_sexp);
+ gcry_sexp_release(*key_sexp);
+
+ free(str_key_buf);
+
+ }
+
+ state->stale = 1;
+
+ int ret = 1;
+
+ unsigned long partial = state->sig_count / (sizeof(uint32_t) * 8);
+
+ for(int i = 0; i < state->results_len; i++) {
+
+ uint32_t mask = 0;
+
+ if (i >= partial) {
+ int remaining = state->sig_count % (sizeof(uint32_t) * 8);
+
+ for (int x = 0; x < remaining; x++) { mask |= 1 << x; }
+
+ } else {
+ mask = UINT32_MAX;
+ }
+
+ if (state->results[i] != mask) {
+ ret = 0;
+ }
+ }
+
+ state->info->in_progress = 0;
+ state->valid = ret;
+
+ clock_gettime(CLOCK_REALTIME, &p2);
+
+ state->t.tv_sec = ( p2.tv_nsec < p1.tv_nsec ? p2.tv_sec - (p1.tv_sec + 1) : p2.tv_sec - p1.tv_sec );
+ state->t.tv_nsec = ( p2.tv_nsec < p1.tv_nsec ? ((999999999 - p1.tv_nsec) + p2.tv_nsec) : (p2.tv_nsec - p1.tv_nsec) ) / 1000;
+
+ // pass results to user
+ state->cls(state->arg, state->valid, state->t, state->results_len, state->results);
+
+}
diff --git a/source/util.h b/source/util.h
@@ -32,7 +32,7 @@
#define NEED_LIBGCRYPT_VERSION "1.9.4"
typedef u_int64_t gpu_register;
-typedef u_int32_t DIGIT_T;
+typedef uint32_t DIGIT_T;
unsigned long gpuvt_estimate_pairs(void);
@@ -64,74 +64,6 @@ create_kernel (cl_program program,
// Implementations of structures, their fields should be private
-//struct gpu_info {
-//
-// cl_platform_id platform;
-// cl_device_id device_id;
-// cl_context context;
-// cl_command_queue commands;
-// cl_program program;
-// cl_kernel kernel;
-//
-//};
-//
-//struct gpu_state {
-//
-// cl_mem invalid;
-//
-// 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 pks_indices;
-//
-// struct timespec t1, t2;
-//
-// unsigned long total; // 0
-// unsigned long result; // 0
-//
-// bool skip;
-//
-//};
-//
-//struct gpu_state_alt {
-//
-// cl_mem x_mem;
-// cl_mem m_mem;
-// cl_mem res_mem;
-// cl_mem n_mem;
-// cl_mem ni_mem;
-// cl_mem msg_mem;
-// cl_mem exp_mem;
-//
-// cl_mem pks_indices;
-//
-// struct timespec t1, t2;
-//
-// struct timespec p;
-//
-// unsigned long total; // 0
-// u_int32_t *results;
-// unsigned long results_len;
-// int valid; // 0
-//
-// bool skip;
-//
-//};
-
-// new
-
-//union buf_ptr {
-// DIGIT_T *regular;
-// gpu_register *montgomery;
-//};
-
struct gpuv_public_key {
@@ -181,12 +113,12 @@ struct gpuv_batch {
unsigned long current;
// this is needed here, to make it easier to load it into the gpu
- u_int32_t *pk_indices;
+ uint32_t *pk_indices;
unsigned long pk_max_count;
unsigned long pk_current;
// a list, containing indices of pubkeys
- u_int32_t *pk_list;
+ uint32_t *pk_list;
};
@@ -198,7 +130,7 @@ struct gpuv_info {
cl_program program;
cl_kernel kernel;
- enum Variant variant;
+ enum GPUV_VARIANT variant;
int in_progress;
@@ -209,11 +141,8 @@ struct gpuv_state {
struct gpuv_info *info;
cl_command_queue queue;
- cl_event event_kernel;
- cl_event event_results;
-
// the callback the user defines
- void (*cls)(void *, int, struct timespec, unsigned long, u_int32_t *); // pointer to the results and how long the array is
+ void (*cls)(void *, int, struct timespec, unsigned long, uint32_t *); // pointer to the results and how long the array is
void * arg;
int stale; // says if we can enqueue it or not (0 at init)
@@ -240,7 +169,7 @@ struct gpuv_state {
int valid; // 0 in the beginning
- u_int32_t *results;
+ uint32_t *results;
unsigned long results_len;
};
@@ -252,4 +181,6 @@ void CL_CALLBACK callback_kernel(cl_event event, cl_int event_command_status, vo
void pk_to_mont(struct gpuv_public_key * pk);
void sig_msg_to_mont(struct gpuv_signature_message * sig_msg);
+void cpu_verify(struct gpuv_batch *batch, struct gpuv_state *state);
+
#endif /* util_h */
diff --git a/xcode/.DS_Store b/xcode/.DS_Store
Binary files differ.
diff --git a/xcode/lib-gpu-verify.xcodeproj/project.pbxproj b/xcode/lib-gpu-verify.xcodeproj/project.pbxproj
@@ -11,11 +11,9 @@
6A9DF2232B4B418400E368BE /* universal.c in Sources */ = {isa = PBXBuildFile; fileRef = 6A9DF2222B4B418400E368BE /* universal.c */; };
6AA38E5B2B0A97FC00E85243 /* main.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AA38E5A2B0A97FC00E85243 /* main.c */; };
6ABC2E842B231DFF00033B90 /* util.c in Sources */ = {isa = PBXBuildFile; fileRef = 6ABC2E832B231DFF00033B90 /* util.c */; };
- 6ABC2E882B231E3D00033B90 /* gpuv-ref.c in Sources */ = {isa = PBXBuildFile; fileRef = 6ABC2E862B231E3D00033B90 /* gpuv-ref.c */; };
6AC553252B2E174900046AB7 /* gpuv-montg.cl in Sources */ = {isa = PBXBuildFile; fileRef = 6AC553242B2E174900046AB7 /* gpuv-montg.cl */; };
6AC553292B2E17C800046AB7 /* gpuv-montg.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AC553282B2E17C800046AB7 /* gpuv-montg.c */; };
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */; };
- 6AF748832ADADF4500D58E08 /* gpuv.c in Sources */ = {isa = PBXBuildFile; fileRef = 6AF7487F2ADADF4500D58E08 /* gpuv.c */; };
C3770EFD0E6F1138009A5A77 /* OpenCL.framework in Frameworks */ = {isa = PBXBuildFile; fileRef = C3770EFC0E6F1138009A5A77 /* OpenCL.framework */; };
/* End PBXBuildFile section */
@@ -50,14 +48,10 @@
6AA38E612B0A9B2100E85243 /* lib-gpu-generate.entitlements */ = {isa = PBXFileReference; lastKnownFileType = text.plist.entitlements; path = "lib-gpu-generate.entitlements"; sourceTree = "<group>"; };
6ABC2E832B231DFF00033B90 /* util.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = util.c; path = ../source/util.c; sourceTree = "<group>"; };
6ABC2E852B231E0400033B90 /* util.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = util.h; path = ../source/util.h; sourceTree = "<group>"; };
- 6ABC2E862B231E3D00033B90 /* gpuv-ref.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "gpuv-ref.c"; path = "../source/gpuv-ref.c"; sourceTree = "<group>"; };
- 6ABC2E872B231E3D00033B90 /* gpuv-ref.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = "gpuv-ref.h"; path = "../source/gpuv-ref.h"; sourceTree = "<group>"; };
6AC553242B2E174900046AB7 /* gpuv-montg.cl */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.opencl; path = "gpuv-montg.cl"; sourceTree = "<group>"; };
6AC553272B2E17C800046AB7 /* gpuv-montg.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; name = "gpuv-montg.h"; path = "../source/gpuv-montg.h"; sourceTree = "<group>"; };
6AC553282B2E17C800046AB7 /* gpuv-montg.c */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.c; name = "gpuv-montg.c"; path = "../source/gpuv-montg.c"; sourceTree = "<group>"; };
6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = "lib-gpu-verify.c"; path = "../source/lib-gpu-verify.c"; sourceTree = "<group>"; };
- 6AF7487F2ADADF4500D58E08 /* gpuv.c */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.c; name = gpuv.c; path = ../source/gpuv.c; sourceTree = "<group>"; };
- 6AF748802ADADF4500D58E08 /* gpuv.h */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = sourcecode.c.h; name = gpuv.h; path = ../source/gpuv.h; sourceTree = "<group>"; };
C3770EFC0E6F1138009A5A77 /* OpenCL.framework */ = {isa = PBXFileReference; lastKnownFileType = wrapper.framework; name = OpenCL.framework; path = /System/Library/Frameworks/OpenCL.framework; sourceTree = "<absolute>"; };
/* End PBXFileReference section */
@@ -102,10 +96,8 @@
6A984F162AC5B18A00F530FD /* Headers */ = {
isa = PBXGroup;
children = (
- 6AF748802ADADF4500D58E08 /* gpuv.h */,
6ABC2E852B231E0400033B90 /* util.h */,
6A9DF2212B4B20F300E368BE /* universal.h */,
- 6ABC2E872B231E3D00033B90 /* gpuv-ref.h */,
6AC553272B2E17C800046AB7 /* gpuv-montg.h */,
);
name = Headers;
@@ -129,8 +121,6 @@
6A9DF2222B4B418400E368BE /* universal.c */,
6ABC2E832B231DFF00033B90 /* util.c */,
6AF748792ADADEBD00D58E08 /* lib-gpu-verify.c */,
- 6AF7487F2ADADF4500D58E08 /* gpuv.c */,
- 6ABC2E862B231E3D00033B90 /* gpuv-ref.c */,
6AC553282B2E17C800046AB7 /* gpuv-montg.c */,
);
name = Sources;
@@ -221,14 +211,12 @@
isa = PBXSourcesBuildPhase;
buildActionMask = 2147483647;
files = (
- 6ABC2E882B231E3D00033B90 /* gpuv-ref.c in Sources */,
6AC553252B2E174900046AB7 /* gpuv-montg.cl in Sources */,
6AC553292B2E17C800046AB7 /* gpuv-montg.c in Sources */,
6AF7487A2ADADEBD00D58E08 /* lib-gpu-verify.c in Sources */,
6A8A795F2A89672700116D7D /* gpuv.cl in Sources */,
6A9DF2232B4B418400E368BE /* universal.c in Sources */,
6ABC2E842B231DFF00033B90 /* util.c in Sources */,
- 6AF748832ADADF4500D58E08 /* gpuv.c in Sources */,
);
runOnlyForDeploymentPostprocessing = 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
@@ -4626,7 +4626,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "202"
endingLineNumber = "202"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -4719,7 +4719,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "253"
endingLineNumber = "253"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -4797,7 +4797,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "277"
endingLineNumber = "277"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -4875,7 +4875,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "248"
endingLineNumber = "248"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -4953,7 +4953,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "245"
endingLineNumber = "245"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -5196,8 +5196,8 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "164"
endingLineNumber = "164"
- landmarkName = "gpuv_test_montg()"
- landmarkType = "9">
+ landmarkName = "unknown"
+ landmarkType = "0">
<Locations>
<Location
uuid = "3ABB73B4-8A1A-4320-9629-FCC9139F5271 - 8470a8d1381f0aa6"
@@ -5274,7 +5274,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "216"
endingLineNumber = "216"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
<Locations>
<Location
@@ -5337,7 +5337,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "150"
endingLineNumber = "150"
- landmarkName = "finally(cls, valid, time, len, res)"
+ landmarkName = "gpuvt_read_files(variant, n_buf, msg_buf, exp_buf, s_buf, pks, n)"
landmarkType = "9">
<Locations>
<Location
@@ -5430,7 +5430,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "218"
endingLineNumber = "218"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
@@ -5700,7 +5700,7 @@
endingColumnNumber = "9223372036854775807"
startingLineNumber = "224"
endingLineNumber = "224"
- landmarkName = "gpuv_test_montg()"
+ landmarkName = "gpuv_test(variant)"
landmarkType = "9">
</BreakpointContent>
</BreakpointProxy>
