chore(gpu): post hackathon cleanup

Author: agnesLeroy

backends/tfhe-cuda-backend/cuda/include/integer/integer.h

@@ -364,8 +364,7 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks, uint32_t message_modulus,
     uint32_t carry_modulus, PBS_TYPE pbs_type, uint32_t requested_flag,
-    uint32_t uses_carry, bool allocate_gpu_memory,
-    PBS_MS_REDUCTION_T noise_reduction_type);
+    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type);
 
 uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
@@ -374,8 +373,7 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks, uint32_t message_modulus,
     uint32_t carry_modulus, PBS_TYPE pbs_type, uint32_t requested_flag,
-    uint32_t uses_carry, bool allocate_gpu_memory,
-    PBS_MS_REDUCTION_T noise_reduction_type);
+    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type);
 
 void cuda_propagate_single_carry_kb_64_inplace(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,

backends/tfhe-cuda-backend/cuda/include/integer/integer_utilities.h

@@ -244,8 +244,6 @@ struct int_radix_params {
   uint32_t carry_modulus;
   PBS_MS_REDUCTION_T noise_reduction_type;
 
-  int_radix_params(){};
-
   int_radix_params(PBS_TYPE pbs_type, uint32_t glwe_dimension,
                    uint32_t polynomial_size, uint32_t big_lwe_dimension,
                    uint32_t small_lwe_dimension, uint32_t ks_level,
@@ -262,6 +260,8 @@ struct int_radix_params {
         message_modulus(message_modulus), carry_modulus(carry_modulus),
         noise_reduction_type(noise_reduction_type){};
 
+  int_radix_params() = default;
+
   void print() {
     printf("pbs_type: %u, glwe_dimension: %u, "
            "polynomial_size: %u, "
@@ -2450,8 +2450,7 @@ template <typename Torus> struct int_sc_prop_memory {
   int_sc_prop_memory(cudaStream_t const *streams, uint32_t const *gpu_indexes,
                      uint32_t gpu_count, int_radix_params params,
                      uint32_t num_radix_blocks, uint32_t requested_flag_in,
-                     uint32_t uses_carry, bool allocate_gpu_memory,
-                     uint64_t &size_tracker) {
+                     bool allocate_gpu_memory, uint64_t &size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     this->params = params;
     auto glwe_dimension = params.glwe_dimension;
@@ -3203,11 +3202,10 @@ template <typename Torus> struct int_mul_memory {
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
         2 * num_radix_blocks, vector_result_sb, small_lwe_vector, luts_array,
         true, allocate_gpu_memory, size_tracker);
-    uint32_t uses_carry = 0;
     uint32_t requested_flag = outputFlag::FLAG_NONE;
     sc_prop_mem = new int_sc_prop_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-        requested_flag, uses_carry, allocate_gpu_memory, size_tracker);
+        requested_flag, allocate_gpu_memory, size_tracker);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -3823,36 +3821,56 @@ template <typename Torus> struct int_comparison_eq_buffer {
     gpu_memory_allocated = allocate_gpu_memory;
     this->params = params;
     this->op = op;
+    Torus total_modulus = params.message_modulus * params.carry_modulus;
+    auto active_gpu_count = get_active_gpu_count(num_radix_blocks, gpu_count);
 
     are_all_block_true_buffer = new int_are_all_block_true_buffer<Torus>(
         streams, gpu_indexes, gpu_count, op, params, num_radix_blocks,
         allocate_gpu_memory, size_tracker);
 
-    // Operator LUT
-    auto operator_f = [op](Torus lhs, Torus rhs) -> Torus {
-      if (op == COMPARISON_TYPE::EQ) {
-        // EQ
-        return (lhs == rhs);
-      } else {
-        // NE
-        return (lhs != rhs);
-      }
-    };
-    operator_lut = new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count,
-                                            params, 1, num_radix_blocks,
-                                            allocate_gpu_memory, size_tracker);
+    if (op == COMPARISON_TYPE::EQ || COMPARISON_TYPE::NE) {
+      // Operator LUT
+      auto operator_f = [op](Torus lhs, Torus rhs) -> Torus {
+        if (op == COMPARISON_TYPE::EQ) {
+          return (lhs == rhs);
+        } else if (op == COMPARISON_TYPE::NE) {
+          return (lhs != rhs);
+          PANIC("Cuda error (eq/ne): invalid comparison type")
+        }
+      };
+      operator_lut = new int_radix_lut<Torus>(
+          streams, gpu_indexes, gpu_count, params, 1, num_radix_blocks,
+          allocate_gpu_memory, size_tracker);
 
-    generate_device_accumulator_bivariate<Torus>(
-        streams[0], gpu_indexes[0], operator_lut->get_lut(0, 0),
-        operator_lut->get_degree(0), operator_lut->get_max_degree(0),
-        params.glwe_dimension, params.polynomial_size, params.message_modulus,
-        params.carry_modulus, operator_f, gpu_memory_allocated);
+      generate_device_accumulator_bivariate<Torus>(
+          streams[0], gpu_indexes[0], operator_lut->get_lut(0, 0),
+          operator_lut->get_degree(0), operator_lut->get_max_degree(0),
+          params.glwe_dimension, params.polynomial_size, params.message_modulus,
+          params.carry_modulus, operator_f, gpu_memory_allocated);
 
-    auto active_gpu_count = get_active_gpu_count(num_radix_blocks, gpu_count);
-    operator_lut->broadcast_lut(streams, gpu_indexes, active_gpu_count);
+      operator_lut->broadcast_lut(streams, gpu_indexes, active_gpu_count);
+      // Scalar may have up to num_radix_blocks blocks
+      scalar_comparison_luts = new int_radix_lut<Torus>(
+          streams, gpu_indexes, gpu_count, params, total_modulus,
+          num_radix_blocks, allocate_gpu_memory, size_tracker);
+
+      for (int i = 0; i < total_modulus; i++) {
+        auto lut_f = [i, operator_f](Torus x) -> Torus {
+          return operator_f(i, x);
+        };
+
+        generate_device_accumulator<Torus>(
+            streams[0], gpu_indexes[0], scalar_comparison_luts->get_lut(0, i),
+            scalar_comparison_luts->get_degree(i),
+            scalar_comparison_luts->get_max_degree(i), params.glwe_dimension,
+            params.polynomial_size, params.message_modulus,
+            params.carry_modulus, lut_f, gpu_memory_allocated);
+      }
+      scalar_comparison_luts->broadcast_lut(streams, gpu_indexes,
+                                            active_gpu_count);
+    }
 
     // f(x) -> x == 0
-    Torus total_modulus = params.message_modulus * params.carry_modulus;
     auto is_non_zero_lut_f = [total_modulus](Torus x) -> Torus {
       return (x % total_modulus) != 0;
     };
@@ -3868,38 +3886,20 @@ template <typename Torus> struct int_comparison_eq_buffer {
         params.carry_modulus, is_non_zero_lut_f, gpu_memory_allocated);
 
     is_non_zero_lut->broadcast_lut(streams, gpu_indexes, active_gpu_count);
-
-    // Scalar may have up to num_radix_blocks blocks
-    scalar_comparison_luts = new int_radix_lut<Torus>(
-        streams, gpu_indexes, gpu_count, params, total_modulus,
-        num_radix_blocks, allocate_gpu_memory, size_tracker);
-
-    for (int i = 0; i < total_modulus; i++) {
-      auto lut_f = [i, operator_f](Torus x) -> Torus {
-        return operator_f(i, x);
-      };
-
-      generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], scalar_comparison_luts->get_lut(0, i),
-          scalar_comparison_luts->get_degree(i),
-          scalar_comparison_luts->get_max_degree(i), params.glwe_dimension,
-          params.polynomial_size, params.message_modulus, params.carry_modulus,
-          lut_f, gpu_memory_allocated);
-    }
-    scalar_comparison_luts->broadcast_lut(streams, gpu_indexes,
-                                          active_gpu_count);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
                uint32_t gpu_count) {
-    operator_lut->release(streams, gpu_indexes, gpu_count);
-    delete operator_lut;
-    is_non_zero_lut->release(streams, gpu_indexes, gpu_count);
-    delete is_non_zero_lut;
-    scalar_comparison_luts->release(streams, gpu_indexes, gpu_count);
-    delete scalar_comparison_luts;
+    if (op == COMPARISON_TYPE::EQ || COMPARISON_TYPE::NE) {
+      operator_lut->release(streams, gpu_indexes, gpu_count);
+      delete operator_lut;
+      scalar_comparison_luts->release(streams, gpu_indexes, gpu_count);
+      delete scalar_comparison_luts;
+    }
     are_all_block_true_buffer->release(streams, gpu_indexes, gpu_count);
     delete are_all_block_true_buffer;
+    is_non_zero_lut->release(streams, gpu_indexes, gpu_count);
+    delete is_non_zero_lut;
   }
 };
 
@@ -4020,8 +4020,7 @@ template <typename Torus> struct int_comparison_diff_buffer {
       case LE:
         return (x == IS_INFERIOR) || (x == IS_EQUAL);
       default:
-        // We don't need a default case but we need to return something
-        return 42;
+        PANIC("Cuda error (comparisons): unknown comparison type")
       }
     };
 
@@ -5069,11 +5068,10 @@ template <typename Torus> struct int_scalar_mul_buffer {
           streams, gpu_indexes, gpu_count, params, num_radix_blocks,
           num_ciphertext_bits, true, allocate_gpu_memory, last_step_mem);
     }
-    uint32_t uses_carry = 0;
     uint32_t requested_flag = outputFlag::FLAG_NONE;
     sc_prop_mem = new int_sc_prop_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-        requested_flag, uses_carry, allocate_gpu_memory, last_step_mem);
+        requested_flag, allocate_gpu_memory, last_step_mem);
     if (anticipated_buffer_drop) {
       size_tracker += std::max(anticipated_drop_mem, last_step_mem);
     } else {
@@ -5132,10 +5130,9 @@ template <typename Torus> struct int_abs_buffer {
         streams, gpu_indexes, gpu_count, SHIFT_OR_ROTATE_TYPE::RIGHT_SHIFT,
         params, num_radix_blocks, allocate_gpu_memory, size_tracker);
     uint32_t requested_flag = outputFlag::FLAG_NONE;
-    uint32_t uses_carry = 0;
     scp_mem = new int_sc_prop_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-        requested_flag, uses_carry, allocate_gpu_memory, size_tracker);
+        requested_flag, allocate_gpu_memory, size_tracker);
     bitxor_mem = new int_bitop_buffer<Torus>(
         streams, gpu_indexes, gpu_count, BITOP_TYPE::BITXOR, params,
         num_radix_blocks, allocate_gpu_memory, size_tracker);
@@ -5216,13 +5213,12 @@ template <typename Torus> struct int_div_rem_memory {
                                             params, num_blocks,
                                             allocate_gpu_memory, size_tracker);
       uint32_t requested_flag = outputFlag::FLAG_NONE;
-      uint32_t uses_carry = 0;
       scp_mem_1 = new int_sc_prop_memory<Torus>(
           streams, gpu_indexes, gpu_count, params, num_blocks, requested_flag,
-          uses_carry, allocate_gpu_memory, size_tracker);
+          allocate_gpu_memory, size_tracker);
       scp_mem_2 = new int_sc_prop_memory<Torus>(
           streams, gpu_indexes, gpu_count, params, num_blocks, requested_flag,
-          uses_carry, allocate_gpu_memory, size_tracker);
+          allocate_gpu_memory, size_tracker);
 
       std::function<uint64_t(uint64_t)> quotient_predicate_lut_f =
           [](uint64_t x) -> uint64_t { return x == 1; };
@@ -5420,7 +5416,7 @@ template <typename Torus> struct int_sub_and_propagate {
 
     this->sc_prop_mem = new int_sc_prop_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-        requested_flag_in, (uint32_t)0, allocate_gpu_memory, size_tracker);
+        requested_flag_in, allocate_gpu_memory, size_tracker);
 
     this->neg_rhs_array = new CudaRadixCiphertextFFI;
     create_zero_radix_ciphertext_async<Torus>(
@@ -5559,7 +5555,7 @@ template <typename Torus> struct int_unsigned_scalar_div_mem {
             scalar_divisor_ffi->active_bits, allocate_gpu_memory, size_tracker);
         scp_mem = new int_sc_prop_memory<Torus>(
             streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-            FLAG_NONE, (uint32_t)0, allocate_gpu_memory, size_tracker);
+            FLAG_NONE, allocate_gpu_memory, size_tracker);
         sub_and_propagate_mem = new int_sub_and_propagate<Torus>(
             streams, gpu_indexes, gpu_count, params, num_radix_blocks,
             FLAG_NONE, allocate_gpu_memory, size_tracker);
@@ -5716,7 +5712,7 @@ template <typename Torus> struct int_signed_scalar_div_mem {
               num_radix_blocks, allocate_gpu_memory, size_tracker);
           scp_mem = new int_sc_prop_memory<Torus>(
               streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-              FLAG_NONE, (uint32_t)0, allocate_gpu_memory, size_tracker);
+              FLAG_NONE, allocate_gpu_memory, size_tracker);
 
         } else {
 
@@ -5737,7 +5733,7 @@ template <typename Torus> struct int_signed_scalar_div_mem {
           if (scalar_divisor_ffi->is_chosen_multiplier_geq_two_pow_numerator) {
             scp_mem = new int_sc_prop_memory<Torus>(
                 streams, gpu_indexes, gpu_count, params, num_radix_blocks,
-                FLAG_NONE, (uint32_t)0, allocate_gpu_memory, size_tracker);
+                FLAG_NONE, allocate_gpu_memory, size_tracker);
           }
         }
       }
@@ -5885,7 +5881,7 @@ template <typename Torus> struct int_signed_scalar_div_rem_buffer {
 
     this->scp_mem = new int_sc_prop_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks, FLAG_NONE,
-        (uint32_t)0, allocate_gpu_memory, size_tracker);
+        allocate_gpu_memory, size_tracker);
 
     bool is_divisor_one = scalar_divisor_ffi->is_abs_divisor_one &&
                           !scalar_divisor_ffi->is_divisor_negative;
@@ -6085,7 +6081,7 @@ template <typename Torus> struct int_count_of_consecutive_bits_buffer {
         num_radix_blocks, true, allocate_gpu_memory, size_tracker);
 
     this->propagate_mem = new int_sc_prop_memory<Torus>(
-        streams, gpu_indexes, gpu_count, params, counter_num_blocks, 0, 0,
+        streams, gpu_indexes, gpu_count, params, counter_num_blocks, 0,
         allocate_gpu_memory, size_tracker);
   }
 

backends/tfhe-cuda-backend/cuda/src/integer/comparison.cuh

@@ -439,7 +439,6 @@ __host__ void tree_sign_reduction(
           "than the number of blocks to operate on")
 
   auto params = tree_buffer->params;
-  auto big_lwe_dimension = params.big_lwe_dimension;
   auto glwe_dimension = params.glwe_dimension;
   auto polynomial_size = params.polynomial_size;
   auto message_modulus = params.message_modulus;

backends/tfhe-cuda-backend/cuda/src/integer/integer.cu

@@ -53,8 +53,7 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks, uint32_t message_modulus,
     uint32_t carry_modulus, PBS_TYPE pbs_type, uint32_t requested_flag,
-    uint32_t uses_carry, bool allocate_gpu_memory,
-    PBS_MS_REDUCTION_T noise_reduction_type) {
+    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type) {
   int_radix_params params(pbs_type, glwe_dimension, polynomial_size,
                           big_lwe_dimension, small_lwe_dimension, ks_level,
                           ks_base_log, pbs_level, pbs_base_log, grouping_factor,
@@ -63,7 +62,7 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
   return scratch_cuda_propagate_single_carry_kb_inplace<uint64_t>(
       (cudaStream_t *)(streams), gpu_indexes, gpu_count,
       (int_sc_prop_memory<uint64_t> **)mem_ptr, num_blocks, params,
-      requested_flag, uses_carry, allocate_gpu_memory);
+      requested_flag, allocate_gpu_memory);
 }
 
 uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
@@ -73,8 +72,7 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks, uint32_t message_modulus,
     uint32_t carry_modulus, PBS_TYPE pbs_type, uint32_t requested_flag,
-    uint32_t uses_carry, bool allocate_gpu_memory,
-    PBS_MS_REDUCTION_T noise_reduction_type) {
+    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type) {
   int_radix_params params(pbs_type, glwe_dimension, polynomial_size,
                           big_lwe_dimension, small_lwe_dimension, ks_level,
                           ks_base_log, pbs_level, pbs_base_log, grouping_factor,
@@ -83,7 +81,7 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
   return scratch_cuda_propagate_single_carry_kb_inplace<uint64_t>(
       (cudaStream_t *)(streams), gpu_indexes, gpu_count,
       (int_sc_prop_memory<uint64_t> **)mem_ptr, num_blocks, params,
-      requested_flag, uses_carry, allocate_gpu_memory);
+      requested_flag, allocate_gpu_memory);
 }
 
 uint64_t scratch_cuda_integer_overflowing_sub_kb_64_inplace(

backends/tfhe-cuda-backend/cuda/src/integer/integer.cuh

@@ -243,8 +243,7 @@ __host__ void host_radix_cumulative_sum_in_groups(cudaStream_t stream,
   auto lwe_size = dest->lwe_dimension + 1;
   cuda_set_device(gpu_index);
   // Each CUDA block is responsible for a single group
-  int num_blocks = (num_radix_blocks + group_size - 1) / group_size,
-      num_threads = 512;
+  int num_blocks = CEIL_DIV(num_radix_blocks, group_size), num_threads = 512;
   device_radix_cumulative_sum_in_groups<Torus>
       <<<num_blocks, num_threads, 0, stream>>>(
           (Torus *)dest->ptr, (Torus *)src->ptr, num_radix_blocks, lwe_size,
@@ -1574,9 +1573,6 @@ void host_full_propagate_inplace(
     void *const *bsks, uint32_t num_blocks) {
   auto params = mem_ptr->lut->params;
 
-  int big_lwe_size = (params.glwe_dimension * params.polynomial_size + 1);
-  int small_lwe_size = (params.small_lwe_dimension + 1);
-
   // In the case of extracting a single LWE this parameters are dummy
   uint32_t num_many_lut = 1;
   uint32_t lut_stride = 0;
@@ -1990,12 +1986,12 @@ uint64_t scratch_cuda_propagate_single_carry_kb_inplace(
     cudaStream_t const *streams, uint32_t const *gpu_indexes,
     uint32_t gpu_count, int_sc_prop_memory<Torus> **mem_ptr,
     uint32_t num_radix_blocks, int_radix_params params, uint32_t requested_flag,
-    uint32_t uses_carry, bool allocate_gpu_memory) {
+    bool allocate_gpu_memory) {
   PUSH_RANGE("scratch add & propagate sc")
   uint64_t size_tracker = 0;
   *mem_ptr = new int_sc_prop_memory<Torus>(
       streams, gpu_indexes, gpu_count, params, num_radix_blocks, requested_flag,
-      uses_carry, allocate_gpu_memory, size_tracker);
+      allocate_gpu_memory, size_tracker);
   POP_RANGE()
   return size_tracker;
 }
@@ -2141,9 +2137,6 @@ void host_add_and_propagate_single_carry(
 
   auto num_radix_blocks = lhs_array->num_radix_blocks;
   auto params = mem->params;
-  auto glwe_dimension = params.glwe_dimension;
-  auto polynomial_size = params.polynomial_size;
-  uint32_t big_lwe_size = glwe_dimension * polynomial_size + 1;
   auto lut_stride = mem->lut_stride;
   auto num_many_lut = mem->num_many_lut;
   CudaRadixCiphertextFFI output_flag;
@@ -2416,7 +2409,6 @@ __host__ void integer_radix_apply_noise_squashing_kb(
 
   PUSH_RANGE("apply noise squashing")
   auto params = lut->params;
-  auto pbs_type = params.pbs_type;
   auto big_lwe_dimension = params.big_lwe_dimension;
   auto small_lwe_dimension = params.small_lwe_dimension;
   auto ks_level = params.ks_level;