feat: Support partial bootstrapping pipeline with level_start=-1 and …

.gitignore

@@ -28,4 +28,6 @@ CTestTestfile.cmake
 *.vcxproj.filters
 *.DS_Store
 
-install/
+install/
+
+*.txt

example/bootstrapping/5_ckks_regular_bootstrapping_v2.cu

@@ -7,16 +7,23 @@
 #include "ckks/precision.cuh"
 #include "../example_util.h"
 
-int main(int argc, char* argv[])
+void run_bootstrapping(size_t poly_modulus_degree, int slot_count,
+                       const std::string& label)
 {
-    cudaSetDevice(0); 
+    std::cout << "\n========================================" << std::endl;
+    std::cout << label << std::endl;
+    std::cout << "  N = " << poly_modulus_degree
+              << ", slot_count = " << slot_count
+              << ", gap = " << (poly_modulus_degree / 2) / slot_count
+              << std::endl;
+    std::cout << "========================================" << std::endl;
 
     heongpu::HEContext<heongpu::Scheme::CKKS> context(
         heongpu::keyswitching_type::KEYSWITCHING_METHOD_II,
         heongpu::sec_level_type::none);
-    size_t poly_modulus_degree = 1 << 16;
     context.set_poly_modulus_degree(poly_modulus_degree);
-
+    context.set_slot_count(slot_count);
+
     context.set_coeff_modulus_values(
         {
             0x10000000006e0001, // 60 Q0
@@ -57,7 +64,7 @@ int main(int argc, char* argv[])
     context.print_parameters();
 
     int h = 192;
-    int ephemeral_secret_weight = 32;  
+    int ephemeral_secret_weight = 32;
 
     double scale = pow(2.0, 40);
 
@@ -97,7 +104,6 @@ int main(int argc, char* argv[])
     heongpu::HEArithmeticOperator<heongpu::Scheme::CKKS> operators(context,
                                                                    encoder);
 
-    const int slot_count = poly_modulus_degree / 2;
     std::vector<Complex64> message;
     for (int i = 0; i < slot_count; i++)
     {
@@ -110,24 +116,26 @@ int main(int argc, char* argv[])
     heongpu::Ciphertext<heongpu::Scheme::CKKS> C1(context);
     encryptor.encrypt(C1, P1);
 
-    heongpu::EvalModConfig eval_mod_config(20);
-
+    heongpu::EvalModConfig eval_mod_config(context.get_key_modulus()[0].value,
+                                           20, 256.0, 16, 30, 3, 0,
+                                           pow(2.0, 60));
+
     heongpu::BootstrappingConfigV2 boot_config(
         heongpu::EncodingMatrixConfig(
-            heongpu::LinearTransformType::SLOTS_TO_COEFFS, 12),
+            heongpu::LinearTransformType::SLOTS_TO_COEFFS, 12, 2.0, 3),
         eval_mod_config,
         heongpu::EncodingMatrixConfig(
-            heongpu::LinearTransformType::COEFFS_TO_SLOTS, 24));
+            heongpu::LinearTransformType::COEFFS_TO_SLOTS, 24, 2.0, 4));
 
     operators.generate_bootstrapping_params_v2(scale, boot_config);
 
     std::vector<int> key_index = operators.bootstrapping_key_indexs();
     std::cout << "Total galois key needed for CKKS bootstrapping: "
               << key_index.size() << std::endl;
     heongpu::Galoiskey<heongpu::Scheme::CKKS> galois_key(context, key_index);
-    keygen.generate_galois_key(galois_key, secret_key); 
+    keygen.generate_galois_key(galois_key, secret_key);
 
-    // Drop all level until one level remain
+    // Drop all levels until one level remains
     for (int i = 0; i < 24; i++)
     {
         operators.mod_drop_inplace(C1);
@@ -175,6 +183,26 @@ int main(int argc, char* argv[])
                 << " - ACTUAL:" << decrypted_1[j] << std::endl;
     }
     std::cout << std::endl;
-
+
+    delete swk_dense_to_sparse;
+    delete swk_sparse_to_dense;
+}
+
+int main(int argc, char* argv[])
+{
+    cudaSetDevice(0);
+
+    size_t poly_modulus_degree = 1 << 16;
+
+    // Full packing: slot_count = N/2, gap = 1
+    run_bootstrapping(poly_modulus_degree, poly_modulus_degree / 2,
+                      "Full Packing Bootstrapping");
+
+    cudaDeviceSynchronize();
+
+    // Sparse packing: slot_count = N/8, gap = 4
+    run_bootstrapping(poly_modulus_degree, poly_modulus_degree / 8,
+                      "Sparse Packing Bootstrapping");
+
     return EXIT_SUCCESS;
 }

src/heongpu/include/host/ckks/context.cuh

@@ -55,6 +55,8 @@ namespace heongpu
         void set_coeff_modulus_values(const std::vector<Data64>& log_Q_bases,
                                       const std::vector<Data64>& log_P_bases);
 
+        void set_slot_count(int slot_count); // @company CipherFlow
+
         void generate();
 
         void print_parameters();
@@ -66,6 +68,8 @@ namespace heongpu
             return n_power;
         }
 
+        inline int get_slot_count() const noexcept { return slot_count; } // @company CipherFlow
+
         inline int get_ciphertext_modulus_count() const noexcept
         {
             return Q_size;
@@ -91,6 +95,7 @@ namespace heongpu
         bool poly_modulus_degree_specified_ = false;
         bool coeff_modulus_specified_ = false;
         bool context_generated_ = false;
+        bool slot_count_specified_ = false; // @company CipherFlow
 
         scheme_type scheme_;
         sec_level_type sec_level_;
@@ -99,6 +104,9 @@ namespace heongpu
         int n;
         int n_power;
 
+        int slot_count; // @company CipherFlow
+        int log_slot_count; // @company CipherFlow
+
         int coeff_modulus;
         int total_coeff_bit_count;
 
@@ -121,6 +129,10 @@ namespace heongpu
         std::shared_ptr<DeviceVector<Root64>> ntt_table_;
         std::shared_ptr<DeviceVector<Root64>> intt_table_;
         std::shared_ptr<DeviceVector<Ninverse64>> n_inverse_;
+
+        std::shared_ptr<DeviceVector<Root64>> ntt_table_slot_; // @company CipherFlow
+        std::shared_ptr<DeviceVector<Root64>> ntt_table_dslot_; // @company CipherFlow
+
         std::shared_ptr<DeviceVector<Data64>> last_q_modinv_;
         std::shared_ptr<DeviceVector<Data64>> half_p_;
         std::shared_ptr<DeviceVector<Data64>> half_mod_;