Skip to content

Commit

Permalink
Add sign_many to verify test
Browse files Browse the repository at this point in the history 
Create 100 keypairs and sample them in the messages to
sign with different keys on the GPU.

Check some number against CPU reference both comparing signature bits
and doing a verify.
  • Loading branch information
@sakridge
sakridge committed Nov 1, 2019
1 parent c00960e commit 9910c4a
Show file tree
Hide file tree
Showing 2 changed files with 84 additions and 41 deletions.
6 changes: 3 additions & 3 deletions src/cuda-ecc-ed25519/gpu_ctx.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -46,7 +46,7 @@ bool ed25519_init() {
gpu_ctx_t* get_gpu_ctx() {
int32_t cur_gpu, cur_queue;

LOG("locking global mutex");
LOG("locking global mutex\n");
pthread_mutex_lock(&g_ctx_mutex);
if (!cuda_crypt_init_locked()) {
pthread_mutex_unlock(&g_ctx_mutex);
Expand All @@ -62,7 +62,7 @@ gpu_ctx_t* get_gpu_ctx() {
pthread_mutex_unlock(&g_ctx_mutex);

gpu_ctx_t* cur_ctx = &g_gpu_ctx[cur_gpu][cur_queue];
LOG("locking contex mutex queue: %d gpu: %d", cur_queue, cur_gpu);
LOG("locking contex mutex queue: %d gpu: %d\n", cur_queue, cur_gpu);
pthread_mutex_lock(&cur_ctx->mutex);

CUDA_CHK(cudaSetDevice(cur_gpu));
Expand Down Expand Up @@ -122,7 +122,7 @@ void setup_gpu_ctx(verify_ctx_t* cur_ctx,
cur_ctx->offsets_len = total_signatures;
}

LOG("Done alloc");
LOG("Done alloc\n");

CUDA_CHK(cudaMemcpyAsync(cur_ctx->public_key_offsets, public_key_offsets, offsets_size, cudaMemcpyHostToDevice, stream));
CUDA_CHK(cudaMemcpyAsync(cur_ctx->signature_offsets, signature_offsets, offsets_size, cudaMemcpyHostToDevice, stream));
Expand Down
119 changes: 81 additions & 38 deletions src/cuda-ecc-ed25519/main.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@ void print_dwords(unsigned char* ptr, int size) {
typedef struct {
uint8_t signature[SIG_SIZE];
uint8_t public_key[PUB_KEY_SIZE];
uint8_t private_key[PRIV_KEY_SIZE];
uint32_t message_len;
uint8_t message[8];
} packet_t;
Expand All @@ -45,6 +46,7 @@ typedef struct {
uint32_t total_signatures;
uint32_t* message_lens;
uint32_t* public_key_offsets;
uint32_t* private_key_offsets;
uint32_t* signature_offsets;
uint32_t* message_start_offsets;
uint8_t* out_h;
Expand Down Expand Up @@ -151,9 +153,8 @@ int main(int argc, const char* argv[]) {

// Host allocate
unsigned char* seed_h = (unsigned char*)calloc(num_signatures_per_elem * SEED_SIZE, sizeof(uint32_t));
unsigned char* private_key_h = (unsigned char*)calloc(num_signatures_per_elem, PRIV_KEY_SIZE);
unsigned char message_h[] = "abcd1234";
int message_h_len = strlen((char*)message_h);
uint32_t message_h_len = strlen((char*)message_h);

uint32_t total_signatures = num_elems * num_signatures_per_elem;

Expand All @@ -169,10 +170,14 @@ int main(int argc, const char* argv[]) {
uint32_t* message_start_offsets = NULL;
ed25519_alloc(&message_start_offsets, total_signatures);

uint32_t* private_key_offsets = NULL;
ed25519_alloc(&private_key_offsets, total_signatures);

for (uint32_t i = 0; i < total_signatures; i++) {
uint32_t base_offset = i * sizeof(streamer_Packet);
signature_offsets[i] = base_offset + offsetof(packet_t, signature);
public_key_offsets[i] = base_offset + offsetof(packet_t, public_key);
private_key_offsets[i] = base_offset + offsetof(packet_t, private_key);
message_start_offsets[i] = base_offset + offsetof(packet_t, message);
message_lens[i] = message_h_len;
}
Expand All @@ -181,6 +186,7 @@ int main(int argc, const char* argv[]) {
vctx[i].message_lens = message_lens;
vctx[i].signature_offsets = signature_offsets;
vctx[i].public_key_offsets = public_key_offsets;
vctx[i].private_key_offsets = private_key_offsets;
vctx[i].message_start_offsets = message_start_offsets;
vctx[i].num_iterations = num_iterations;
vctx[i].use_non_default_stream = use_non_default_stream;
Expand All @@ -199,18 +205,10 @@ int main(int argc, const char* argv[]) {
total_packets += num_signatures_per_elem;
}

LOG("initing signatures..\n");
LOG("initing messages..\n");
for (int i = 0; i < num_signatures_per_elem; i++) {
packet_t* packet = (packet_t*)packets_h[i].data;
memcpy(packet->message, message_h, message_h_len);

LOG("message_len: %d\n",
message_h_len);
}

for (uint32_t i = 0; i < total_signatures; i++) {
LOG("sig_offset: %d pub_key_offset: %d message_start_offset: %d message_len: %d\n",
signature_offsets[i], public_key_offsets[i], message_start_offsets[i], message_lens[i]);
}

int out_size = total_signatures * sizeof(uint8_t);
Expand All @@ -222,42 +220,88 @@ int main(int argc, const char* argv[]) {
vctx[i].total_packets = total_packets;
}

LOG("creating seed..\n");
int ret = ed25519_create_seed(seed_h);
LOG("create_seed: %d\n", ret);
packet_t* first_packet_h = (packet_t*)packets_h[0].data;
ed25519_create_keypair(first_packet_h->public_key, private_key_h, seed_h);
ed25519_sign(first_packet_h->signature, first_packet_h->message, message_h_len, first_packet_h->public_key, private_key_h);
ret = ed25519_verify(first_packet_h->signature, message_h, message_h_len, first_packet_h->public_key);
LOG("verify: %d\n", ret);
LOG("creating %d keypairs..\n", num_signatures_per_elem);
int num_keypairs_to_create = std::min(100, num_signatures_per_elem);
uint8_t* public_keys = (uint8_t*)calloc(num_keypairs_to_create, PUB_KEY_SIZE);
uint8_t* private_keys = (uint8_t*)calloc(num_keypairs_to_create, PRIV_KEY_SIZE);
for (int i = 0; i < num_keypairs_to_create; i++) {
int ret = ed25519_create_seed(seed_h);
if (ret != 0) {
fprintf(stderr, "Invalid seed!");
exit(1);
}

for (int i = 1; i < num_signatures_per_elem; i++) {
packet_t* packet_h = (packet_t*)packets_h[i].data;
memcpy(packet_h->signature, first_packet_h->signature, SIG_SIZE);
memcpy(packet_h->public_key, first_packet_h->public_key, PUB_KEY_SIZE);
ed25519_create_keypair(&public_keys[PUB_KEY_SIZE * i], &private_keys[PRIV_KEY_SIZE * i], seed_h);
}

for (int i = 0; i < num_signatures_per_elem; i++ ) {
for (int i = 0; i < num_signatures_per_elem; i++) {
packet_t* packet = (packet_t*)packets_h[i].data;
int j = rand() % num_keypairs_to_create;
memcpy(packet->public_key, &public_keys[j * PUB_KEY_SIZE], PUB_KEY_SIZE);
memcpy(packet->private_key, &private_keys[j * PRIV_KEY_SIZE], PRIV_KEY_SIZE);
}

free(public_keys);
free(private_keys);

uint8_t* signatures_h = (uint8_t*)calloc(SIG_SIZE, total_signatures);

perftime_t start, end;
get_time(&start);

ed25519_sign_many(&vctx[0].elems_h[0],
vctx[0].num_elems,
sizeof(streamer_Packet),
vctx[0].total_packets,
vctx[0].total_signatures,
vctx[0].message_lens,
vctx[0].public_key_offsets,
vctx[0].private_key_offsets,
vctx[0].message_start_offsets,
signatures_h,
1);
get_time(&end);

double diff = get_diff(&start, &end);
printf("time diff: %f total: %d signs/sec: %f\n",
diff,
vctx[0].total_signatures,
(double)vctx[0].total_signatures / (diff / 1e6));

for (int i = 0; i < num_signatures_per_elem; i++) {
packet_t* packet_h = (packet_t*)packets_h[i].data;
unsigned char* sig_ptr = packet_h->signature;
unsigned char* messages_ptr = packet_h->message;
LOG("sig:");
print_dwords(sig_ptr, SIG_SIZE);
LOG("\nmessage: ");
print_dwords(messages_ptr, message_h_len);
LOG("\n\n");
memcpy(packet_h->signature, &signatures_h[i * SIG_SIZE], SIG_SIZE);
}
free(signatures_h);

int num_sigs_to_check = std::min(100, num_signatures_per_elem);
LOG("checking %d signatures\n", num_sigs_to_check);
for (int i = 0; i < num_sigs_to_check; i++) {
int j = rand() % num_signatures_per_elem;
packet_t* packet = (packet_t*)packets_h[j].data;

uint8_t signature[SIG_SIZE];
ed25519_sign(signature, message_h, message_h_len, packet->public_key, packet->private_key);
if (0 != memcmp(packet->signature, signature, SIG_SIZE)) {
fprintf(stderr, "Invalid signature!\n");
exit(1);
}

int ret = ed25519_verify(packet->signature, message_h, message_h_len, packet->public_key);
if (ret != 1) {
fprintf(stderr, "Invalid signature!\n");
exit(1);
}
}
LOG("\n");

std::vector<pthread_t> threads = std::vector<pthread_t>(num_threads);
pthread_attr_t attr;
ret = pthread_attr_init(&attr);
int ret = pthread_attr_init(&attr);
if (ret != 0) {
LOG("ERROR: pthread_attr_init: %d\n", ret);
return 1;
}

perftime_t start, end;
get_time(&start);
for (int i = 0; i < num_threads; i++) {
ret = pthread_create(&threads[i],
Expand All @@ -281,8 +325,8 @@ int main(int argc, const char* argv[]) {
get_time(&end);

int total = (num_threads * total_signatures * num_iterations);
double diff = get_diff(&start, &end);
printf("time diff: %f total: %d sigs/sec: %f\n",
diff = get_diff(&start, &end);
printf("time diff: %f total: %d verifies/sec: %f\n",
diff,
total,
(double)total / (diff / 1e6));
Expand All @@ -291,7 +335,6 @@ int main(int argc, const char* argv[]) {
LOG("ret:\n");
bool verify_failed = false;
for (int i = 0; i < out_size / (int)sizeof(uint8_t); i++) {
LOG("%x ", vctx[thread].out_h[i]);
if (vctx[thread].out_h[i] != 1) {
verify_failed = true;
}
Expand All @@ -306,12 +349,12 @@ int main(int argc, const char* argv[]) {
ed25519_free(message_lens);
ed25519_free(signature_offsets);
ed25519_free(public_key_offsets);
ed25519_free(private_key_offsets);
ed25519_free(message_start_offsets);
for (int thread = 0; thread < num_threads; thread++) {
ed25519_free(vctx[thread].out_h);
}
free(seed_h);
free(private_key_h);
ed25519_free_gpu_mem();
return 0;
}

0 comments on commit 9910c4a

Please sign in to comment.