From f04a29aea4770bfc3007013e2454a15f5f19b4f6 Mon Sep 17 00:00:00 2001 From: Agnes Leroy Date: Mon, 28 Nov 2022 11:56:29 +0100 Subject: [PATCH] chore(cuda): fix asserts regarding the base log value --- src/bootstrap_amortized.cu | 4 ++-- src/bootstrap_low_latency.cu | 8 ++++---- src/bootstrap_wop.cu | 12 ++++++------ src/negation.cuh | 1 - 4 files changed, 12 insertions(+), 13 deletions(-) diff --git a/src/bootstrap_amortized.cu b/src/bootstrap_amortized.cu index a6997fbc2..d6147ae6c 100644 --- a/src/bootstrap_amortized.cu +++ b/src/bootstrap_amortized.cu @@ -65,7 +65,7 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_32( uint32_t num_lut_vectors, uint32_t lwe_idx, uint32_t max_shared_memory) { assert( - ("Error (GPU amortized PBS): base log should be <= 16", base_log <= 16)); + ("Error (GPU amortized PBS): base log should be <= 32", base_log <= 32)); assert(("Error (GPU amortized PBS): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU amortized PBS): polynomial size should be one of 512, " @@ -123,7 +123,7 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_64( uint32_t num_lut_vectors, uint32_t lwe_idx, uint32_t max_shared_memory) { assert( - ("Error (GPU amortized PBS): base log should be <= 16", base_log <= 16)); + ("Error (GPU amortized PBS): base log should be <= 64", base_log <= 64)); assert(("Error (GPU amortized PBS): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU amortized PBS): polynomial size should be one of 512, " diff --git a/src/bootstrap_low_latency.cu b/src/bootstrap_low_latency.cu index e1a12e549..17e4ca4db 100644 --- a/src/bootstrap_low_latency.cu +++ b/src/bootstrap_low_latency.cu @@ -63,8 +63,8 @@ void cuda_bootstrap_low_latency_lwe_ciphertext_vector_32( uint32_t base_log, uint32_t level_count, uint32_t num_samples, uint32_t num_lut_vectors, uint32_t lwe_idx, uint32_t max_shared_memory) { - assert(("Error (GPU low latency PBS): base log should be <= 16", - base_log <= 16)); + assert(("Error (GPU low latency PBS): base log should be <= 32", + base_log <= 32)); assert(("Error (GPU low latency PBS): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU low latency PBS): polynomial size should be one of 512, " @@ -116,8 +116,8 @@ void cuda_bootstrap_low_latency_lwe_ciphertext_vector_64( uint32_t base_log, uint32_t level_count, uint32_t num_samples, uint32_t num_lut_vectors, uint32_t lwe_idx, uint32_t max_shared_memory) { - assert(("Error (GPU low latency PBS): base log should be <= 16", - base_log <= 16)); + assert(("Error (GPU low latency PBS): base log should be <= 64", + base_log <= 64)); assert(("Error (GPU low latency PBS): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU low latency PBS): polynomial size should be one of 512, " diff --git a/src/bootstrap_wop.cu b/src/bootstrap_wop.cu index 9411d407a..8041f61a7 100644 --- a/src/bootstrap_wop.cu +++ b/src/bootstrap_wop.cu @@ -6,7 +6,7 @@ void cuda_cmux_tree_32(void *v_stream, void *glwe_array_out, void *ggsw_in, uint32_t level_count, uint32_t r, uint32_t max_shared_memory) { - assert(("Error (GPU Cmux tree): base log should be <= 16", base_log <= 16)); + assert(("Error (GPU Cmux tree): base log should be <= 32", base_log <= 32)); assert(("Error (GPU Cmux tree): polynomial size should be one of 512, 1024, " "2048, 4096, 8192", polynomial_size == 512 || polynomial_size == 1024 || @@ -61,7 +61,7 @@ void cuda_cmux_tree_64(void *v_stream, void *glwe_array_out, void *ggsw_in, uint32_t level_count, uint32_t r, uint32_t max_shared_memory) { - assert(("Error (GPU Cmux tree): base log should be <= 16", base_log <= 16)); + assert(("Error (GPU Cmux tree): base log should be <= 64", base_log <= 64)); assert(("Error (GPU Cmux tree): polynomial size should be one of 512, 1024, " "2048, 4096, 8192", polynomial_size == 512 || polynomial_size == 1024 || @@ -119,8 +119,8 @@ void cuda_extract_bits_32( uint32_t lwe_dimension_out, uint32_t glwe_dimension, uint32_t base_log_bsk, uint32_t level_count_bsk, uint32_t base_log_ksk, uint32_t level_count_ksk, uint32_t number_of_samples) { - assert(("Error (GPU extract bits): base log should be <= 16", - base_log_bsk <= 16)); + assert(("Error (GPU extract bits): base log should be <= 32", + base_log_bsk <= 32)); assert(("Error (GPU extract bits): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU extract bits): lwe_dimension_in should be one of " @@ -190,8 +190,8 @@ void cuda_extract_bits_64( uint32_t lwe_dimension_out, uint32_t glwe_dimension, uint32_t base_log_bsk, uint32_t level_count_bsk, uint32_t base_log_ksk, uint32_t level_count_ksk, uint32_t number_of_samples) { - assert(("Error (GPU extract bits): base log should be <= 16", - base_log_bsk <= 16)); + assert(("Error (GPU extract bits): base log should be <= 64", + base_log_bsk <= 64)); assert(("Error (GPU extract bits): glwe_dimension should be equal to 1", glwe_dimension == 1)); assert(("Error (GPU extract bits): lwe_dimension_in should be one of " diff --git a/src/negation.cuh b/src/negation.cuh index fd701ea3f..301fddfa0 100644 --- a/src/negation.cuh +++ b/src/negation.cuh @@ -11,7 +11,6 @@ #include "linear_algebra.h" #include "utils/kernel_dimensions.cuh" - template __global__ void negation(T *output, T *input, uint32_t num_entries) {