Al/natural decomposition order

Makefile

@@ -21,7 +21,7 @@ BENCH_OP_FLAVOR?=DEFAULT
 NODE_VERSION=22.6
 FORWARD_COMPAT?=OFF
 BACKWARD_COMPAT_DATA_URL=https://github.com/zama-ai/tfhe-backward-compat-data.git
-BACKWARD_COMPAT_DATA_BRANCH?=v0.3
+BACKWARD_COMPAT_DATA_BRANCH?=v0.4
 BACKWARD_COMPAT_DATA_PROJECT=tfhe-backward-compat-data
 BACKWARD_COMPAT_DATA_DIR=$(BACKWARD_COMPAT_DATA_PROJECT)
 TFHE_SPEC:=tfhe

backends/tfhe-cuda-backend/cuda/src/crypto/keyswitch.cuh

@@ -75,8 +75,7 @@ keyswitch(Torus *lwe_array_out, const Torus *__restrict__ lwe_output_indexes,
                                             level_count);
       Torus state = a_i >> (sizeof(Torus) * 8 - base_log * level_count);
 
-      for (int j = level_count - 1; j >= 0; j--) {
-        // Levels are stored in reverse order
+      for (int j = 0; j < level_count; j++) {
         auto ksk_block =
             get_ith_block(ksk, i, j, lwe_dimension_out, level_count);
         Torus decomposed = decompose_one<Torus>(state, mask_mod_b, base_log);
@@ -209,8 +208,7 @@ __device__ void packing_keyswitch_lwe_ciphertext_into_glwe_ciphertext(
 
       // block of key for current lwe coefficient (cur_input_lwe[i])
       auto ksk_block = &fp_ksk[i * ksk_block_size];
-      for (int j = level_count - 1; j >= 0; j--) {
-        // Levels are stored in reverse order
+      for (int j = 0; j < level_count; j++) {
         auto ksk_glwe = &ksk_block[j * glwe_size * polynomial_size];
         // Iterate through each level and multiply by the ksk piece
         auto ksk_glwe_chunk = &ksk_glwe[poly_id * coef_per_block];

backends/tfhe-cuda-backend/cuda/src/pbs/bootstrapping_key.cuh

@@ -24,8 +24,8 @@ __device__ const T *get_ith_mask_kth_block(const T *ptr, int i, int k,
                                            uint32_t level_count) {
   return &ptr[get_start_ith_ggsw(i, polynomial_size, glwe_dimension,
                                  level_count) +
-              level * polynomial_size / 2 * (glwe_dimension + 1) *
-                  (glwe_dimension + 1) +
+              (level_count - level - 1) * polynomial_size / 2 *
+                  (glwe_dimension + 1) * (glwe_dimension + 1) +
               k * polynomial_size / 2 * (glwe_dimension + 1)];
 }
 
@@ -35,8 +35,8 @@ __device__ T *get_ith_mask_kth_block(T *ptr, int i, int k, int level,
                                      int glwe_dimension, uint32_t level_count) {
   return &ptr[get_start_ith_ggsw(i, polynomial_size, glwe_dimension,
                                  level_count) +
-              level * polynomial_size / 2 * (glwe_dimension + 1) *
-                  (glwe_dimension + 1) +
+              (level_count - level - 1) * polynomial_size / 2 *
+                  (glwe_dimension + 1) * (glwe_dimension + 1) +
               k * polynomial_size / 2 * (glwe_dimension + 1)];
 }
 template <typename T>
@@ -45,8 +45,8 @@ __device__ T *get_ith_body_kth_block(T *ptr, int i, int k, int level,
                                      int glwe_dimension, uint32_t level_count) {
   return &ptr[get_start_ith_ggsw(i, polynomial_size, glwe_dimension,
                                  level_count) +
-              level * polynomial_size / 2 * (glwe_dimension + 1) *
-                  (glwe_dimension + 1) +
+              (level_count - level - 1) * polynomial_size / 2 *
+                  (glwe_dimension + 1) * (glwe_dimension + 1) +
               k * polynomial_size / 2 * (glwe_dimension + 1) +
               glwe_dimension * polynomial_size / 2];
 }

backends/tfhe-cuda-backend/cuda/src/pbs/programmable_bootstrap_amortized.cuh

@@ -137,7 +137,7 @@ __global__ void device_programmable_bootstrap_amortized(
     // Now that the rotation is done, decompose the resulting polynomial
     // coefficients so as to multiply each decomposed level with the
     // corresponding part of the bootstrapping key
-    for (int level = level_count - 1; level >= 0; level--) {
+    for (int level = 0; level < level_count; level++) {
       for (int i = 0; i < (glwe_dimension + 1); i++) {
         gadget.decompose_and_compress_next_polynomial(accumulator_fft, i);
 

tfhe/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "tfhe"
-version = "0.9.0"
+version = "0.10.0"
 edition = "2021"
 readme = "../README.md"
 keywords = ["fully", "homomorphic", "encryption", "fhe", "cryptography"]
@@ -46,7 +46,7 @@ hex = "0.4.3"
 # End regex-engine deps
 # Used for backward compatibility test metadata
 ron = "0.8"
-tfhe-backward-compat-data = { git = "https://github.com/zama-ai/tfhe-backward-compat-data.git", branch = "v0.3", default-features = false, features = [
+tfhe-backward-compat-data = { git = "https://github.com/zama-ai/tfhe-backward-compat-data.git", branch = "v0.4", default-features = false, features = [
     "load",
 ] }
 

tfhe/docs/fundamentals/serialization.md

@@ -158,7 +158,7 @@ In the following example, we use [bincode](https://crates.io/crates/bincode) for
 
 [dependencies]
 # ...
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 bincode = "1.3.3"
 ```
 

tfhe/docs/getting_started/benchmarks/summary.md

@@ -2,7 +2,7 @@
 
 This document summarizes the timings of some homomorphic operations over 64-bit encrypted integers, depending on the hardware. More details are given for [the CPU](cpu\_benchmarks.md), [the GPU](gpu\_benchmarks.md), or [zeros-knowledge proofs](zk\_proof\_benchmarks.md).
 
-You can get the parameters used for benchmarks by cloning the repository and checking out the commit you want to use (starting with the v0.9.0 release) and run the following make command:
+You can get the parameters used for benchmarks by cloning the repository and checking out the commit you want to use (starting with the v0.10.0 release) and run the following make command:
 
 ```console
 make print_doc_bench_parameters

tfhe/docs/getting_started/installation.md

@@ -9,13 +9,13 @@ First, add **TFHE-rs** as a dependency in your `Cargo.toml`.
 **For `x86_64` machine running a Unix-like OS:**
 
 ```toml
-tfhe = { version = "0.9.0", features = ["boolean", "shortint", "integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["boolean", "shortint", "integer", "x86_64-unix"] }
 ```
 
 **For `ARM` machine running a Unix-like OS:**
 
 ```toml
-tfhe = { version = "0.9.0", features = ["boolean", "shortint", "integer", "aarch64-unix"] }
+tfhe = { version = "0.10.0", features = ["boolean", "shortint", "integer", "aarch64-unix"] }
 ```
 
 **For `x86_64` machines with the** [**`rdseed instruction`**](https://en.wikipedia.org/wiki/RDRAND) **running Windows:**

tfhe/docs/getting_started/quick_start.md

@@ -59,7 +59,7 @@ edition = "2021"
 For x86 Unix systems, add the following configuration to include **TFHE-rs**:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 ```
 
 Your updated `Cargo.toml` file should look like this:
@@ -71,7 +71,7 @@ version = "0.1.0"
 edition = "2021"
 
 [dependencies]
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 ```
 
 If you are on a different platform please refer to the [installation documentation](installation.md) for configuration options of other supported platforms.

tfhe/docs/guides/array.md

@@ -19,7 +19,7 @@ The following example shows a complete workflow of working with encrypted arrays
 # Cargo.toml
 
 [dependencies]
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 ```
 
 ```rust

tfhe/docs/guides/data_versioning.md

@@ -16,7 +16,7 @@ You can load serialized data with the `unversionize` function, even in newer ver
 
 [dependencies]
 # ...
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 tfhe-versionable = "0.2.0"
 bincode = "1.3.3"
 ```

tfhe/docs/guides/run_on_gpu.md

@@ -20,13 +20,13 @@ To use the **TFHE-rs** GPU backend in your project, add the following dependency
 If you are using an `x86` machine:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["boolean", "shortint", "integer", "x86_64-unix", "gpu"] }
+tfhe = { version = "0.10.0", features = ["boolean", "shortint", "integer", "x86_64-unix", "gpu"] }
 ```
 
 If you are using an `ARM` machine:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["boolean", "shortint", "integer", "aarch64-unix", "gpu"] }
+tfhe = { version = "0.10.0", features = ["boolean", "shortint", "integer", "aarch64-unix", "gpu"] }
 ```
 
 {% hint style="success" %}

tfhe/docs/references/core-crypto-api/tutorial.md

@@ -9,7 +9,7 @@ Welcome to this tutorial about `TFHE-rs` `core_crypto` module.
 To use `TFHE-rs`, it first has to be added as a dependency in the `Cargo.toml`:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["x86_64-unix"] }
 ```
 
 This enables the `x86_64-unix` feature to have efficient implementations of various algorithms for `x86_64` CPUs on a Unix-like system. The 'unix' suffix indicates that the `UnixSeeder`, which uses `/dev/random` to generate random numbers, is activated as a fallback if no hardware number generator is available (like `rdseed` on `x86_64` or if the [`Randomization Services`](https://developer.apple.com/documentation/security/1399291-secrandomcopybytes?language=objc) on Apple platforms are not available). To avoid having the `UnixSeeder` as a potential fallback or to run on non-Unix systems (e.g., Windows), the `x86_64` feature is sufficient.
@@ -19,19 +19,19 @@ For Apple Silicon, the `aarch64-unix` or `aarch64` feature should be enabled. `a
 In short: For `x86_64`-based machines running Unix-like OSes:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["x86_64-unix"] }
 ```
 
 For Apple Silicon or aarch64-based machines running Unix-like OSes:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["aarch64-unix"] }
+tfhe = { version = "0.10.0", features = ["aarch64-unix"] }
 ```
 
 For `x86_64`-based machines with the [`rdseed instruction`](https://en.wikipedia.org/wiki/RDRAND) running Windows:
 
 ```toml
-tfhe = { version = "0.9.0", features = ["x86_64"] }
+tfhe = { version = "0.10.0", features = ["x86_64"] }
 ```
 
 ### Commented code to double a 2-bit message in a leveled fashion and using a PBS with the `core_crypto` module.

tfhe/docs/tutorials/ascii_fhe_string.md

@@ -25,7 +25,7 @@ To use the `FheUint8` type, enable the `integer` feature:
 
 [dependencies]
 # Default configuration for x86 Unix machines:
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 ```
 
 Refer to the [installation guide](../getting\_started/installation.md) for other configurations.

tfhe/docs/tutorials/parity_bit.md

@@ -18,7 +18,7 @@ This function returns a Boolean (`true` or `false`) so that the total count of `
 # Cargo.toml
 
 # Default configuration for x86 Unix machines:
-tfhe = { version = "0.9.0", features = ["integer", "x86_64-unix"] }
+tfhe = { version = "0.10.0", features = ["integer", "x86_64-unix"] }
 ```
 
 Refer to the [installation](../getting\_started/installation.md) for other configurations.

tfhe/src/core_crypto/algorithms/ggsw_encryption.rs

@@ -136,12 +136,13 @@ pub fn encrypt_constant_ggsw_ciphertext<Scalar, NoiseDistribution, KeyCont, Outp
         .expect("Failed to split generator into ggsw levels");
 
     let decomp_base_log = output.decomposition_base_log();
+    let decomp_level_count = output.decomposition_level_count();
     let ciphertext_modulus = output.ciphertext_modulus();
 
-    for (level_index, (mut level_matrix, mut generator)) in
+    for (output_index, (mut level_matrix, mut generator)) in
         output.iter_mut().zip(gen_iter).enumerate()
     {
-        let decomp_level = DecompositionLevel(level_index + 1);
+        let decomp_level = DecompositionLevel(decomp_level_count.0 - output_index);
         let factor = ggsw_encryption_multiplicative_factor(
             ciphertext_modulus,
             decomp_level,
@@ -269,11 +270,12 @@ pub fn par_encrypt_constant_ggsw_ciphertext<Scalar, NoiseDistribution, KeyCont,
         .expect("Failed to split generator into ggsw levels");
 
     let decomp_base_log = output.decomposition_base_log();
+    let decomp_level_count = output.decomposition_level_count();
     let ciphertext_modulus = output.ciphertext_modulus();
 
     output.par_iter_mut().zip(gen_iter).enumerate().for_each(
-        |(level_index, (mut level_matrix, mut generator))| {
-            let decomp_level = DecompositionLevel(level_index + 1);
+        |(output_index, (mut level_matrix, mut generator))| {
+            let decomp_level = DecompositionLevel(decomp_level_count.0 - output_index);
             let factor = ggsw_encryption_multiplicative_factor(
                 ciphertext_modulus,
                 decomp_level,
@@ -401,12 +403,13 @@ pub fn encrypt_constant_seeded_ggsw_ciphertext_with_existing_generator<
         .expect("Failed to split generator into ggsw levels");
 
     let decomp_base_log = output.decomposition_base_log();
+    let decomp_level_count = output.decomposition_level_count();
     let ciphertext_modulus = output.ciphertext_modulus();
 
-    for (level_index, (mut level_matrix, mut loop_generator)) in
+    for (output_index, (mut level_matrix, mut loop_generator)) in
         output.iter_mut().zip(gen_iter).enumerate()
     {
-        let decomp_level = DecompositionLevel(level_index + 1);
+        let decomp_level = DecompositionLevel(decomp_level_count.0 - output_index);
         let factor = ggsw_encryption_multiplicative_factor(
             ciphertext_modulus,
             decomp_level,
@@ -581,11 +584,12 @@ pub fn par_encrypt_constant_seeded_ggsw_ciphertext_with_existing_generator<
         .expect("Failed to split generator into ggsw levels");
 
     let decomp_base_log = output.decomposition_base_log();
+    let decomp_level_count = output.decomposition_level_count();
     let ciphertext_modulus = output.ciphertext_modulus();
 
     output.par_iter_mut().zip(gen_iter).enumerate().for_each(
-        |(level_index, (mut level_matrix, mut generator))| {
-            let decomp_level = DecompositionLevel(level_index + 1);
+        |(output_index, (mut level_matrix, mut generator))| {
+            let decomp_level = DecompositionLevel(decomp_level_count.0 - output_index);
             let factor = ggsw_encryption_multiplicative_factor(
                 ciphertext_modulus,
                 decomp_level,
@@ -881,7 +885,7 @@ where
         glwe_secret_key.glwe_dimension()
     );
 
-    let level_matrix = ggsw_ciphertext.last().unwrap();
+    let level_matrix = ggsw_ciphertext.first().unwrap();
     let level_matrix_as_glwe_list = level_matrix.as_glwe_list();
     let last_row = level_matrix_as_glwe_list.last().unwrap();
     let decomp_level = ggsw_ciphertext.decomposition_level_count();

tfhe/src/core_crypto/algorithms/lwe_keyswitch.rs

@@ -216,8 +216,7 @@ pub fn keyswitch_lwe_ciphertext_native_mod_compatible<Scalar, KSKCont, InputCont
     {
         let decomposition_iter = decomposer.decompose(input_mask_element);
         // Loop over the levels
-        for (level_key_ciphertext, decomposed) in
-            keyswitch_key_block.iter().rev().zip(decomposition_iter)
+        for (level_key_ciphertext, decomposed) in keyswitch_key_block.iter().zip(decomposition_iter)
         {
             slice_wrapping_sub_scalar_mul_assign(
                 output_lwe_ciphertext.as_mut(),
@@ -305,8 +304,7 @@ pub fn keyswitch_lwe_ciphertext_other_mod<Scalar, KSKCont, InputCont, OutputCont
     {
         let decomposition_iter = decomposer.decompose(input_mask_element);
         // Loop over the levels
-        for (level_key_ciphertext, decomposed) in
-            keyswitch_key_block.iter().rev().zip(decomposition_iter)
+        for (level_key_ciphertext, decomposed) in keyswitch_key_block.iter().zip(decomposition_iter)
         {
             slice_wrapping_sub_scalar_mul_assign_custom_modulus(
                 output_lwe_ciphertext.as_mut(),
@@ -438,8 +436,7 @@ pub fn keyswitch_lwe_ciphertext_with_scalar_change<
     {
         let decomposition_iter = input_decomposer.decompose(input_mask_element);
         // Loop over the levels
-        for (level_key_ciphertext, decomposed) in
-            keyswitch_key_block.iter().rev().zip(decomposition_iter)
+        for (level_key_ciphertext, decomposed) in keyswitch_key_block.iter().zip(decomposition_iter)
         {
             slice_wrapping_sub_scalar_mul_assign(
                 output_lwe_ciphertext.as_mut(),
@@ -802,7 +799,7 @@ pub fn par_keyswitch_lwe_ciphertext_with_thread_count_native_mod_compatible<
                 let decomposition_iter = decomposer.decompose(input_mask_element);
                 // Loop over the levels
                 for (level_key_ciphertext, decomposed) in
-                    keyswitch_key_block.iter().rev().zip(decomposition_iter)
+                    keyswitch_key_block.iter().zip(decomposition_iter)
                 {
                     slice_wrapping_sub_scalar_mul_assign(
                         buffer.as_mut(),
@@ -949,7 +946,7 @@ pub fn par_keyswitch_lwe_ciphertext_with_thread_count_other_mod<
                 let decomposition_iter = decomposer.decompose(input_mask_element);
                 // Loop over the levels
                 for (level_key_ciphertext, decomposed) in
-                    keyswitch_key_block.iter().rev().zip(decomposition_iter)
+                    keyswitch_key_block.iter().zip(decomposition_iter)
                 {
                     slice_wrapping_sub_scalar_mul_assign_custom_modulus(
                         buffer.as_mut(),

tfhe/src/core_crypto/algorithms/lwe_keyswitch_key_generation.rs

@@ -160,6 +160,7 @@ pub fn generate_lwe_keyswitch_key_native_mod_compatible<
         // We fill the buffer with the powers of the key elements
         for (level, message) in (1..=decomp_level_count.0)
             .map(DecompositionLevel)
+            .rev()
             .zip(decomposition_plaintexts_buffer.iter_mut())
         {
             // Here  we take the decomposition term from the native torus, bring it to the torus we
@@ -234,6 +235,7 @@ pub fn generate_lwe_keyswitch_key_other_mod<
         // We fill the buffer with the powers of the key elements
         for (level, message) in (1..=decomp_level_count.0)
             .map(DecompositionLevel)
+            .rev()
             .zip(decomposition_plaintexts_buffer.iter_mut())
         {
             // Here  we take the decomposition term from the native torus, bring it to the torus we
@@ -415,6 +417,7 @@ pub fn generate_seeded_lwe_keyswitch_key<
         // We fill the buffer with the powers of the key elmements
         for (level, message) in (1..=decomp_level_count.0)
             .map(DecompositionLevel)
+            .rev()
             .zip(decomposition_plaintexts_buffer.iter_mut())
         {
             // Here  we take the decomposition term from the native torus, bring it to the torus we

tfhe/src/core_crypto/algorithms/lwe_packing_keyswitch.rs

@@ -177,7 +177,7 @@ pub fn keyswitch_lwe_ciphertext_into_glwe_ciphertext<Scalar, KeyCont, InputCont,
         // Loop over the number of levels:
         // We compute the multiplication of a ciphertext from the private functional
         // keyswitching key with a piece of the decomposition and subtract it to the buffer
-        for (level_key_cipher, decomposed) in keyswitch_key_block.iter().rev().zip(decomp) {
+        for (level_key_cipher, decomposed) in keyswitch_key_block.iter().zip(decomp) {
             slice_wrapping_sub_scalar_mul_assign(
                 output_glwe_ciphertext.as_mut(),
                 level_key_cipher.as_ref(),

tfhe/src/core_crypto/algorithms/lwe_packing_keyswitch_key_generation.rs

@@ -135,6 +135,7 @@ pub fn generate_lwe_packing_keyswitch_key<
         // We fill the buffer with the powers of the key elements
         for (level, mut messages) in (1..=decomp_level_count.0)
             .map(DecompositionLevel)
+            .rev()
             .zip(decomposition_plaintexts_buffer.chunks_exact_mut(polynomial_size.0))
         {
             // Here  we take the decomposition term from the native torus, bring it to the torus we
@@ -330,6 +331,7 @@ pub fn generate_seeded_lwe_packing_keyswitch_key<
         // We fill the buffer with the powers of the key elements
         for (level, mut messages) in (1..=decomp_level_count.0)
             .map(DecompositionLevel)
+            .rev()
             .zip(decomposition_plaintexts_buffer.chunks_exact_mut(polynomial_size.0))
         {
             // Here  we take the decomposition term from the native torus, bring it to the torus we