From 62135791bf16ec088912b4ab6d640cad864aa181 Mon Sep 17 00:00:00 2001 From: Agnes Leroy Date: Mon, 19 Feb 2024 11:48:58 +0100 Subject: [PATCH] chore(gpu): panic when polynomial size is not supported --- backends/tfhe-cuda-backend/cuda/src/device.cu | 20 +++--- .../cuda/src/integer/comparison.cu | 2 +- .../cuda/src/integer/integer.cu | 4 +- .../cuda/src/integer/multiplication.cu | 6 +- .../cuda/src/integer/scalar_comparison.cu | 2 +- .../cuda/src/integer/scalar_comparison.cuh | 4 +- .../cuda/src/pbs/bootstrap.cuh | 14 ++-- .../cuda/src/pbs/bootstrap_amortized.cu | 48 +++++-------- .../cuda/src/pbs/bootstrap_low_latency.cu | 71 ++++++------------- .../cuda/src/pbs/bootstrap_multibit.cu | 22 +++--- 10 files changed, 77 insertions(+), 116 deletions(-) diff --git a/backends/tfhe-cuda-backend/cuda/src/device.cu b/backends/tfhe-cuda-backend/cuda/src/device.cu index 8ec8bb549c..988c6b76be 100644 --- a/backends/tfhe-cuda-backend/cuda/src/device.cu +++ b/backends/tfhe-cuda-backend/cuda/src/device.cu @@ -56,7 +56,7 @@ void cuda_check_valid_malloc(uint64_t size, uint32_t gpu_index) { if (size > free_mem) { PANIC("Cuda error: not enough memory on device. " "Available: %zu vs Requested: %lu", - free_mem, size); + free_mem, size) } } @@ -79,7 +79,7 @@ void cuda_memcpy_async_to_gpu(void *dest, void *src, uint64_t size, cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, dest)); if (attr.device != stream->gpu_index && attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid device pointer in async copy to GPU."); + PANIC("Cuda error: invalid device pointer in async copy to GPU.") } check_cuda_error(cudaSetDevice(stream->gpu_index)); @@ -94,7 +94,7 @@ void cuda_memcpy_to_gpu(void *dest, void *src, uint64_t size) { cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, dest)); if (attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid device pointer in copy to GPU."); + PANIC("Cuda error: invalid device pointer in copy to GPU.") } check_cuda_error(cudaMemcpy(dest, src, size, cudaMemcpyHostToDevice)); } @@ -106,7 +106,7 @@ void cuda_memcpy_to_cpu(void *dest, void *src, uint64_t size) { cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, src)); if (attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid device pointer in copy to CPU."); + PANIC("Cuda error: invalid device pointer in copy to CPU.") } check_cuda_error(cudaMemcpy(dest, src, size, cudaMemcpyDeviceToHost)); } @@ -120,16 +120,16 @@ void cuda_memcpy_async_gpu_to_gpu(void *dest, void *src, uint64_t size, check_cuda_error(cudaPointerGetAttributes(&attr_dest, dest)); if (attr_dest.device != stream->gpu_index && attr_dest.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid dest device pointer in copy from GPU to GPU."); + PANIC("Cuda error: invalid dest device pointer in copy from GPU to GPU.") } cudaPointerAttributes attr_src; check_cuda_error(cudaPointerGetAttributes(&attr_src, src)); if (attr_src.device != stream->gpu_index && attr_src.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid src device pointer in copy from GPU to GPU."); + PANIC("Cuda error: invalid src device pointer in copy from GPU to GPU.") } if (attr_src.device != attr_dest.device) { - PANIC("Cuda error: different devices specified in copy from GPU to GPU."); + PANIC("Cuda error: different devices specified in copy from GPU to GPU.") } check_cuda_error(cudaSetDevice(stream->gpu_index)); @@ -150,7 +150,7 @@ void cuda_memset_async(void *dest, uint64_t val, uint64_t size, cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, dest)); if (attr.device != stream->gpu_index && attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid dest device pointer in cuda memset."); + PANIC("Cuda error: invalid dest device pointer in cuda memset.") } check_cuda_error(cudaSetDevice(stream->gpu_index)); check_cuda_error(cudaMemsetAsync(dest, val, size, stream->stream)); @@ -169,7 +169,7 @@ void cuda_set_value_async(cudaStream_t *stream, Torus *d_array, Torus value, cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, d_array)); if (attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid dest device pointer in cuda set value."); + PANIC("Cuda error: invalid dest device pointer in cuda set value.") } int block_size = 256; int num_blocks = (n + block_size - 1) / block_size; @@ -194,7 +194,7 @@ void cuda_memcpy_async_to_cpu(void *dest, const void *src, uint64_t size, cudaPointerAttributes attr; check_cuda_error(cudaPointerGetAttributes(&attr, src)); if (attr.device != stream->gpu_index && attr.type != cudaMemoryTypeDevice) { - PANIC("Cuda error: invalid src device pointer in copy to CPU async."); + PANIC("Cuda error: invalid src device pointer in copy to CPU async.") } check_cuda_error(cudaSetDevice(stream->gpu_index)); diff --git a/backends/tfhe-cuda-backend/cuda/src/integer/comparison.cu b/backends/tfhe-cuda-backend/cuda/src/integer/comparison.cu index 7662617add..2cda7dc004 100644 --- a/backends/tfhe-cuda-backend/cuda/src/integer/comparison.cu +++ b/backends/tfhe-cuda-backend/cuda/src/integer/comparison.cu @@ -70,7 +70,7 @@ void cuda_comparison_integer_radix_ciphertext_kb_64( static_cast(ksk), lwe_ciphertext_count); break; default: - PANIC("Cuda error: integer operation not supported"); + PANIC("Cuda error: integer operation not supported") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/integer/integer.cu b/backends/tfhe-cuda-backend/cuda/src/integer/integer.cu index 569fb1dda2..90c4bfc706 100644 --- a/backends/tfhe-cuda-backend/cuda/src/integer/integer.cu +++ b/backends/tfhe-cuda-backend/cuda/src/integer/integer.cu @@ -59,7 +59,9 @@ void cuda_full_propagation_64_inplace( ks_level, pbs_base_log, pbs_level, grouping_factor, num_blocks); break; default: - break; + PANIC("Cuda error (full propagation inplace): unsupported polynomial size. " + "Supported N's are powers of two" + " in the interval [256..16384].") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/integer/multiplication.cu b/backends/tfhe-cuda-backend/cuda/src/integer/multiplication.cu index 24dd6a08bb..772364fb97 100644 --- a/backends/tfhe-cuda-backend/cuda/src/integer/multiplication.cu +++ b/backends/tfhe-cuda-backend/cuda/src/integer/multiplication.cu @@ -24,7 +24,8 @@ void scratch_cuda_integer_mult_radix_ciphertext_kb_64( allocate_gpu_memory); break; default: - break; + PANIC("Cuda error (integer multiplication): unsupported polynomial size. " + "Only N = 2048 is supported") } } @@ -75,7 +76,8 @@ void cuda_integer_mult_radix_ciphertext_kb_64( num_blocks); break; default: - break; + PANIC("Cuda error (integer multiplication): unsupported polynomial size. " + "Only N = 2048 is supported") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cu b/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cu index 8301554c64..db05f43c84 100644 --- a/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cu +++ b/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cu @@ -36,6 +36,6 @@ void cuda_scalar_comparison_integer_radix_ciphertext_kb_64( static_cast(ksk), lwe_ciphertext_count, num_scalar_blocks); break; default: - PANIC("Cuda error: integer operation not supported"); + PANIC("Cuda error: integer operation not supported") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cuh b/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cuh index 26975176ae..56d4497320 100644 --- a/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cuh +++ b/backends/tfhe-cuda-backend/cuda/src/integer/scalar_comparison.cuh @@ -363,7 +363,7 @@ __host__ void host_integer_radix_scalar_equality_check_kb( msb_lut = mem_ptr->eq_buffer->is_non_zero_lut; break; default: - PANIC("Cuda error: integer operation not supported"); + PANIC("Cuda error: integer operation not supported") } host_compare_with_zero_equality(msb_stream, lwe_array_msb_out, msb, @@ -388,7 +388,7 @@ __host__ void host_integer_radix_scalar_equality_check_kb( num_halved_scalar_blocks + (num_msb_radix_blocks > 0)); break; default: - PANIC("Cuda error: integer operation not supported"); + PANIC("Cuda error: integer operation not supported") } // The result will be in the two first block. Everything else is diff --git a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap.cuh b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap.cuh index 6da130412a..9e4fe1abe9 100644 --- a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap.cuh +++ b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap.cuh @@ -20,7 +20,7 @@ void execute_pbs(cuda_stream_t *stream, Torus *lwe_array_out, // 32 bits switch (pbs_type) { case MULTI_BIT: - PANIC("Error: 32-bit multibit PBS is not supported.\n"); + PANIC("Error: 32-bit multibit PBS is not supported.\n") case LOW_LAT: cuda_bootstrap_low_latency_lwe_ciphertext_vector_32( stream, lwe_array_out, lwe_output_indexes, lut_vector, @@ -69,12 +69,12 @@ void execute_pbs(cuda_stream_t *stream, Torus *lwe_array_out, num_luts, lwe_idx, max_shared_memory); break; default: - PANIC("Error: unsupported cuda PBS type."); + PANIC("Error: unsupported cuda PBS type.") } break; default: PANIC("Cuda error: unsupported modulus size: only 32 and 64 bit integer " - "moduli are supported."); + "moduli are supported.") } } @@ -91,7 +91,7 @@ void execute_scratch_pbs(cuda_stream_t *stream, int8_t **pbs_buffer, // 32 bits switch (pbs_type) { case MULTI_BIT: - PANIC("Error: 32-bit multibit PBS is not supported.\n"); + PANIC("Error: 32-bit multibit PBS is not supported.\n") case LOW_LAT: scratch_cuda_bootstrap_low_latency_32( stream, pbs_buffer, glwe_dimension, polynomial_size, level_count, @@ -103,7 +103,7 @@ void execute_scratch_pbs(cuda_stream_t *stream, int8_t **pbs_buffer, input_lwe_ciphertext_count, max_shared_memory, allocate_gpu_memory); break; default: - PANIC("Error: unsupported cuda PBS type."); + PANIC("Error: unsupported cuda PBS type.") } break; case sizeof(uint64_t): @@ -126,11 +126,11 @@ void execute_scratch_pbs(cuda_stream_t *stream, int8_t **pbs_buffer, input_lwe_ciphertext_count, max_shared_memory, allocate_gpu_memory); break; default: - PANIC("Error: unsupported cuda PBS type."); + PANIC("Error: unsupported cuda PBS type.") } break; default: PANIC("Cuda error: unsupported modulus size: only 32 and 64 bit integer " - "moduli are supported."); + "moduli are supported.") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_amortized.cu b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_amortized.cu index 7b4ebee024..3ecbf8b2d1 100644 --- a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_amortized.cu +++ b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_amortized.cu @@ -11,28 +11,6 @@ uint64_t get_buffer_size_bootstrap_amortized_64( max_shared_memory); } -/* - * Runs standard checks to validate the inputs - */ -void checks_fast_bootstrap_amortized(int polynomial_size) { - assert( - ("Error (GPU amortized PBS): polynomial size should be one of 256, 512, " - "1024, 2048, 4096, 8192, 16384", - polynomial_size == 256 || polynomial_size == 512 || - polynomial_size == 1024 || polynomial_size == 2048 || - polynomial_size == 4096 || polynomial_size == 8192 || - polynomial_size == 16384)); -} - -/* - * Runs standard checks to validate the inputs - */ -void checks_bootstrap_amortized(int nbits, int base_log, int polynomial_size) { - assert(("Error (GPU amortized PBS): base log should be <= nbits", - base_log <= nbits)); - checks_fast_bootstrap_amortized(polynomial_size); -} - /* * This scratch function allocates the necessary amount of data on the GPU for * the amortized PBS on 32 bits inputs, into `pbs_buffer`. It also @@ -43,7 +21,6 @@ void scratch_cuda_bootstrap_amortized_32( cuda_stream_t *stream, int8_t **pbs_buffer, uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t input_lwe_ciphertext_count, uint32_t max_shared_memory, bool allocate_gpu_memory) { - checks_fast_bootstrap_amortized(polynomial_size); switch (polynomial_size) { case 256: @@ -82,7 +59,9 @@ void scratch_cuda_bootstrap_amortized_32( input_lwe_ciphertext_count, max_shared_memory, allocate_gpu_memory); break; default: - break; + PANIC("Cuda error (amortized PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } } @@ -96,7 +75,6 @@ void scratch_cuda_bootstrap_amortized_64( cuda_stream_t *stream, int8_t **pbs_buffer, uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t input_lwe_ciphertext_count, uint32_t max_shared_memory, bool allocate_gpu_memory) { - checks_fast_bootstrap_amortized(polynomial_size); switch (polynomial_size) { case 256: @@ -135,7 +113,9 @@ void scratch_cuda_bootstrap_amortized_64( input_lwe_ciphertext_count, max_shared_memory, allocate_gpu_memory); break; default: - break; + PANIC("Cuda error (amortized PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } } @@ -150,7 +130,9 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_32( uint32_t base_log, uint32_t level_count, uint32_t num_samples, uint32_t num_luts, uint32_t lwe_idx, uint32_t max_shared_memory) { - checks_bootstrap_amortized(32, base_log, polynomial_size); + if (base_log > 32) + PANIC("Cuda error (amortized PBS): base log should be > number of bits in " + "the ciphertext representation (32)"); switch (polynomial_size) { case 256: @@ -217,7 +199,9 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_32( max_shared_memory); break; default: - break; + PANIC("Cuda error (amortized PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } } @@ -294,7 +278,9 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_64( uint32_t base_log, uint32_t level_count, uint32_t num_samples, uint32_t num_luts, uint32_t lwe_idx, uint32_t max_shared_memory) { - checks_bootstrap_amortized(64, base_log, polynomial_size); + if (base_log > 64) + PANIC("Cuda error (amortized PBS): base log should be > number of bits in " + "the ciphertext representation (64)"); switch (polynomial_size) { case 256: @@ -361,7 +347,9 @@ void cuda_bootstrap_amortized_lwe_ciphertext_vector_64( max_shared_memory); break; default: - break; + PANIC("Cuda error (amortized PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_low_latency.cu b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_low_latency.cu index 4f180dc4ad..d8ad38195b 100644 --- a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_low_latency.cu +++ b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_low_latency.cu @@ -20,7 +20,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 512: if (verify_cuda_bootstrap_fast_low_latency_grid_size>( @@ -33,7 +32,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 1024: if (verify_cuda_bootstrap_fast_low_latency_grid_size>( @@ -46,7 +44,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 2048: if (verify_cuda_bootstrap_fast_low_latency_grid_size>( @@ -59,7 +56,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 4096: if (verify_cuda_bootstrap_fast_low_latency_grid_size>( @@ -72,7 +68,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 8192: if (verify_cuda_bootstrap_fast_low_latency_grid_size>( @@ -85,7 +80,6 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; case 16384: if (verify_cuda_bootstrap_fast_low_latency_grid_size< uint64_t, AmortizedDegree<16384>>(glwe_dimension, level_count, @@ -98,40 +92,13 @@ uint64_t get_buffer_size_bootstrap_low_latency_64( return get_buffer_size_bootstrap_low_latency( glwe_dimension, polynomial_size, level_count, input_lwe_ciphertext_count, max_shared_memory); - break; default: - return 0; - break; + PANIC("Cuda error (low latency PBS): unsupported polynomial size. " + "Supported N's are powers of two" + " in the interval [256..16384].") } } -/* - * Runs standard checks to validate the inputs - */ -void checks_fast_bootstrap_low_latency(int glwe_dimension, int level_count, - int polynomial_size, int num_samples) { - - assert(( - "Error (GPU low latency PBS): polynomial size should be one of 256, 512, " - "1024, 2048, 4096, 8192, 16384", - polynomial_size == 256 || polynomial_size == 512 || - polynomial_size == 1024 || polynomial_size == 2048 || - polynomial_size == 4096 || polynomial_size == 8192 || - polynomial_size == 16384)); -} - -/* - * Runs standard checks to validate the inputs - */ -void checks_bootstrap_low_latency(int nbits, int glwe_dimension, - int level_count, int base_log, - int polynomial_size, int num_samples) { - assert(("Error (GPU low latency PBS): base log should be <= nbits", - base_log <= nbits)); - checks_fast_bootstrap_low_latency(glwe_dimension, level_count, - polynomial_size, num_samples); -} - /* * This scratch function allocates the necessary amount of data on the GPU for * the low latency PBS on 32 bits inputs, into `pbs_buffer`. It also @@ -143,8 +110,6 @@ void scratch_cuda_bootstrap_low_latency_32( uint32_t polynomial_size, uint32_t level_count, uint32_t input_lwe_ciphertext_count, uint32_t max_shared_memory, bool allocate_gpu_memory) { - checks_fast_bootstrap_low_latency( - glwe_dimension, level_count, polynomial_size, input_lwe_ciphertext_count); switch (polynomial_size) { case 256: @@ -232,7 +197,9 @@ void scratch_cuda_bootstrap_low_latency_32( input_lwe_ciphertext_count, max_shared_memory, allocate_gpu_memory); break; default: - break; + PANIC("Cuda error (low latency PBS): unsupported polynomial size. " + "Supported N's are powers of two" + " in the interval [256..16384].") } } @@ -248,9 +215,6 @@ void scratch_cuda_bootstrap_low_latency_64( uint32_t input_lwe_ciphertext_count, uint32_t max_shared_memory, bool allocate_gpu_memory) { - checks_fast_bootstrap_low_latency( - glwe_dimension, level_count, polynomial_size, input_lwe_ciphertext_count); - switch (polynomial_size) { case 256: if (verify_cuda_bootstrap_fast_low_latency_grid_size 32) + PANIC("Cuda error (low latency PBS): base log should be > number of bits " + "in the ciphertext representation (32)"); switch (polynomial_size) { case 256: @@ -557,7 +524,9 @@ void cuda_bootstrap_low_latency_lwe_ciphertext_vector_32( num_luts, max_shared_memory); break; default: - break; + PANIC("Cuda error (low latency PBS): unsupported polynomial size. " + "Supported N's are powers of two" + " in the interval [256..16384].") } } @@ -644,8 +613,9 @@ void cuda_bootstrap_low_latency_lwe_ciphertext_vector_64( uint32_t lwe_dimension, uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t base_log, uint32_t level_count, uint32_t num_samples, uint32_t num_luts, uint32_t lwe_idx, uint32_t max_shared_memory) { - checks_bootstrap_low_latency(64, glwe_dimension, level_count, base_log, - polynomial_size, num_samples); + if (base_log > 64) + PANIC("Cuda error (low latency PBS): base log should be > number of bits " + "in the ciphertext representation (64)"); switch (polynomial_size) { case 256: @@ -829,8 +799,11 @@ void cuda_bootstrap_low_latency_lwe_ciphertext_vector_64( static_cast(bootstrapping_key), pbs_buffer, glwe_dimension, lwe_dimension, polynomial_size, base_log, level_count, num_samples, num_luts, max_shared_memory); - default: break; + default: + PANIC("Cuda error (low latency PBS): unsupported polynomial size. " + "Supported N's are powers of two" + " in the interval [256..16384].") } } diff --git a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_multibit.cu b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_multibit.cu index aca961155f..13cf2bc743 100644 --- a/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_multibit.cu +++ b/backends/tfhe-cuda-backend/cuda/src/pbs/bootstrap_multibit.cu @@ -3,16 +3,6 @@ #include "bootstrap_multibit.cuh" #include "bootstrap_multibit.h" -void checks_multi_bit_pbs(int polynomial_size) { - assert( - ("Error (GPU multi-bit PBS): polynomial size should be one of 256, 512, " - "1024, 2048, 4096, 8192, 16384", - polynomial_size == 256 || polynomial_size == 512 || - polynomial_size == 1024 || polynomial_size == 2048 || - polynomial_size == 4096 || polynomial_size == 8192 || - polynomial_size == 16384)); -} - void cuda_multi_bit_pbs_lwe_ciphertext_vector_64( cuda_stream_t *stream, void *lwe_array_out, void *lwe_output_indexes, void *lut_vector, void *lut_vector_indexes, void *lwe_array_in, @@ -22,7 +12,9 @@ void cuda_multi_bit_pbs_lwe_ciphertext_vector_64( uint32_t num_samples, uint32_t num_luts, uint32_t lwe_idx, uint32_t max_shared_memory, uint32_t lwe_chunk_size) { - checks_multi_bit_pbs(polynomial_size); + if (base_log > 64) + PANIC("Cuda error (multi-bit PBS): base log should be > number of bits in " + "the ciphertext representation (64)"); switch (polynomial_size) { case 256: @@ -229,7 +221,9 @@ void cuda_multi_bit_pbs_lwe_ciphertext_vector_64( } break; default: - break; + PANIC("Cuda error (multi-bit PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } } @@ -354,7 +348,9 @@ void scratch_cuda_multi_bit_pbs_64( } break; default: - break; + PANIC("Cuda error (multi-bit PBS): unsupported polynomial size. Supported " + "N's are powers of two" + " in the interval [256..16384].") } }