chore(zk): add a shortint wrapper for the CompactPkeCrs

This allows to detect unused dispatch enums

.github/workflows/aws_tfhe_fast_tests.yml

@@ -54,10 +54,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |
@@ -198,7 +199,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/restore@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         with:
           path: |
             ~/.nvm
@@ -211,7 +212,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/save@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/aws_tfhe_integer_tests.yml

@@ -42,11 +42,12 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
           persist-credentials: "false"
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |
@@ -58,6 +59,7 @@ jobs:
               - tfhe/src/core_crypto/**
               - tfhe/src/shortint/**
               - tfhe/src/integer/**
+              - .github/workflows/aws_tfhe_integer_tests.yml
 
   setup-instance:
     name: Setup instance (unsigned-integer-tests)

.github/workflows/aws_tfhe_signed_integer_tests.yml

@@ -42,11 +42,12 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
           persist-credentials: "false"
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |
@@ -58,6 +59,7 @@ jobs:
               - tfhe/src/core_crypto/**
               - tfhe/src/shortint/**
               - tfhe/src/integer/**
+              - .github/workflows/aws_tfhe_signed_integer_tests.yml
 
   setup-instance:
     name: Setup instance (unsigned-integer-tests)

.github/workflows/aws_tfhe_tests.yml

@@ -63,10 +63,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/aws_tfhe_wasm_tests.yml

@@ -61,7 +61,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/restore@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         with:
           path: |
             ~/.nvm
@@ -74,7 +74,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/save@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/benchmark_gpu_4090.yml

@@ -116,6 +116,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Get benchmark details
         run: |

.github/workflows/benchmark_gpu_integer_common.yml

@@ -236,7 +236,8 @@ jobs:
           --commit-date "${{ env.COMMIT_DATE }}" \
           --bench-date "${{ env.BENCH_DATE }}" \
           --walk-subdirs \
-          --name-suffix avx512
+          --name-suffix avx512 \
+          --bench-type ${{ matrix.bench_type }}
 
       - name: Upload parsed results artifact
         uses: actions/upload-artifact@b4b15b8c7c6ac21ea08fcf65892d2ee8f75cf882

.github/workflows/benchmark_integer.yml

@@ -35,7 +35,6 @@ env:
   SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
   SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
   FAST_BENCH: TRUE
-  BENCH_TYPE: latency
 
 jobs:
   prepare-matrix:
@@ -168,7 +167,7 @@ jobs:
           --bench-date "${{ env.BENCH_DATE }}" \
           --walk-subdirs \
           --name-suffix avx512 \
-          --bench-type ${{ env.BENCH_TYPE }}
+          --bench-type ${{ matrix.bench_type }}
 
       - name: Upload parsed results artifact
         uses: actions/upload-artifact@b4b15b8c7c6ac21ea08fcf65892d2ee8f75cf882

.github/workflows/benchmark_tfhe_fft.yml

@@ -16,6 +16,9 @@ on:
   push:
     branches:
       - "main"
+    paths:
+      - tfhe-fft/**
+      - .github/workflows/benchmark_tfhe_fft.yml
   schedule:
     # Job will be triggered each Thursday at 11p.m.
     - cron: '0 23 * * 4'
@@ -47,7 +50,7 @@ jobs:
     runs-on: ${{ needs.setup-ec2.outputs.runner-name }}
     steps:
       - name: Checkout tfhe-rs repo with tags
-        uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
 
@@ -87,7 +90,7 @@ jobs:
           path: ${{ env.RESULTS_FILENAME }}
 
       - name: Checkout Slab repo
-        uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           repository: zama-ai/slab
           path: slab

.github/workflows/benchmark_tfhe_ntt.yml

@@ -16,6 +16,9 @@ on:
   push:
     branches:
       - "main"
+    paths:
+      - tfhe-ntt/**
+      - .github/workflows/benchmark_tfhe_ntt.yml
   schedule:
     # Job will be triggered each Friday at 11p.m.
     - cron: "0 23 * * 5"
@@ -47,7 +50,7 @@ jobs:
     runs-on: ${{ needs.setup-ec2.outputs.runner-name }}
     steps:
       - name: Checkout tfhe-rs repo with tags
-        uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
 
@@ -87,7 +90,7 @@ jobs:
           path: ${{ env.RESULTS_FILENAME }}
 
       - name: Checkout Slab repo
-        uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           repository: zama-ai/slab
           path: slab

.github/workflows/benchmark_tfhe_zk_pok.yml

@@ -36,7 +36,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@c3a1bb2c992d77180ae65be6ae6c166cf40f857c
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/benchmark_wasm_client.yml

@@ -36,10 +36,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |
@@ -108,7 +109,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/restore@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         with:
           path: |
             ~/.nvm
@@ -121,7 +122,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@6849a6489940f00c2f30c0fb92c6274307ccb58a #v4.1.2
+        uses: actions/cache/save@1bd1e32a3bdc45362d1e726936510720a7c30a57 #v4.2.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/benchmark_zk_pke.yml

@@ -4,10 +4,14 @@ name: PKE ZK benchmarks
 on:
   workflow_dispatch:
     inputs:
-      run_throughput:
-        description: "Run throughput benchmarks"
-        type: boolean
-        default: false
+      bench_type:
+        description: "Benchmarks type"
+        type: choice
+        default: latency
+        options:
+          - latency
+          - throughput
+          - both
 
   push:
     branches:
@@ -26,7 +30,6 @@ env:
   SLACK_ICON: https://pbs.twimg.com/profile_images/1274014582265298945/OjBKP9kn_400x400.png
   SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
   SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
-  BENCH_TYPE: latency
 
 jobs:
   should-run:
@@ -40,10 +43,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |
@@ -59,10 +63,37 @@ jobs:
               - tfhe/benches/integer/zk_pke.rs
               - .github/workflows/zk_pke_benchmark.yml
 
+  prepare-matrix:
+    name: Prepare operations matrix
+    runs-on: ubuntu-latest
+    if: github.event_name != 'schedule' ||
+      (github.event_name == 'schedule' && github.repository == 'zama-ai/tfhe-rs')
+    outputs:
+      bench_type: ${{ steps.set_bench_type.outputs.bench_type }}
+    steps:
+      - name: Set benchmark types
+        if: github.event_name == 'workflow_dispatch'
+        run: |
+          if [[ "${{ inputs.bench_type }}" == "both" ]]; then
+            echo "BENCH_TYPE=[\"latency\", \"throughput\"]" >> "${GITHUB_ENV}"
+          else
+            echo "BENCH_TYPE=[\"${{ inputs.bench_type }}\"]" >> "${GITHUB_ENV}"
+          fi
+
+      - name: Default benchmark type
+        if: github.event_name != 'workflow_dispatch'
+        run: |
+          echo "BENCH_TYPE=[\"latency\"]" >> "${GITHUB_ENV}"
+
+      - name: Set benchmark types output
+        id: set_bench_type
+        run: |
+          echo "bench_type=${{ toJSON(env.BENCH_TYPE) }}" >> "${GITHUB_OUTPUT}"
+
   setup-instance:
     name: Setup instance (pke-zk-benchmarks)
     runs-on: ubuntu-latest
-    needs: should-run
+    needs: [ should-run, prepare-matrix ]
     if: github.event_name == 'workflow_dispatch' ||
       (github.event_name == 'schedule' && github.repository == 'zama-ai/tfhe-rs') ||
       (github.event_name == 'push' &&
@@ -85,11 +116,15 @@ jobs:
   pke-zk-benchmarks:
     name: Execute PKE ZK benchmarks
     if: needs.setup-instance.result != 'skipped'
-    needs: setup-instance
+    needs: [ prepare-matrix, setup-instance ]
     concurrency:
       group: ${{ github.workflow }}_${{github.event_name}}_${{ github.ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
       cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
     runs-on: ${{ needs.setup-instance.outputs.runner-name }}
+    strategy:
+      max-parallel: 1
+      matrix:
+        bench_type: ${{ fromJSON(needs.prepare-matrix.outputs.bench_type) }}
     steps:
       - name: Checkout tfhe-rs repo with tags
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
@@ -117,14 +152,9 @@ jobs:
           path: slab
           token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
-      - name: Should run throughput benchmarks
-        if: inputs.run_throughput
-        run: |
-          echo "BENCH_TYPE=throughput" >> "${GITHUB_ENV}"
-
       - name: Run benchmarks with AVX512
         run: |
-          make bench_integer_zk
+          make BENCH_TYPE=${{ matrix.bench_type }} bench_integer_zk
 
       - name: Parse results
         run: |
@@ -138,7 +168,7 @@ jobs:
           --bench-date "${{ env.BENCH_DATE }}" \
           --walk-subdirs \
           --name-suffix avx512 \
-          --bench-type ${{ env.BENCH_TYPE }}
+          --bench-type ${{ matrix.bench_type }}
 
       - name: Parse CRS sizes results
         run: |

.github/workflows/cargo_build_tfhe_fft.yml

@@ -21,7 +21,7 @@ jobs:
       fail-fast: false
 
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af

.github/workflows/cargo_build_tfhe_ntt.yml

@@ -19,7 +19,7 @@ jobs:
         os: [ubuntu-latest, macos-latest, windows-latest]
       fail-fast: false
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af

.github/workflows/cargo_test_fft.yml

@@ -19,7 +19,7 @@ jobs:
         runner_type: [ubuntu-latest, macos-latest, windows-latest]
       fail-fast: false
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af
@@ -44,7 +44,7 @@ jobs:
       matrix:
         runner_type: [ubuntu-latest, macos-latest, windows-latest]
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af
@@ -63,7 +63,7 @@ jobs:
   cargo-tests-node-js:
     runs-on: "ubuntu-latest"
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Test node js
         run: |

.github/workflows/cargo_test_ntt.yml

@@ -19,7 +19,7 @@ jobs:
         os: [ubuntu-latest, macos-latest, windows-latest]
       fail-fast: false
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af
@@ -39,7 +39,7 @@ jobs:
       matrix:
         os: [ubuntu-latest, macos-latest, windows-latest]
     steps:
-      - uses: actions/checkout@ac593985615ec2ede58e132d2e21d2b1cbd6127c
+      - uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
 
       - name: Install Rust
         uses: actions-rs/toolchain@16499b5e05bf2e26879000db0c1d13f7e13fa3af

.github/workflows/code_coverage.yml

@@ -53,7 +53,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           files_yaml: |
             tfhe:
@@ -83,7 +83,7 @@ jobs:
           make test_shortint_cov
 
       - name: Upload tfhe coverage to Codecov
-        uses: codecov/codecov-action@5c47607acb93fed5485fdbf7232e8a31425f672a
+        uses: codecov/codecov-action@7f8b4b4bde536c465e797be725718b88c5d95e0e
         if: steps.changed-files.outputs.tfhe_any_changed == 'true'
         with:
           token: ${{ secrets.CODECOV_TOKEN }}
@@ -97,7 +97,7 @@ jobs:
           make test_integer_cov
 
       - name: Upload tfhe coverage to Codecov
-        uses: codecov/codecov-action@5c47607acb93fed5485fdbf7232e8a31425f672a
+        uses: codecov/codecov-action@7f8b4b4bde536c465e797be725718b88c5d95e0e
         if: steps.changed-files.outputs.tfhe_any_changed == 'true'
         with:
           token: ${{ secrets.CODECOV_TOKEN }}

.github/workflows/gpu_fast_h100_tests.yml

@@ -31,10 +31,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_fast_tests.yml

@@ -30,10 +30,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_full_multi_gpu_tests.yml

@@ -31,10 +31,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_integer_long_run_tests.yml

@@ -29,14 +29,14 @@ jobs:
     steps:
       - name: Start instance
         id: start-instance
-        uses: zama-ai/slab-github-runner@801df0b8db5ea2b06128b7476c652f5ed5f193a8
+        uses: zama-ai/slab-github-runner@98f0788261a7323d5d695a883e20df36591a92b7
         with:
           mode: start
           github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
           slab-url: ${{ secrets.SLAB_BASE_URL }}
           job-secret: ${{ secrets.JOB_SECRET }}
           backend: hyperstack
-          profile: single-h100
+          profile: 2-h100
 
   cuda-tests:
     name: Long run GPU H100 tests
@@ -77,7 +77,7 @@ jobs:
           echo "HOME=/home/ubuntu" >> "${GITHUB_ENV}"
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@7b1c307e0dcbda6122208f10795a713336a9b35a
+        uses: dtolnay/rust-toolchain@315e265cd78dad1e1dcf3a5074f6d6c47029d5aa
         with:
           toolchain: stable
 
@@ -129,7 +129,7 @@ jobs:
     steps:
       - name: Stop instance
         id: stop-instance
-        uses: zama-ai/slab-github-runner@801df0b8db5ea2b06128b7476c652f5ed5f193a8
+        uses: zama-ai/slab-github-runner@98f0788261a7323d5d695a883e20df36591a92b7
         with:
           mode: stop
           github-token: ${{ secrets.SLAB_ACTION_TOKEN }}

.github/workflows/gpu_signed_integer_classic_tests.yml

@@ -35,7 +35,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_signed_integer_h100_tests.yml

@@ -35,7 +35,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_signed_integer_tests.yml

@@ -38,10 +38,11 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_unsigned_integer_classic_tests.yml

@@ -35,7 +35,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_unsigned_integer_h100_tests.yml

@@ -35,7 +35,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/gpu_unsigned_integer_tests.yml

@@ -41,7 +41,7 @@ jobs:
 
       - name: Check for file changes
         id: changed-files
-        uses: tj-actions/changed-files@4edd678ac3f81e2dc578756871e4d00c19191daf
+        uses: tj-actions/changed-files@bab30c2299617f6615ec02a68b9a40d10bd21366
         with:
           since_last_remote_commit: true
           files_yaml: |

.github/workflows/integer_long_run_tests.yml

@@ -29,7 +29,7 @@ jobs:
     steps:
       - name: Start instance
         id: start-instance
-        uses: zama-ai/slab-github-runner@801df0b8db5ea2b06128b7476c652f5ed5f193a8
+        uses: zama-ai/slab-github-runner@98f0788261a7323d5d695a883e20df36591a92b7
         with:
           mode: start
           github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
@@ -53,7 +53,7 @@ jobs:
           token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@7b1c307e0dcbda6122208f10795a713336a9b35a
+        uses: dtolnay/rust-toolchain@315e265cd78dad1e1dcf3a5074f6d6c47029d5aa
         with:
           toolchain: stable
 
@@ -77,7 +77,7 @@ jobs:
     steps:
       - name: Stop instance
         id: stop-instance
-        uses: zama-ai/slab-github-runner@801df0b8db5ea2b06128b7476c652f5ed5f193a8
+        uses: zama-ai/slab-github-runner@98f0788261a7323d5d695a883e20df36591a92b7
         with:
           mode: stop
           github-token: ${{ secrets.SLAB_ACTION_TOKEN }}

.github/workflows/make_release.yml

@@ -46,6 +46,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
       - name: Prepare package
         run: |
           cargo package -p tfhe
@@ -84,6 +85,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
       - name: Create NPM version tag
         if: ${{ inputs.npm_latest_tag }}
         run: |

.github/workflows/make_release_concrete_csprng.yml

@@ -27,6 +27,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Publish crate.io package
         env:

.github/workflows/make_release_cuda.yml

@@ -64,6 +64,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Set up home
         run: |

.github/workflows/make_release_tfhe_fft.yml

@@ -25,7 +25,7 @@ jobs:
     needs: verify_tag
     steps:
       - name: Checkout
-        uses: actions/checkout@b4ffde65f46336ab88eb53be808477a3936bae11
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
 

.github/workflows/make_release_tfhe_ntt.yml

@@ -25,7 +25,7 @@ jobs:
     needs: verify_tag
     steps:
       - name: Checkout
-        uses: actions/checkout@b4ffde65f46336ab88eb53be808477a3936bae11
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
 

.github/workflows/make_release_tfhe_versionable.yml

@@ -27,6 +27,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Publish proc-macro crate
         env:

.github/workflows/make_release_zk_pok.yml

@@ -28,6 +28,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
 
       - name: Publish crate.io package
         env:

.github/workflows/sync_on_push.yml

@@ -16,6 +16,7 @@ jobs:
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
         with:
           fetch-depth: 0
+          token: ${{ secrets.FHE_ACTIONS_TOKEN }}
       - name: git-sync
         uses: wei/git-sync@55c6b63b4f21607da0e9877ca9b4d11a29fc6d83
         with:

.gitignore

@@ -12,8 +12,8 @@ target/
 **/*.bin
 
 # Some of our bench outputs
-/tfhe/benchmarks_parameters
-/tfhe-zk-pok/benchmarks_parameters
+/crates/tfhe/benchmarks_parameters
+/crates/tfhe-zk-pok/benchmarks_parameters
 **/*.csv
 
 # dieharder run log
@@ -26,11 +26,11 @@ dieharder_run.log
 backends/tfhe-cuda-backend/cuda/cmake-build-debug/
 
 # WASM tests
-tfhe/web_wasm_parallel_tests/server.PID
+crates/tfhe/tests/web_wasm_parallel/server.PID
 venv/
 web-test-runner/
 node_modules/
 package-lock.json
 
 # Dir used for backward compatibility test data
-tfhe/tfhe-backward-compat-data/
+crates/tfhe/tfhe-backward-compat-data/

Cargo.toml

@@ -1,23 +1,18 @@
 [workspace]
 resolver = "2"
 members = [
-    "tfhe",
-    "tfhe-fft",
-    "tfhe-ntt",
-    "tfhe-zk-pok",
+    "crates/*",
     "tasks",
     "apps/trivium",
-    "tfhe-csprng",
     "backends/tfhe-cuda-backend",
-    "utils/tfhe-versionable",
-    "utils/tfhe-versionable-derive",
 ]
 
 exclude = [
-    "tfhe/backward_compatibility_tests",
+    "crates/tfhe/backward_compatibility_tests",
     "utils/cargo-tfhe-lints-inner",
     "utils/cargo-tfhe-lints"
 ]
+
 [workspace.dependencies]
 aligned-vec = { version = "0.5", default-features = false }
 bytemuck = "1.14.3"
@@ -27,6 +22,19 @@ pulp = { version = "0.18.22", default-features = false }
 serde = { version = "1.0", default-features = false }
 wasm-bindgen = ">=0.2.86,<0.2.94"
 
+[workspace.package]
+version = "0.11.0"
+license = "BSD-3-Clause-Clear"
+repository = "https://github.com/zama-ai/tfhe-rs"
+documentation = "https://docs.zama.ai/tfhe-rs"
+
+[workspace.lints.rust]
+unexpected_cfgs = { level = "warn", check-cfg = [
+    'cfg(bench)',
+    'cfg(tarpaulin)',
+    'cfg(tfhe_lints)',
+] }
+
 [profile.bench]
 lto = "fat"
 

Makefile

@@ -5,9 +5,10 @@ CARGO_RS_CHECK_TOOLCHAIN:=+$(RS_CHECK_TOOLCHAIN)
 TARGET_ARCH_FEATURE:=$(shell ./scripts/get_arch_feature.sh)
 CPU_COUNT=$(shell ./scripts/cpu_count.sh)
 RS_BUILD_TOOLCHAIN:=stable
+TFHE_SRC:=crates/tfhe
 CARGO_RS_BUILD_TOOLCHAIN:=+$(RS_BUILD_TOOLCHAIN)
 CARGO_PROFILE?=release
-MIN_RUST_VERSION:=$(shell grep '^rust-version[[:space:]]*=' tfhe/Cargo.toml | cut -d '=' -f 2 | xargs)
+MIN_RUST_VERSION:=$(shell grep '^rust-version[[:space:]]*=' $(TFHE_SRC)/Cargo.toml | cut -d '=' -f 2 | xargs)
 AVX512_SUPPORT?=OFF
 WASM_RUSTFLAGS:=
 BIG_TESTS_INSTANCE?=FALSE
@@ -28,7 +29,7 @@ TFHE_SPEC:=tfhe
 # We are kind of hacking the cut here, the version cannot contain a quote '"'
 WASM_BINDGEN_VERSION:=$(shell grep '^wasm-bindgen[[:space:]]*=' Cargo.toml | cut -d '"' -f 2 | xargs)
 WEB_RUNNER_DIR=web-test-runner
-WEB_SERVER_DIR=tfhe/web_wasm_parallel_tests
+WEB_SERVER_DIR=tfhe/tests/web_wasm_parallel
 # This is done to avoid forgetting it, we still precise the RUSTFLAGS in the commands to be able to
 # copy paste the command in the terminal and change them if required without forgetting the flags
 export RUSTFLAGS?=-C target-cpu=native
@@ -243,7 +244,7 @@ fmt_js: check_nvm_installed
 	source ~/.nvm/nvm.sh && \
 	nvm install $(NODE_VERSION) && \
 	nvm use $(NODE_VERSION) && \
-	$(MAKE) -C tfhe/web_wasm_parallel_tests fmt
+	$(MAKE) -C $(TFHE_SRC)/tests/web_wasm_parallel fmt
 
 .PHONY: fmt_gpu # Format rust and cuda code
 fmt_gpu: install_rs_check_toolchain
@@ -252,7 +253,7 @@ fmt_gpu: install_rs_check_toolchain
 
 .PHONY: fmt_c_tests # Format c tests
 fmt_c_tests:
-	find tfhe/c_api_tests/ -regex '.*\.\(cpp\|hpp\|cu\|c\|h\)' -exec clang-format -style=file -i {} \;
+	find $(TFHE_SRC)/tests/c_api/ -regex '.*\.\(cpp\|hpp\|cu\|c\|h\)' -exec clang-format -style=file -i {} \;
 
 .PHONY: check_fmt # Check rust code format
 check_fmt: install_rs_check_toolchain
@@ -260,7 +261,7 @@ check_fmt: install_rs_check_toolchain
 
 .PHONY: check_fmt_c_tests  # Check C tests format
 check_fmt_c_tests:
-	find tfhe/c_api_tests/ -regex '.*\.\(cpp\|hpp\|cu\|c\|h\)' -exec clang-format --dry-run --Werror -style=file {} \;
+	find $(TFHE_SRC)/tests/c_api/ -regex '.*\.\(cpp\|hpp\|cu\|c\|h\)' -exec clang-format --dry-run --Werror -style=file {} \;
 
 .PHONY: check_fmt_gpu # Check rust and cuda code format
 check_fmt_gpu: install_rs_check_toolchain
@@ -272,7 +273,7 @@ check_fmt_js: check_nvm_installed
 	source ~/.nvm/nvm.sh && \
 	nvm install $(NODE_VERSION) && \
 	nvm use $(NODE_VERSION) && \
-	$(MAKE) -C tfhe/web_wasm_parallel_tests check_fmt
+	$(MAKE) -C $(TFHE_SRC)/tests/web_wasm_parallel check_fmt
 
 .PHONY: check_typos # Check for typos in codebase
 check_typos: install_typos_checker
@@ -442,7 +443,7 @@ check_rust_bindings_did_not_change:
 
 .PHONY: tfhe_lints # Run custom tfhe-rs lints
 tfhe_lints: install_tfhe_lints
-	cd tfhe && RUSTFLAGS="$(RUSTFLAGS)" cargo tfhe-lints \
+	cd $(TFHE_SRC) && RUSTFLAGS="$(RUSTFLAGS)" cargo tfhe-lints \
 		--features=$(TARGET_ARCH_FEATURE),boolean,shortint,integer,zk-pok -- -D warnings
 
 .PHONY: build_core # Build core_crypto without experimental features
@@ -508,25 +509,25 @@ build_c_api_experimental_deterministic_fft: install_rs_check_toolchain
 
 .PHONY: build_web_js_api # Build the js API targeting the web browser
 build_web_js_api: install_rs_build_toolchain install_wasm_pack
-	cd tfhe && \
+	cd $(TFHE_SRC) && \
 	RUSTFLAGS="$(WASM_RUSTFLAGS)" rustup run "$(RS_BUILD_TOOLCHAIN)" \
 		wasm-pack build --release --target=web \
 		-- --features=boolean-client-js-wasm-api,shortint-client-js-wasm-api,integer-client-js-wasm-api,zk-pok
 
 .PHONY: build_web_js_api_parallel # Build the js API targeting the web browser with parallelism support
 build_web_js_api_parallel: install_rs_check_toolchain install_wasm_pack
-	cd tfhe && \
+	cd $(TFHE_SRC) && \
 	rustup component add rust-src --toolchain $(RS_CHECK_TOOLCHAIN) && \
 	RUSTFLAGS="$(WASM_RUSTFLAGS) -C target-feature=+atomics,+bulk-memory,+mutable-globals" rustup run $(RS_CHECK_TOOLCHAIN) \
 		wasm-pack build --release --target=web \
 		-- --features=boolean-client-js-wasm-api,shortint-client-js-wasm-api,integer-client-js-wasm-api,parallel-wasm-api,zk-pok \
 		-Z build-std=panic_abort,std && \
 	find pkg/snippets -type f -iname workerHelpers.worker.js -exec sed -i "s|from '..\/..\/..\/';|from '..\/..\/..\/tfhe.js';|" {} \;
-	jq '.files += ["snippets"]' tfhe/pkg/package.json > tmp_pkg.json && mv -f tmp_pkg.json tfhe/pkg/package.json
+	jq '.files += ["snippets"]' $(TFHE_SRC)/pkg/package.json > tmp_pkg.json && mv -f tmp_pkg.json $(TFHE_SRC)/pkg/package.json
 
 .PHONY: build_node_js_api # Build the js API targeting nodejs
 build_node_js_api: install_rs_build_toolchain install_wasm_pack
-	cd tfhe && \
+	cd $(TFHE_SRC) && \
 	RUSTFLAGS="$(WASM_RUSTFLAGS)" rustup run "$(RS_BUILD_TOOLCHAIN)" \
 		wasm-pack build --release --target=nodejs \
 		-- --features=boolean-client-js-wasm-api,shortint-client-js-wasm-api,integer-client-js-wasm-api,zk-pok
@@ -877,7 +878,7 @@ test_zk_wasm_x86_compat_ci: check_nvm_installed
 
 .PHONY: test_zk_wasm_x86_compat # Check compatibility between wasm and x86_64 proofs
 test_zk_wasm_x86_compat: install_rs_build_toolchain build_node_js_api
-	cd tfhe/tests/zk_wasm_x86_test && npm install
+	cd $(TFHE_SRC)/tests/zk_wasm_x86_test && npm install
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
 		-p tfhe --test zk_wasm_x86_test --features=$(TARGET_ARCH_FEATURE),integer,zk-pok
 
@@ -891,11 +892,11 @@ test_versionable: install_rs_build_toolchain
 .PHONY: test_backward_compatibility_ci
 test_backward_compatibility_ci: install_rs_build_toolchain
 	TFHE_BACKWARD_COMPAT_DATA_DIR="$(BACKWARD_COMPAT_DATA_DIR)" RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		--config "patch.'$(BACKWARD_COMPAT_DATA_URL)'.$(BACKWARD_COMPAT_DATA_PROJECT).path=\"tfhe/$(BACKWARD_COMPAT_DATA_DIR)\"" \
+		--config "patch.'$(BACKWARD_COMPAT_DATA_URL)'.$(BACKWARD_COMPAT_DATA_PROJECT).path=\"$(TFHE_SRC)/$(BACKWARD_COMPAT_DATA_DIR)\"" \
 		--features=$(TARGET_ARCH_FEATURE),shortint,integer,zk-pok -p $(TFHE_SPEC) test_backward_compatibility -- --nocapture
 
 .PHONY: test_backward_compatibility # Same as test_backward_compatibility_ci but tries to clone the data repo first if needed
-test_backward_compatibility: tfhe/$(BACKWARD_COMPAT_DATA_DIR) test_backward_compatibility_ci
+test_backward_compatibility: $(TFHE_SRC)/$(BACKWARD_COMPAT_DATA_DIR) test_backward_compatibility_ci
 
 .PHONY: backward_compat_branch # Prints the required backward compatibility branch
 backward_compat_branch:
@@ -907,7 +908,7 @@ doc: install_rs_check_toolchain
 	DOCS_RS=1 \
 	RUSTDOCFLAGS="--html-in-header katex-header.html" \
 	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" doc \
-		--features=$(TARGET_ARCH_FEATURE),boolean,shortint,integer,gpu,internal-keycache,experimental,zk-pok --no-deps -p $(TFHE_SPEC)
+		--features=$(TARGET_ARCH_FEATURE),boolean,shortint,integer,strings,gpu,internal-keycache,experimental,zk-pok --no-deps -p $(TFHE_SPEC)
 
 .PHONY: docs # Build rust doc alias for doc
 docs: doc
@@ -918,7 +919,7 @@ lint_doc: install_rs_check_toolchain
 	DOCS_RS=1 \
 	RUSTDOCFLAGS="--html-in-header katex-header.html -Dwarnings" \
 	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" doc \
-		--features=$(TARGET_ARCH_FEATURE),boolean,shortint,integer,gpu,internal-keycache,experimental,zk-pok -p $(TFHE_SPEC) --no-deps
+		--features=$(TARGET_ARCH_FEATURE),boolean,shortint,integer,strings,gpu,internal-keycache,experimental,zk-pok -p $(TFHE_SPEC) --no-deps
 
 .PHONY: lint_docs # Build rust doc with linting enabled alias for lint_doc
 lint_docs: lint_doc
@@ -938,11 +939,11 @@ check_md_docs_are_tested:
 
 .PHONY: check_intra_md_links # Checks broken internal links in Markdown docs
 check_intra_md_links: install_mlc
-	mlc --offline --match-file-extension tfhe/docs
+	mlc --offline --match-file-extension $(TFHE_SRC)/docs
 
 .PHONY: check_md_links # Checks all broken links in Markdown docs
 check_md_links: install_mlc
-	mlc --match-file-extension tfhe/docs
+	mlc --match-file-extension $(TFHE_SRC)/docs
 
 .PHONY: check_compile_tests # Build tests in debug without running them
 check_compile_tests: install_rs_build_toolchain
@@ -967,7 +968,7 @@ check_compile_tests_benches_gpu: install_rs_build_toolchain
 
 .PHONY: test_nodejs_wasm_api # Run tests for the nodejs on wasm API
 test_nodejs_wasm_api: build_node_js_api
-	cd tfhe/js_on_wasm_tests && npm install && npm run test
+	cd $(TFHE_SRC)/tests/js_on_wasm && npm install && npm run test
 
 .PHONY: test_nodejs_wasm_api_ci # Run tests for the nodejs on wasm API
 test_nodejs_wasm_api_ci: build_node_js_api
@@ -1275,9 +1276,9 @@ write_params_to_file: install_rs_check_toolchain
 
 .PHONY: clone_backward_compat_data # Clone the data repo needed for backward compatibility tests
 clone_backward_compat_data:
-	./scripts/clone_backward_compat_data.sh $(BACKWARD_COMPAT_DATA_URL) $(BACKWARD_COMPAT_DATA_BRANCH) tfhe/$(BACKWARD_COMPAT_DATA_DIR)
+	./scripts/clone_backward_compat_data.sh $(BACKWARD_COMPAT_DATA_URL) $(BACKWARD_COMPAT_DATA_BRANCH) $(TFHE_SRC)/$(BACKWARD_COMPAT_DATA_DIR)
 
-tfhe/$(BACKWARD_COMPAT_DATA_DIR): clone_backward_compat_data
+$(TFHE_SRC)/$(BACKWARD_COMPAT_DATA_DIR): clone_backward_compat_data
 
 #
 # Real use case examples

apps/trivium/Cargo.toml

@@ -9,11 +9,11 @@ edition = "2021"
 rayon = { version = "1.7.0"}
 
 [target.'cfg(target_arch = "x86_64")'.dependencies.tfhe]
-path = "../../tfhe"
+path = "../../crates/tfhe"
 features = [ "boolean", "shortint", "integer", "x86_64" ]
 
 [target.'cfg(target_arch = "aarch64")'.dependencies.tfhe]
-path = "../../tfhe"
+path = "../../crates/tfhe"
 features = [ "boolean", "shortint", "integer", "aarch64-unix" ]
 
 [dev-dependencies]

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

@@ -42,6 +42,8 @@ void cuda_destroy_stream(cudaStream_t stream, uint32_t gpu_index);
 
 void cuda_synchronize_stream(cudaStream_t stream, uint32_t gpu_index);
 
+uint32_t cuda_is_available();
+
 void *cuda_malloc(uint64_t size, uint32_t gpu_index);
 
 void *cuda_malloc_async(uint64_t size, cudaStream_t stream, uint32_t gpu_index);

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

@@ -102,13 +102,12 @@ template <typename Torus> struct int_decompression {
       };
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0],
-          carry_extract_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], carry_extract_lut->get_lut(0, 0),
           encryption_params.glwe_dimension, encryption_params.polynomial_size,
           encryption_params.message_modulus, encryption_params.carry_modulus,
           carry_extract_f);
 
-      carry_extract_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      carry_extract_lut->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,

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

@@ -559,7 +559,7 @@ template <typename Torus> struct int_bit_extract_luts_buffer {
         };
 
         generate_device_accumulator<Torus>(
-            streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], i),
+            streams[0], gpu_indexes[0], lut->get_lut(0, i),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, operator_f);
       }
@@ -574,11 +574,11 @@ template <typename Torus> struct int_bit_extract_luts_buffer {
         for (int i = 0; i < bits_per_block; i++)
           h_lut_indexes[i + j * bits_per_block] = i;
       }
-      cuda_memcpy_async_to_gpu(
-          lut->get_lut_indexes(gpu_indexes[0], 0), h_lut_indexes,
-          num_radix_blocks * bits_per_block * sizeof(Torus), streams[0],
-          gpu_indexes[0]);
-      lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      cuda_memcpy_async_to_gpu(lut->get_lut_indexes(0, 0), h_lut_indexes,
+                               num_radix_blocks * bits_per_block *
+                                   sizeof(Torus),
+                               streams[0], gpu_indexes[0]);
+      lut->broadcast_lut(streams, gpu_indexes, 0);
 
       /**
        * the input indexes should take the first bits_per_block PBS to target
@@ -757,17 +757,17 @@ template <typename Torus> struct int_shift_and_rotate_buffer {
       };
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], mux_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], mux_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, mux_lut_f);
-      mux_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      mux_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       auto cleaning_lut_f = [](Torus x) -> Torus { return x; };
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], cleaning_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], cleaning_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, cleaning_lut_f);
-      cleaning_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      cleaning_lut->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 
@@ -818,8 +818,8 @@ template <typename Torus> struct int_fullprop_buffer {
       };
 
       //
-      Torus *lut_buffer_message = lut->get_lut(gpu_indexes[0], 0);
-      Torus *lut_buffer_carry = lut->get_lut(gpu_indexes[0], 1);
+      Torus *lut_buffer_message = lut->get_lut(0, 0);
+      Torus *lut_buffer_carry = lut->get_lut(0, 1);
 
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut_buffer_message, params.glwe_dimension,
@@ -835,11 +835,11 @@ template <typename Torus> struct int_fullprop_buffer {
       Torus *h_lwe_indexes = (Torus *)malloc(lwe_indexes_size);
       for (int i = 0; i < 2; i++)
         h_lwe_indexes[i] = i;
-      Torus *lwe_indexes = lut->get_lut_indexes(gpu_indexes[0], 0);
+      Torus *lwe_indexes = lut->get_lut_indexes(0, 0);
       cuda_memcpy_async_to_gpu(lwe_indexes, h_lwe_indexes, lwe_indexes_size,
                                streams[0], gpu_indexes[0]);
 
-      lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      lut->broadcast_lut(streams, gpu_indexes, 0);
 
       // Temporary arrays
       Torus small_vector_size =
@@ -940,9 +940,8 @@ template <typename Torus> struct int_legacy_sc_prop_memory {
         new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count, params, 1,
                                  num_radix_blocks, luts_array);
 
-    auto lut_does_block_generate_carry = luts_array->get_lut(gpu_indexes[0], 0);
-    auto lut_does_block_generate_or_propagate =
-        luts_array->get_lut(gpu_indexes[0], 1);
+    auto lut_does_block_generate_carry = luts_array->get_lut(0, 0);
+    auto lut_does_block_generate_or_propagate = luts_array->get_lut(0, 1);
 
     // generate luts (aka accumulators)
     generate_device_accumulator<Torus>(
@@ -954,24 +953,21 @@ template <typename Torus> struct int_legacy_sc_prop_memory {
         glwe_dimension, polynomial_size, message_modulus, carry_modulus,
         f_lut_does_block_generate_or_propagate);
     cuda_set_value_async<Torus>(streams[0], gpu_indexes[0],
-                                luts_array->get_lut_indexes(gpu_indexes[0], 1),
-                                1, num_radix_blocks - 1);
+                                luts_array->get_lut_indexes(0, 1), 1,
+                                num_radix_blocks - 1);
 
     generate_device_accumulator_bivariate<Torus>(
-        streams[0], gpu_indexes[0],
-        luts_carry_propagation_sum->get_lut(gpu_indexes[0], 0), glwe_dimension,
-        polynomial_size, message_modulus, carry_modulus,
+        streams[0], gpu_indexes[0], luts_carry_propagation_sum->get_lut(0, 0),
+        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
         f_luts_carry_propagation_sum);
 
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], message_acc->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
-        f_message_acc);
+        streams[0], gpu_indexes[0], message_acc->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus, f_message_acc);
 
-    luts_array->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
-    luts_carry_propagation_sum->broadcast_lut(streams, gpu_indexes,
-                                              gpu_indexes[0]);
-    message_acc->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array->broadcast_lut(streams, gpu_indexes, 0);
+    luts_carry_propagation_sum->broadcast_lut(streams, gpu_indexes, 0);
+    message_acc->broadcast_lut(streams, gpu_indexes, 0);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -1059,9 +1055,8 @@ template <typename Torus> struct int_overflowing_sub_memory {
         new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count, params, 1,
                                  num_radix_blocks, luts_array);
 
-    auto lut_does_block_generate_carry = luts_array->get_lut(gpu_indexes[0], 0);
-    auto lut_does_block_generate_or_propagate =
-        luts_array->get_lut(gpu_indexes[0], 1);
+    auto lut_does_block_generate_carry = luts_array->get_lut(0, 0);
+    auto lut_does_block_generate_or_propagate = luts_array->get_lut(0, 1);
 
     // generate luts (aka accumulators)
     generate_device_accumulator<Torus>(
@@ -1073,24 +1068,21 @@ template <typename Torus> struct int_overflowing_sub_memory {
         glwe_dimension, polynomial_size, message_modulus, carry_modulus,
         f_lut_does_block_generate_or_propagate);
     cuda_set_value_async<Torus>(streams[0], gpu_indexes[0],
-                                luts_array->get_lut_indexes(gpu_indexes[0], 1),
-                                1, num_radix_blocks - 1);
+                                luts_array->get_lut_indexes(0, 1), 1,
+                                num_radix_blocks - 1);
 
     generate_device_accumulator_bivariate<Torus>(
-        streams[0], gpu_indexes[0],
-        luts_borrow_propagation_sum->get_lut(gpu_indexes[0], 0), glwe_dimension,
-        polynomial_size, message_modulus, carry_modulus,
+        streams[0], gpu_indexes[0], luts_borrow_propagation_sum->get_lut(0, 0),
+        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
         f_luts_borrow_propagation_sum);
 
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], message_acc->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
-        f_message_acc);
+        streams[0], gpu_indexes[0], message_acc->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus, f_message_acc);
 
-    luts_array->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
-    luts_borrow_propagation_sum->broadcast_lut(streams, gpu_indexes,
-                                               gpu_indexes[0]);
-    message_acc->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array->broadcast_lut(streams, gpu_indexes, 0);
+    luts_borrow_propagation_sum->broadcast_lut(streams, gpu_indexes, 0);
+    message_acc->broadcast_lut(streams, gpu_indexes, 0);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -1252,20 +1244,18 @@ template <typename Torus> struct int_seq_group_prop_memory {
       auto f_lut_sequential = [index](Torus propa_cum_sum_block) {
         return (propa_cum_sum_block >> (index + 1)) & 1;
       };
-      auto seq_lut = lut_sequential_algorithm->get_lut(gpu_indexes[0], index);
+      auto seq_lut = lut_sequential_algorithm->get_lut(0, index);
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], seq_lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_lut_sequential);
       h_seq_lut_indexes[index] = index;
     }
-    Torus *seq_lut_indexes =
-        lut_sequential_algorithm->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *seq_lut_indexes = lut_sequential_algorithm->get_lut_indexes(0, 0);
     cuda_memcpy_async_to_gpu(seq_lut_indexes, h_seq_lut_indexes,
                              num_seq_luts * sizeof(Torus), streams[0],
                              gpu_indexes[0]);
 
-    lut_sequential_algorithm->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
+    lut_sequential_algorithm->broadcast_lut(streams, gpu_indexes, 0);
     free(h_seq_lut_indexes);
   };
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -1312,12 +1302,12 @@ template <typename Torus> struct int_hs_group_prop_memory {
         new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count, params, 1,
                                  num_groups, allocate_gpu_memory);
 
-    auto hillis_steele_lut = lut_hillis_steele->get_lut(gpu_indexes[0], 0);
+    auto hillis_steele_lut = lut_hillis_steele->get_lut(0, 0);
     generate_device_accumulator_bivariate<Torus>(
         streams[0], gpu_indexes[0], hillis_steele_lut, glwe_dimension,
         polynomial_size, message_modulus, carry_modulus, f_lut_hillis_steele);
 
-    lut_hillis_steele->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut_hillis_steele->broadcast_lut(streams, gpu_indexes, 0);
   };
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
                uint32_t gpu_count) {
@@ -1382,7 +1372,7 @@ template <typename Torus> struct int_shifted_blocks_and_states_memory {
     std::vector<std::function<Torus(Torus)>> f_first_grouping_luts = {
         f_first_block_state, f_shift_block};
 
-    auto first_block_lut = luts_array_first_step->get_lut(gpu_indexes[0], 0);
+    auto first_block_lut = luts_array_first_step->get_lut(0, 0);
 
     generate_many_lut_device_accumulator<Torus>(
         streams[0], gpu_indexes[0], first_block_lut, glwe_dimension,
@@ -1403,7 +1393,7 @@ template <typename Torus> struct int_shifted_blocks_and_states_memory {
       };
       std::vector<std::function<Torus(Torus)>> f_grouping_luts = {
           f_state, f_shift_block};
-      auto lut = luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_first_step->get_lut(0, lut_id);
       generate_many_lut_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_grouping_luts);
@@ -1426,7 +1416,7 @@ template <typename Torus> struct int_shifted_blocks_and_states_memory {
       std::vector<std::function<Torus(Torus)>> f_grouping_luts = {
           f_state, f_shift_block};
 
-      auto lut = luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_first_step->get_lut(0, lut_id);
 
       generate_many_lut_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
@@ -1443,8 +1433,7 @@ template <typename Torus> struct int_shifted_blocks_and_states_memory {
 
     uint32_t lut_id = num_luts_first_step - 1; // The last lut of the first step
 
-    auto last_block_lut =
-        luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+    auto last_block_lut = luts_array_first_step->get_lut(0, lut_id);
 
     std::vector<std::function<Torus(Torus)>> f_last_grouping_luts = {
         f_last_block_state, f_shift_block};
@@ -1476,13 +1465,12 @@ template <typename Torus> struct int_shifted_blocks_and_states_memory {
     }
 
     // copy the indexes to the gpu
-    Torus *lut_indexes =
-        luts_array_first_step->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *lut_indexes = luts_array_first_step->get_lut_indexes(0, 0);
     cuda_memcpy_async_to_gpu(lut_indexes, h_lut_indexes, lut_indexes_size,
                              streams[0], gpu_indexes[0]);
     // Do I need to do something else for the multi-gpu?
 
-    luts_array_first_step->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array_first_step->broadcast_lut(streams, gpu_indexes, 0);
 
     free(h_lut_indexes);
   };
@@ -1604,7 +1592,7 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
         }
       };
 
-      auto lut = luts_array_second_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_second_step->get_lut(0, lut_id);
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_first_grouping_inner_propagation);
@@ -1616,8 +1604,7 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
     };
 
     int lut_id = grouping_size - 1;
-    auto lut_first_group_outer =
-        luts_array_second_step->get_lut(gpu_indexes[0], lut_id);
+    auto lut_first_group_outer = luts_array_second_step->get_lut(0, lut_id);
     generate_device_accumulator<Torus>(
         streams[0], gpu_indexes[0], lut_first_group_outer, glwe_dimension,
         polynomial_size, message_modulus, carry_modulus,
@@ -1639,7 +1626,7 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
         }
       };
 
-      auto lut = luts_array_second_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_second_step->get_lut(0, lut_id);
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_other_groupings_inner_propagation);
@@ -1658,7 +1645,7 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
           }
         };
 
-        auto lut = luts_array_second_step->get_lut(gpu_indexes[0], lut_id);
+        auto lut = luts_array_second_step->get_lut(0, lut_id);
         generate_device_accumulator<Torus>(
             streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
             message_modulus, carry_modulus, f_group_propagation);
@@ -1673,7 +1660,7 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
         }
       };
 
-      auto lut = luts_array_second_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_second_step->get_lut(0, lut_id);
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_group_propagation);
@@ -1718,15 +1705,14 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
     }
 
     // copy the indexes to the gpu
-    Torus *second_lut_indexes =
-        luts_array_second_step->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *second_lut_indexes = luts_array_second_step->get_lut_indexes(0, 0);
     cuda_memcpy_async_to_gpu(second_lut_indexes, h_second_lut_indexes,
                              lut_indexes_size, streams[0], gpu_indexes[0]);
 
     cuda_memcpy_async_to_gpu(scalar_array_cum_sum, h_scalar_array_cum_sum,
                              num_radix_blocks * sizeof(Torus), streams[0],
                              gpu_indexes[0]);
-    luts_array_second_step->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array_second_step->broadcast_lut(streams, gpu_indexes, 0);
 
     if (use_sequential_algorithm_to_resolver_group_carries) {
 
@@ -1748,13 +1734,12 @@ template <typename Torus> struct int_prop_simu_group_carries_memory {
   void update_lut_indexes(cudaStream_t const *streams,
                           uint32_t const *gpu_indexes, Torus *new_lut_indexes,
                           Torus *new_scalars, uint32_t new_num_blocks) {
-    Torus *lut_indexes =
-        luts_array_second_step->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *lut_indexes = luts_array_second_step->get_lut_indexes(0, 0);
     cuda_memcpy_async_gpu_to_gpu(lut_indexes, new_lut_indexes,
                                  new_num_blocks * sizeof(Torus), streams[0],
                                  gpu_indexes[0]);
 
-    luts_array_second_step->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array_second_step->broadcast_lut(streams, gpu_indexes, 0);
 
     cuda_memcpy_async_gpu_to_gpu(scalar_array_cum_sum, new_scalars,
                                  new_num_blocks * sizeof(Torus), streams[0],
@@ -1857,13 +1842,13 @@ template <typename Torus> struct int_sc_prop_memory {
       return (block >> 1) % message_modulus;
     };
 
-    auto extract_lut = lut_message_extract->get_lut(gpu_indexes[0], 0);
+    auto extract_lut = lut_message_extract->get_lut(0, 0);
 
     generate_device_accumulator<Torus>(
         streams[0], gpu_indexes[0], extract_lut, glwe_dimension,
         polynomial_size, message_modulus, carry_modulus, f_message_extract);
 
-    lut_message_extract->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut_message_extract->broadcast_lut(streams, gpu_indexes, 0);
 
     // This store a single block that with be used to store the overflow or
     // carry results
@@ -1914,15 +1899,13 @@ template <typename Torus> struct int_sc_prop_memory {
         return output1 << 3 | output2 << 2;
       };
 
-      auto overflow_flag_prep_lut =
-          lut_overflow_flag_prep->get_lut(gpu_indexes[0], 0);
+      auto overflow_flag_prep_lut = lut_overflow_flag_prep->get_lut(0, 0);
 
       generate_device_accumulator_bivariate<Torus>(
           streams[0], gpu_indexes[0], overflow_flag_prep_lut, glwe_dimension,
           polynomial_size, message_modulus, carry_modulus, f_overflow_fp);
 
-      lut_overflow_flag_prep->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
+      lut_overflow_flag_prep->broadcast_lut(streams, gpu_indexes, 0);
     }
 
     // For the final cleanup in case of overflow or carry (it seems that I can)
@@ -1947,15 +1930,13 @@ template <typename Torus> struct int_sc_prop_memory {
           return does_overflow_if_carry_is_0;
         }
       };
-      auto overflow_flag_last =
-          lut_overflow_flag_last->get_lut(gpu_indexes[0], 0);
+      auto overflow_flag_last = lut_overflow_flag_last->get_lut(0, 0);
 
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], overflow_flag_last, glwe_dimension,
           polynomial_size, message_modulus, carry_modulus, f_overflow_last);
 
-      lut_overflow_flag_last->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
+      lut_overflow_flag_last->broadcast_lut(streams, gpu_indexes, 0);
     }
     if (requested_flag == outputFlag::FLAG_CARRY) { // Carry case
       lut_carry_flag_last = new int_radix_lut<Torus>(
@@ -1964,13 +1945,13 @@ template <typename Torus> struct int_sc_prop_memory {
       auto f_carry_last = [](Torus block) -> Torus {
         return ((block >> 2) & 1);
       };
-      auto carry_flag_last = lut_carry_flag_last->get_lut(gpu_indexes[0], 0);
+      auto carry_flag_last = lut_carry_flag_last->get_lut(0, 0);
 
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], carry_flag_last, glwe_dimension,
           polynomial_size, message_modulus, carry_modulus, f_carry_last);
 
-      lut_carry_flag_last->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      lut_carry_flag_last->broadcast_lut(streams, gpu_indexes, 0);
     }
 
     active_gpu_count = get_active_gpu_count(num_radix_blocks, gpu_count);
@@ -2110,7 +2091,7 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
     std::vector<std::function<Torus(Torus)>> f_first_grouping_luts = {
         f_first_block_state, f_shift_block};
 
-    auto first_block_lut = luts_array_first_step->get_lut(gpu_indexes[0], 0);
+    auto first_block_lut = luts_array_first_step->get_lut(0, 0);
 
     generate_many_lut_device_accumulator<Torus>(
         streams[0], gpu_indexes[0], first_block_lut, glwe_dimension,
@@ -2131,7 +2112,7 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
       };
       std::vector<std::function<Torus(Torus)>> f_grouping_luts = {
           f_state, f_shift_block};
-      auto lut = luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_first_step->get_lut(0, lut_id);
       generate_many_lut_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
           message_modulus, carry_modulus, f_grouping_luts);
@@ -2154,7 +2135,7 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
       std::vector<std::function<Torus(Torus)>> f_grouping_luts = {
           f_state, f_shift_block};
 
-      auto lut = luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+      auto lut = luts_array_first_step->get_lut(0, lut_id);
 
       generate_many_lut_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], lut, glwe_dimension, polynomial_size,
@@ -2170,8 +2151,7 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
 
     uint32_t lut_id = num_luts_first_step - 1; // The last lut of the first step
 
-    auto last_block_lut =
-        luts_array_first_step->get_lut(gpu_indexes[0], lut_id);
+    auto last_block_lut = luts_array_first_step->get_lut(0, lut_id);
 
     std::vector<std::function<Torus(Torus)>> f_last_grouping_luts = {
         f_last_block_state, f_shift_block};
@@ -2202,13 +2182,12 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
       }
     }
     // copy the indexes to the gpu
-    Torus *lut_indexes =
-        luts_array_first_step->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *lut_indexes = luts_array_first_step->get_lut_indexes(0, 0);
     cuda_memcpy_async_to_gpu(lut_indexes, h_lut_indexes, lut_indexes_size,
                              streams[0], gpu_indexes[0]);
     // Do I need to do something else for the multi-gpu?
 
-    luts_array_first_step->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array_first_step->broadcast_lut(streams, gpu_indexes, 0);
 
     free(h_lut_indexes);
   };
@@ -2217,12 +2196,11 @@ template <typename Torus> struct int_shifted_blocks_and_borrow_states_memory {
   void update_lut_indexes(cudaStream_t const *streams,
                           uint32_t const *gpu_indexes, Torus *new_lut_indexes,
                           uint32_t new_num_blocks) {
-    Torus *lut_indexes =
-        luts_array_first_step->get_lut_indexes(gpu_indexes[0], 0);
+    Torus *lut_indexes = luts_array_first_step->get_lut_indexes(0, 0);
     cuda_memcpy_async_gpu_to_gpu(lut_indexes, new_lut_indexes,
                                  new_num_blocks * sizeof(Torus), streams[0],
                                  gpu_indexes[0]);
-    luts_array_first_step->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array_first_step->broadcast_lut(streams, gpu_indexes, 0);
   }
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
                uint32_t gpu_count) {
@@ -2309,13 +2287,13 @@ template <typename Torus> struct int_borrow_prop_memory {
       return (block >> 1) % message_modulus;
     };
 
-    auto extract_lut = lut_message_extract->get_lut(gpu_indexes[0], 0);
+    auto extract_lut = lut_message_extract->get_lut(0, 0);
 
     generate_device_accumulator<Torus>(
         streams[0], gpu_indexes[0], extract_lut, glwe_dimension,
         polynomial_size, message_modulus, carry_modulus, f_message_extract);
 
-    lut_message_extract->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut_message_extract->broadcast_lut(streams, gpu_indexes, 0);
 
     if (compute_overflow) {
       lut_borrow_flag =
@@ -2326,13 +2304,13 @@ template <typename Torus> struct int_borrow_prop_memory {
         return ((block >> 2) & 1);
       };
 
-      auto borrow_flag_lut = lut_borrow_flag->get_lut(gpu_indexes[0], 0);
+      auto borrow_flag_lut = lut_borrow_flag->get_lut(0, 0);
 
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0], borrow_flag_lut, glwe_dimension,
           polynomial_size, message_modulus, carry_modulus, f_borrow_flag);
 
-      lut_borrow_flag->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      lut_borrow_flag->broadcast_lut(streams, gpu_indexes, 0);
     }
 
     active_gpu_count = get_active_gpu_count(num_radix_blocks, gpu_count);
@@ -2485,12 +2463,10 @@ template <typename Torus> struct int_mul_memory {
           new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count, params, 1,
                                    num_radix_blocks, allocate_gpu_memory);
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0],
-          zero_out_predicate_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], zero_out_predicate_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, zero_out_predicate_lut_f);
-      zero_out_predicate_lut->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
+      zero_out_predicate_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       zero_out_mem = new int_zero_out_if_buffer<Torus>(
           streams, gpu_indexes, gpu_count, params, num_radix_blocks,
@@ -2533,8 +2509,8 @@ template <typename Torus> struct int_mul_memory {
     luts_array =
         new int_radix_lut<Torus>(streams, gpu_indexes, gpu_count, params, 2,
                                  total_block_count, allocate_gpu_memory);
-    auto lsb_acc = luts_array->get_lut(gpu_indexes[0], 0);
-    auto msb_acc = luts_array->get_lut(gpu_indexes[0], 1);
+    auto lsb_acc = luts_array->get_lut(0, 0);
+    auto msb_acc = luts_array->get_lut(0, 1);
 
     // define functions for each accumulator
     auto lut_f_lsb = [message_modulus](Torus x, Torus y) -> Torus {
@@ -2558,10 +2534,10 @@ template <typename Torus> struct int_mul_memory {
     // for message and carry default lut_indexes_vec is fine
     cuda_set_value_async<Torus>(
         streams[0], gpu_indexes[0],
-        luts_array->get_lut_indexes(gpu_indexes[0], lsb_vector_block_count), 1,
+        luts_array->get_lut_indexes(0, lsb_vector_block_count), 1,
         msb_vector_block_count);
 
-    luts_array->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_array->broadcast_lut(streams, gpu_indexes, 0);
     // create memory object for sum ciphertexts
     sum_ciphertexts_mem = new int_sum_ciphertexts_vec_memory<Torus>(
         streams, gpu_indexes, gpu_count, params, num_radix_blocks,
@@ -2690,11 +2666,10 @@ template <typename Torus> struct int_logical_scalar_shift_buffer {
 
         // right shift
         generate_device_accumulator_bivariate<Torus>(
-            streams[0], gpu_indexes[0],
-            cur_lut_bivariate->get_lut(gpu_indexes[0], 0),
+            streams[0], gpu_indexes[0], cur_lut_bivariate->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, shift_lut_f);
-        cur_lut_bivariate->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        cur_lut_bivariate->broadcast_lut(streams, gpu_indexes, 0);
 
         lut_buffers_bivariate.push_back(cur_lut_bivariate);
       }
@@ -2777,11 +2752,10 @@ template <typename Torus> struct int_logical_scalar_shift_buffer {
 
         // right shift
         generate_device_accumulator_bivariate<Torus>(
-            streams[0], gpu_indexes[0],
-            cur_lut_bivariate->get_lut(gpu_indexes[0], 0),
+            streams[0], gpu_indexes[0], cur_lut_bivariate->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, shift_lut_f);
-        cur_lut_bivariate->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        cur_lut_bivariate->broadcast_lut(streams, gpu_indexes, 0);
 
         lut_buffers_bivariate.push_back(cur_lut_bivariate);
       }
@@ -2883,11 +2857,10 @@ template <typename Torus> struct int_arithmetic_scalar_shift_buffer {
 
         generate_device_accumulator<Torus>(
             streams[0], gpu_indexes[0],
-            shift_last_block_lut_univariate->get_lut(gpu_indexes[0], 0),
+            shift_last_block_lut_univariate->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, last_block_lut_f);
-        shift_last_block_lut_univariate->broadcast_lut(streams, gpu_indexes,
-                                                       gpu_indexes[0]);
+        shift_last_block_lut_univariate->broadcast_lut(streams, gpu_indexes, 0);
 
         lut_buffers_univariate.push_back(shift_last_block_lut_univariate);
       }
@@ -2907,11 +2880,10 @@ template <typename Torus> struct int_arithmetic_scalar_shift_buffer {
 
       generate_device_accumulator<Torus>(
           streams[0], gpu_indexes[0],
-          padding_block_lut_univariate->get_lut(gpu_indexes[0], 0),
-          params.glwe_dimension, params.polynomial_size, params.message_modulus,
-          params.carry_modulus, padding_block_lut_f);
-      padding_block_lut_univariate->broadcast_lut(streams, gpu_indexes,
-                                                  gpu_indexes[0]);
+          padding_block_lut_univariate->get_lut(0, 0), params.glwe_dimension,
+          params.polynomial_size, params.message_modulus, params.carry_modulus,
+          padding_block_lut_f);
+      padding_block_lut_univariate->broadcast_lut(streams, gpu_indexes, 0);
 
       lut_buffers_univariate.push_back(padding_block_lut_univariate);
 
@@ -2948,11 +2920,10 @@ template <typename Torus> struct int_arithmetic_scalar_shift_buffer {
 
         generate_device_accumulator_bivariate<Torus>(
             streams[0], gpu_indexes[0],
-            shift_blocks_lut_bivariate->get_lut(gpu_indexes[0], 0),
-            params.glwe_dimension, params.polynomial_size,
-            params.message_modulus, params.carry_modulus, blocks_lut_f);
-        shift_blocks_lut_bivariate->broadcast_lut(streams, gpu_indexes,
-                                                  gpu_indexes[0]);
+            shift_blocks_lut_bivariate->get_lut(0, 0), params.glwe_dimension,
+            params.polynomial_size, params.message_modulus,
+            params.carry_modulus, blocks_lut_f);
+        shift_blocks_lut_bivariate->broadcast_lut(streams, gpu_indexes, 0);
 
         lut_buffers_bivariate.push_back(shift_blocks_lut_bivariate);
       }
@@ -3043,26 +3014,23 @@ template <typename Torus> struct int_cmux_buffer {
                                    num_radix_blocks, allocate_gpu_memory);
 
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0], predicate_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], predicate_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, lut_f);
 
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0],
-          inverted_predicate_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], inverted_predicate_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, inverted_lut_f);
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0],
-          message_extract_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], message_extract_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, message_extract_lut_f);
 
-      predicate_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
-      inverted_predicate_lut->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
-      message_extract_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      predicate_lut->broadcast_lut(streams, gpu_indexes, 0);
+      inverted_predicate_lut->broadcast_lut(streams, gpu_indexes, 0);
+      message_extract_lut->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 
@@ -3171,11 +3139,11 @@ template <typename Torus> struct int_comparison_eq_buffer {
                                    num_radix_blocks, allocate_gpu_memory);
 
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0], operator_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], operator_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, operator_f);
 
-      operator_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      operator_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       // f(x) -> x == 0
       Torus total_modulus = params.message_modulus * params.carry_modulus;
@@ -3188,12 +3156,11 @@ template <typename Torus> struct int_comparison_eq_buffer {
                                    num_radix_blocks, allocate_gpu_memory);
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0],
-          is_non_zero_lut->get_lut(gpu_indexes[0], 0), params.glwe_dimension,
-          params.polynomial_size, params.message_modulus, params.carry_modulus,
-          is_non_zero_lut_f);
+          streams[0], gpu_indexes[0], is_non_zero_lut->get_lut(0, 0),
+          params.glwe_dimension, params.polynomial_size, params.message_modulus,
+          params.carry_modulus, is_non_zero_lut_f);
 
-      is_non_zero_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      is_non_zero_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       // Scalar may have up to num_radix_blocks blocks
       scalar_comparison_luts = new int_radix_lut<Torus>(
@@ -3205,7 +3172,7 @@ template <typename Torus> struct int_comparison_eq_buffer {
           return operator_f(i, x);
         };
 
-        Torus *lut = scalar_comparison_luts->get_lut(gpu_indexes[0], i);
+        Torus *lut = scalar_comparison_luts->get_lut(0, i);
 
         generate_device_accumulator<Torus>(
             streams[0], gpu_indexes[0], lut, params.glwe_dimension,
@@ -3213,8 +3180,7 @@ template <typename Torus> struct int_comparison_eq_buffer {
             params.carry_modulus, lut_f);
       }
 
-      scalar_comparison_luts->broadcast_lut(streams, gpu_indexes,
-                                            gpu_indexes[0]);
+      scalar_comparison_luts->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 
@@ -3278,12 +3244,11 @@ template <typename Torus> struct int_tree_sign_reduction_buffer {
       tree_last_leaf_scalar_lut = new int_radix_lut<Torus>(
           streams, gpu_indexes, gpu_count, params, 1, 1, allocate_gpu_memory);
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0],
-          tree_inner_leaf_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], tree_inner_leaf_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, block_selector_f);
 
-      tree_inner_leaf_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      tree_inner_leaf_lut->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 
@@ -3456,11 +3421,11 @@ template <typename Torus> struct int_comparison_buffer {
                                    num_radix_blocks, allocate_gpu_memory);
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], identity_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], identity_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, identity_lut_f);
 
-      identity_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      identity_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       uint32_t total_modulus = params.message_modulus * params.carry_modulus;
       auto is_zero_f = [total_modulus](Torus x) -> Torus {
@@ -3472,11 +3437,11 @@ template <typename Torus> struct int_comparison_buffer {
                                    num_radix_blocks, allocate_gpu_memory);
 
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], is_zero_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], is_zero_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, is_zero_f);
 
-      is_zero_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      is_zero_lut->broadcast_lut(streams, gpu_indexes, 0);
 
       switch (op) {
       case COMPARISON_TYPE::MAX:
@@ -3550,11 +3515,11 @@ template <typename Torus> struct int_comparison_buffer {
         };
 
         generate_device_accumulator_bivariate<Torus>(
-            streams[0], gpu_indexes[0], signed_lut->get_lut(gpu_indexes[0], 0),
+            streams[0], gpu_indexes[0], signed_lut->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, signed_lut_f);
 
-        signed_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        signed_lut->broadcast_lut(streams, gpu_indexes, 0);
       }
     }
   }
@@ -3728,10 +3693,10 @@ template <typename Torus> struct unsigned_int_div_rem_memory {
 
       for (int j = 0; j < 2; j++) {
         generate_device_accumulator<Torus>(
-            streams[0], gpu_indexes[0], luts[j]->get_lut(gpu_indexes[0], 0),
+            streams[0], gpu_indexes[0], luts[j]->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, lut_f_masking);
-        luts[j]->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        luts[j]->broadcast_lut(streams, gpu_indexes, 0);
       }
     }
 
@@ -3752,10 +3717,10 @@ template <typename Torus> struct unsigned_int_div_rem_memory {
                                      message_extract_lut_2};
     for (int j = 0; j < 2; j++) {
       generate_device_accumulator<Torus>(
-          streams[0], gpu_indexes[0], luts[j]->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], luts[j]->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, lut_f_message_extract);
-      luts[j]->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      luts[j]->broadcast_lut(streams, gpu_indexes, 0);
     }
 
     // Give name to closures to improve readability
@@ -3783,14 +3748,14 @@ template <typename Torus> struct unsigned_int_div_rem_memory {
 
     generate_device_accumulator_bivariate_with_factor<Torus>(
         streams[0], gpu_indexes[0],
-        zero_out_if_overflow_did_not_happen[0]->get_lut(gpu_indexes[0], 0),
+        zero_out_if_overflow_did_not_happen[0]->get_lut(0, 0),
         params.glwe_dimension, params.polynomial_size, params.message_modulus,
         params.carry_modulus, cur_lut_f, 2);
     zero_out_if_overflow_did_not_happen[0]->broadcast_lut(streams, gpu_indexes,
                                                           0);
     generate_device_accumulator_bivariate_with_factor<Torus>(
         streams[0], gpu_indexes[0],
-        zero_out_if_overflow_did_not_happen[1]->get_lut(gpu_indexes[0], 0),
+        zero_out_if_overflow_did_not_happen[1]->get_lut(0, 0),
         params.glwe_dimension, params.polynomial_size, params.message_modulus,
         params.carry_modulus, cur_lut_f, 3);
     zero_out_if_overflow_did_not_happen[1]->broadcast_lut(streams, gpu_indexes,
@@ -3813,18 +3778,16 @@ template <typename Torus> struct unsigned_int_div_rem_memory {
 
     generate_device_accumulator_bivariate_with_factor<Torus>(
         streams[0], gpu_indexes[0],
-        zero_out_if_overflow_happened[0]->get_lut(gpu_indexes[0], 0),
-        params.glwe_dimension, params.polynomial_size, params.message_modulus,
-        params.carry_modulus, overflow_happened_f, 2);
-    zero_out_if_overflow_happened[0]->broadcast_lut(streams, gpu_indexes,
-                                                    gpu_indexes[0]);
+        zero_out_if_overflow_happened[0]->get_lut(0, 0), params.glwe_dimension,
+        params.polynomial_size, params.message_modulus, params.carry_modulus,
+        overflow_happened_f, 2);
+    zero_out_if_overflow_happened[0]->broadcast_lut(streams, gpu_indexes, 0);
     generate_device_accumulator_bivariate_with_factor<Torus>(
         streams[0], gpu_indexes[0],
-        zero_out_if_overflow_happened[1]->get_lut(gpu_indexes[0], 0),
-        params.glwe_dimension, params.polynomial_size, params.message_modulus,
-        params.carry_modulus, overflow_happened_f, 3);
-    zero_out_if_overflow_happened[1]->broadcast_lut(streams, gpu_indexes,
-                                                    gpu_indexes[0]);
+        zero_out_if_overflow_happened[1]->get_lut(0, 0), params.glwe_dimension,
+        params.polynomial_size, params.message_modulus, params.carry_modulus,
+        overflow_happened_f, 3);
+    zero_out_if_overflow_happened[1]->broadcast_lut(streams, gpu_indexes, 0);
 
     // merge_overflow_flags_luts
     merge_overflow_flags_luts = new int_radix_lut<Torus> *[num_bits_in_message];
@@ -3838,11 +3801,10 @@ template <typename Torus> struct unsigned_int_div_rem_memory {
 
       generate_device_accumulator_bivariate<Torus>(
           streams[0], gpu_indexes[0],
-          merge_overflow_flags_luts[i]->get_lut(gpu_indexes[0], 0),
-          params.glwe_dimension, params.polynomial_size, params.message_modulus,
-          params.carry_modulus, lut_f_bit);
-      merge_overflow_flags_luts[i]->broadcast_lut(streams, gpu_indexes,
-                                                  gpu_indexes[0]);
+          merge_overflow_flags_luts[i]->get_lut(0, 0), params.glwe_dimension,
+          params.polynomial_size, params.message_modulus, params.carry_modulus,
+          lut_f_bit);
+      merge_overflow_flags_luts[i]->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 
@@ -4156,11 +4118,10 @@ template <typename Torus> struct int_last_block_inner_propagate_memory {
 
     generate_device_accumulator_bivariate<Torus>(
         streams[0], gpu_indexes[0],
-        last_block_inner_propagation_lut->get_lut(gpu_indexes[0], 0),
-        params.glwe_dimension, params.polynomial_size, message_modulus,
-        params.carry_modulus, f_last_block_inner_propagation_lut);
-    last_block_inner_propagation_lut->broadcast_lut(streams, gpu_indexes,
-                                                    gpu_indexes[0]);
+        last_block_inner_propagation_lut->get_lut(0, 0), params.glwe_dimension,
+        params.polynomial_size, message_modulus, params.carry_modulus,
+        f_last_block_inner_propagation_lut);
+    last_block_inner_propagation_lut->broadcast_lut(streams, gpu_indexes, 0);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -4217,11 +4178,10 @@ template <typename Torus> struct int_resolve_signed_overflow_memory {
         streams, gpu_indexes, gpu_count, params, 1, 1, allocate_gpu_memory);
 
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0],
-        resolve_overflow_lut->get_lut(gpu_indexes[0], 0), params.glwe_dimension,
-        params.polynomial_size, message_modulus, params.carry_modulus,
-        f_resolve_overflow_lut);
-    resolve_overflow_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        streams[0], gpu_indexes[0], resolve_overflow_lut->get_lut(0, 0),
+        params.glwe_dimension, params.polynomial_size, message_modulus,
+        params.carry_modulus, f_resolve_overflow_lut);
+    resolve_overflow_lut->broadcast_lut(streams, gpu_indexes, 0);
   }
 
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
@@ -4264,10 +4224,10 @@ template <typename Torus> struct int_bitop_buffer {
         };
 
         generate_device_accumulator_bivariate<Torus>(
-            streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
+            streams[0], gpu_indexes[0], lut->get_lut(0, 0),
             params.glwe_dimension, params.polynomial_size,
             params.message_modulus, params.carry_modulus, lut_bivariate_f);
-        lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        lut->broadcast_lut(streams, gpu_indexes, 0);
       }
       break;
     default:
@@ -4277,7 +4237,7 @@ template <typename Torus> struct int_bitop_buffer {
                                      allocate_gpu_memory);
 
       for (int i = 0; i < params.message_modulus; i++) {
-        auto lut_block = lut->get_lut(gpu_indexes[0], i);
+        auto lut_block = lut->get_lut(0, i);
         auto rhs = i;
 
         auto lut_univariate_scalar_f = [op, rhs](Torus x) -> Torus {
@@ -4296,7 +4256,7 @@ template <typename Torus> struct int_bitop_buffer {
             streams[0], gpu_indexes[0], lut_block, params.glwe_dimension,
             params.polynomial_size, params.message_modulus,
             params.carry_modulus, lut_univariate_scalar_f);
-        lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        lut->broadcast_lut(streams, gpu_indexes, 0);
       }
     }
   }
@@ -4539,12 +4499,10 @@ template <typename Torus> struct int_div_rem_memory {
           streams, gpu_indexes, gpu_count, params, 1, 1, true);
 
       generate_device_accumulator_bivariate<Torus>(
-          streams[0], gpu_indexes[0],
-          compare_signed_bits_lut->get_lut(gpu_indexes[0], 0),
+          streams[0], gpu_indexes[0], compare_signed_bits_lut->get_lut(0, 0),
           params.glwe_dimension, params.polynomial_size, params.message_modulus,
           params.carry_modulus, f_compare_extracted_signed_bits);
-      compare_signed_bits_lut->broadcast_lut(streams, gpu_indexes,
-                                             gpu_indexes[0]);
+      compare_signed_bits_lut->broadcast_lut(streams, gpu_indexes, 0);
     }
   }
 

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

@@ -45,6 +45,9 @@ void cuda_synchronize_stream(cudaStream_t stream, uint32_t gpu_index) {
   check_cuda_error(cudaStreamSynchronize(stream));
 }
 
+// Determine if a CUDA device is available at runtime
+uint32_t cuda_is_available() { return cudaSetDevice(0) == cudaSuccess; }
+
 /// Unsafe function that will try to allocate even if gpu_index is invalid
 /// or if there's not enough memory. A safe wrapper around it must call
 /// cuda_check_valid_malloc() first

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

@@ -125,11 +125,11 @@ __host__ void are_all_comparisons_block_true(
           return x == chunk_length;
         };
         generate_device_accumulator<Torus>(
-            streams[0], gpu_indexes[0], new_lut->get_lut(gpu_indexes[0], 0),
-            glwe_dimension, polynomial_size, message_modulus, carry_modulus,
+            streams[0], gpu_indexes[0], new_lut->get_lut(0, 0), glwe_dimension,
+            polynomial_size, message_modulus, carry_modulus,
             is_equal_to_num_blocks_lut_f);
 
-        new_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        new_lut->broadcast_lut(streams, gpu_indexes, 0);
 
         (*is_equal_to_num_blocks_map)[chunk_length] = new_lut;
         lut = new_lut;
@@ -449,9 +449,9 @@ __host__ void tree_sign_reduction(
     f = sign_handler_f;
   }
   generate_device_accumulator<Torus>(
-      streams[0], gpu_indexes[0], last_lut->get_lut(gpu_indexes[0], 0),
-      glwe_dimension, polynomial_size, message_modulus, carry_modulus, f);
-  last_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+      streams[0], gpu_indexes[0], last_lut->get_lut(0, 0), glwe_dimension,
+      polynomial_size, message_modulus, carry_modulus, f);
+  last_lut->broadcast_lut(streams, gpu_indexes, 0);
 
   // Last leaf
   integer_radix_apply_univariate_lookup_table_kb<Torus>(

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

@@ -1463,10 +1463,10 @@ reduce_signs(cudaStream_t const *streams, uint32_t const *gpu_indexes,
   if (num_sign_blocks > 2) {
     auto lut = diff_buffer->reduce_signs_lut;
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
+        streams[0], gpu_indexes[0], lut->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus,
         reduce_two_orderings_function);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     while (num_sign_blocks > 2) {
       pack_blocks<Torus>(streams[0], gpu_indexes[0], signs_b, signs_a,
@@ -1497,10 +1497,9 @@ reduce_signs(cudaStream_t const *streams, uint32_t const *gpu_indexes,
 
     auto lut = diff_buffer->reduce_signs_lut;
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
-        final_lut_f);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        streams[0], gpu_indexes[0], lut->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus, final_lut_f);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     pack_blocks<Torus>(streams[0], gpu_indexes[0], signs_b, signs_a,
                        big_lwe_dimension, 2, 4);
@@ -1517,10 +1516,9 @@ reduce_signs(cudaStream_t const *streams, uint32_t const *gpu_indexes,
 
     auto lut = mem_ptr->diff_buffer->reduce_signs_lut;
     generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
-        final_lut_f);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        streams[0], gpu_indexes[0], lut->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus, final_lut_f);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     integer_radix_apply_univariate_lookup_table_kb<Torus>(
         streams, gpu_indexes, gpu_count, signs_array_out, signs_a, bsks, ksks,
@@ -1539,11 +1537,11 @@ void scratch_cuda_apply_univariate_lut_kb(
                                       1, num_radix_blocks, allocate_gpu_memory);
   // It is safe to do this copy on GPU 0, because all LUTs always reside on GPU
   // 0
-  cuda_memcpy_async_to_gpu(
-      (*mem_ptr)->get_lut(gpu_indexes[0], 0), (void *)input_lut,
-      (params.glwe_dimension + 1) * params.polynomial_size * sizeof(Torus),
-      streams[0], gpu_indexes[0]);
-  (*mem_ptr)->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+  cuda_memcpy_async_to_gpu((*mem_ptr)->get_lut(0, 0), (void *)input_lut,
+                           (params.glwe_dimension + 1) *
+                               params.polynomial_size * sizeof(Torus),
+                           streams[0], gpu_indexes[0]);
+  (*mem_ptr)->broadcast_lut(streams, gpu_indexes, 0);
 }
 
 template <typename Torus>
@@ -1582,11 +1580,11 @@ void scratch_cuda_apply_bivariate_lut_kb(
                                       1, num_radix_blocks, allocate_gpu_memory);
   // It is safe to do this copy on GPU 0, because all LUTs always reside on GPU
   // 0
-  cuda_memcpy_async_to_gpu(
-      (*mem_ptr)->get_lut(gpu_indexes[0], 0), (void *)input_lut,
-      (params.glwe_dimension + 1) * params.polynomial_size * sizeof(Torus),
-      streams[0], gpu_indexes[0]);
-  (*mem_ptr)->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+  cuda_memcpy_async_to_gpu((*mem_ptr)->get_lut(0, 0), (void *)input_lut,
+                           (params.glwe_dimension + 1) *
+                               params.polynomial_size * sizeof(Torus),
+                           streams[0], gpu_indexes[0]);
+  (*mem_ptr)->broadcast_lut(streams, gpu_indexes, 0);
 }
 
 template <typename Torus>

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

@@ -267,8 +267,8 @@ __host__ void host_integer_partial_sum_ciphertexts_vec_kb(
         streams, gpu_indexes, gpu_count, mem_ptr->params, 2,
         2 * ch_amount * num_blocks, reused_lut);
   }
-  auto message_acc = luts_message_carry->get_lut(gpu_indexes[0], 0);
-  auto carry_acc = luts_message_carry->get_lut(gpu_indexes[0], 1);
+  auto message_acc = luts_message_carry->get_lut(0, 0);
+  auto carry_acc = luts_message_carry->get_lut(0, 1);
 
   // define functions for each accumulator
   auto lut_f_message = [message_modulus](Torus x) -> Torus {
@@ -285,7 +285,7 @@ __host__ void host_integer_partial_sum_ciphertexts_vec_kb(
   generate_device_accumulator<Torus>(
       streams[0], gpu_indexes[0], carry_acc, glwe_dimension, polynomial_size,
       message_modulus, carry_modulus, lut_f_carry);
-  luts_message_carry->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+  luts_message_carry->broadcast_lut(streams, gpu_indexes, 0);
 
   while (r > 2) {
     size_t cur_total_blocks = r * num_blocks;
@@ -334,10 +334,10 @@ __host__ void host_integer_partial_sum_ciphertexts_vec_kb(
     if (carry_count > 0)
       cuda_set_value_async<Torus>(
           streams[0], gpu_indexes[0],
-          luts_message_carry->get_lut_indexes(gpu_indexes[0], message_count), 1,
+          luts_message_carry->get_lut_indexes(0, message_count), 1,
           carry_count);
 
-    luts_message_carry->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    luts_message_carry->broadcast_lut(streams, gpu_indexes, 0);
 
     /// For multi GPU execution we create vectors of pointers for inputs and
     /// outputs
@@ -579,7 +579,7 @@ __host__ void host_integer_mult_radix_kb(
       2 * num_blocks, mem_ptr->luts_array);
 
   uint32_t block_modulus = message_modulus * carry_modulus;
-  uint32_t num_bits_in_block = std::log2(block_modulus);
+  uint32_t num_bits_in_block = log2_int(block_modulus);
 
   auto scp_mem_ptr = mem_ptr->sc_prop_mem;
   uint32_t requested_flag = outputFlag::FLAG_NONE;

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

@@ -129,7 +129,7 @@ __host__ void host_integer_overflowing_sub(
   // of num_blocks changes
   uint32_t block_modulus =
       radix_params.message_modulus * radix_params.carry_modulus;
-  uint32_t num_bits_in_block = std::log2(block_modulus);
+  uint32_t num_bits_in_block = log2_int(block_modulus);
   uint32_t grouping_size = num_bits_in_block;
   uint32_t num_groups = (num_blocks + grouping_size - 1) / grouping_size;
 

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

@@ -31,10 +31,10 @@ __host__ void host_integer_radix_scalar_bitop_kb(
   } else {
     // We have all possible LUTs pre-computed and we use the decomposed scalar
     // as index to recover the right one
-    cuda_memcpy_async_gpu_to_gpu(lut->get_lut_indexes(gpu_indexes[0], 0),
-                                 clear_blocks, num_clear_blocks * sizeof(Torus),
-                                 streams[0], gpu_indexes[0]);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    cuda_memcpy_async_gpu_to_gpu(lut->get_lut_indexes(0, 0), clear_blocks,
+                                 num_clear_blocks * sizeof(Torus), streams[0],
+                                 gpu_indexes[0]);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     integer_radix_apply_univariate_lookup_table_kb<Torus>(
         streams, gpu_indexes, gpu_count, lwe_array_out, lwe_array_input, bsks,

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

@@ -110,11 +110,11 @@ __host__ void integer_radix_unsigned_scalar_difference_check_kb(
     };
 
     auto lut = mem_ptr->diff_buffer->tree_buffer->tree_last_leaf_scalar_lut;
-    generate_device_accumulator<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
-        scalar_last_leaf_lut_f);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    generate_device_accumulator<Torus>(streams[0], gpu_indexes[0],
+                                       lut->get_lut(0, 0), glwe_dimension,
+                                       polynomial_size, message_modulus,
+                                       carry_modulus, scalar_last_leaf_lut_f);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     integer_radix_apply_univariate_lookup_table_kb<Torus>(
         streams, gpu_indexes, gpu_count, lwe_array_out,
@@ -194,10 +194,10 @@ __host__ void integer_radix_unsigned_scalar_difference_check_kb(
 
     auto lut = diff_buffer->tree_buffer->tree_last_leaf_scalar_lut;
     generate_device_accumulator_bivariate<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
+        streams[0], gpu_indexes[0], lut->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus,
         scalar_bivariate_last_leaf_lut_f);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     integer_radix_apply_bivariate_lookup_table_kb<Torus>(
         streams, gpu_indexes, gpu_count, lwe_array_out, lwe_array_lsb_out,
@@ -292,7 +292,7 @@ __host__ void integer_radix_signed_scalar_difference_check_kb(
     Torus const *sign_block =
         lwe_array_in + (total_num_radix_blocks - 1) * big_lwe_size;
 
-    auto sign_bit_pos = (int)std::log2(message_modulus) - 1;
+    auto sign_bit_pos = (int)log2_int(message_modulus) - 1;
 
     auto scalar_last_leaf_with_respect_to_zero_lut_f =
         [sign_handler_f, sign_bit_pos,
@@ -329,10 +329,10 @@ __host__ void integer_radix_signed_scalar_difference_check_kb(
 
     auto lut = mem_ptr->diff_buffer->tree_buffer->tree_last_leaf_scalar_lut;
     generate_device_accumulator_bivariate<Torus>(
-        streams[0], gpu_indexes[0], lut->get_lut(gpu_indexes[0], 0),
-        glwe_dimension, polynomial_size, message_modulus, carry_modulus,
+        streams[0], gpu_indexes[0], lut->get_lut(0, 0), glwe_dimension,
+        polynomial_size, message_modulus, carry_modulus,
         scalar_bivariate_last_leaf_lut_f);
-    lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+    lut->broadcast_lut(streams, gpu_indexes, 0);
 
     integer_radix_apply_bivariate_lookup_table_kb<Torus>(
         streams, gpu_indexes, gpu_count, lwe_array_out, are_all_msb_zeros,
@@ -422,11 +422,10 @@ __host__ void integer_radix_signed_scalar_difference_check_kb(
 
     auto signed_msb_lut = mem_ptr->signed_msb_lut;
     generate_device_accumulator_bivariate<Torus>(
-        msb_streams[0], gpu_indexes[0],
-        signed_msb_lut->get_lut(gpu_indexes[0], 0), params.glwe_dimension,
-        params.polynomial_size, params.message_modulus, params.carry_modulus,
-        lut_f);
-    signed_msb_lut->broadcast_lut(streams, gpu_indexes, gpu_indexes[0]);
+        msb_streams[0], gpu_indexes[0], signed_msb_lut->get_lut(0, 0),
+        params.glwe_dimension, params.polynomial_size, params.message_modulus,
+        params.carry_modulus, lut_f);
+    signed_msb_lut->broadcast_lut(streams, gpu_indexes, 0);
 
     Torus const *sign_block = msb + (num_msb_radix_blocks - 1) * big_lwe_size;
     integer_radix_apply_bivariate_lookup_table_kb<Torus>(
@@ -676,10 +675,10 @@ __host__ void host_integer_radix_scalar_equality_check_kb(
     pack_blocks<Torus>(lsb_streams[0], gpu_indexes[0], packed_scalar,
                        scalar_blocks, 0, num_scalar_blocks, message_modulus);
 
-    cuda_memcpy_async_gpu_to_gpu(
-        scalar_comparison_luts->get_lut_indexes(gpu_indexes[0], 0),
-        packed_scalar, num_halved_scalar_blocks * sizeof(Torus), lsb_streams[0],
-        gpu_indexes[0]);
+    cuda_memcpy_async_gpu_to_gpu(scalar_comparison_luts->get_lut_indexes(0, 0),
+                                 packed_scalar,
+                                 num_halved_scalar_blocks * sizeof(Torus),
+                                 lsb_streams[0], gpu_indexes[0]);
     scalar_comparison_luts->broadcast_lut(lsb_streams, gpu_indexes, 0);
 
     integer_radix_apply_univariate_lookup_table_kb<Torus>(

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

@@ -54,7 +54,7 @@ __host__ void host_integer_scalar_mul_radix(
   // whereas lwe_dimension is the number of elements in the mask
   uint32_t lwe_size = input_lwe_dimension + 1;
   uint32_t lwe_size_bytes = lwe_size * sizeof(T);
-  uint32_t msg_bits = (uint32_t)std::log2(message_modulus);
+  uint32_t msg_bits = log2_int(message_modulus);
   uint32_t num_ciphertext_bits = msg_bits * num_radix_blocks;
 
   T *preshifted_buffer = mem->preshifted_buffer;

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

@@ -38,7 +38,7 @@ __host__ void host_integer_radix_scalar_rotate_kb_inplace(
   size_t big_lwe_size = glwe_dimension * polynomial_size + 1;
   size_t big_lwe_size_bytes = big_lwe_size * sizeof(Torus);
 
-  size_t num_bits_in_message = (size_t)log2(message_modulus);
+  size_t num_bits_in_message = (size_t)log2_int(message_modulus);
   size_t total_num_bits = num_bits_in_message * num_blocks;
   n = n % total_num_bits;
 

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

@@ -38,7 +38,7 @@ __host__ void host_integer_radix_logical_scalar_shift_kb_inplace(
   size_t big_lwe_size = glwe_dimension * polynomial_size + 1;
   size_t big_lwe_size_bytes = big_lwe_size * sizeof(Torus);
 
-  size_t num_bits_in_block = (size_t)log2(message_modulus);
+  size_t num_bits_in_block = (size_t)log2_int(message_modulus);
   size_t total_num_bits = num_bits_in_block * num_blocks;
   shift = shift % total_num_bits;
 
@@ -141,7 +141,7 @@ __host__ void host_integer_radix_arithmetic_scalar_shift_kb_inplace(
   size_t big_lwe_size = glwe_dimension * polynomial_size + 1;
   size_t big_lwe_size_bytes = big_lwe_size * sizeof(Torus);
 
-  size_t num_bits_in_block = (size_t)log2(message_modulus);
+  size_t num_bits_in_block = (size_t)log2_int(message_modulus);
   size_t total_num_bits = num_bits_in_block * num_blocks;
   shift = shift % total_num_bits;
 

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

@@ -29,7 +29,7 @@ __host__ void host_integer_radix_shift_and_rotate_kb_inplace(
     uint32_t gpu_count, Torus *lwe_array, Torus const *lwe_shift,
     int_shift_and_rotate_buffer<Torus> *mem, void *const *bsks,
     Torus *const *ksks, uint32_t num_radix_blocks) {
-  uint32_t bits_per_block = std::log2(mem->params.message_modulus);
+  uint32_t bits_per_block = log2_int(mem->params.message_modulus);
   uint32_t total_nb_bits = bits_per_block * num_radix_blocks;
   if (total_nb_bits == 0)
     return;
@@ -55,7 +55,7 @@ __host__ void host_integer_radix_shift_and_rotate_kb_inplace(
   // then the behaviour of shifting won't be the same
   // if shift >= total_nb_bits compared to when total_nb_bits
   // is a power of two, as will 'capture' more bits in `shift_bits`
-  uint32_t max_num_bits_that_tell_shift = std::log2(total_nb_bits);
+  uint32_t max_num_bits_that_tell_shift = log2_int(total_nb_bits);
   if (!is_power_of_two(total_nb_bits))
     max_num_bits_that_tell_shift += 1;
   // Extracts bits and put them in the bit index 2 (=> bit number 3)

backends/tfhe-cuda-backend/cuda/src/pbs/programmable_bootstrap_classic.cu

@@ -656,6 +656,8 @@ void cuda_programmable_bootstrap_lwe_ciphertext_vector_64(
     uint32_t num_samples, uint32_t num_many_lut, uint32_t lut_stride) {
   if (base_log > 64)
     PANIC("Cuda error (classical PBS): base log should be <= 64")
+  if ((glwe_dimension + 1) * level_count > 8)
+    PANIC("Cuda error (multi-bit PBS): (k + 1)*l should be <= 8")
 
   pbs_buffer<uint64_t, CLASSICAL> *buffer =
       (pbs_buffer<uint64_t, CLASSICAL> *)mem_ptr;

backends/tfhe-cuda-backend/cuda/src/pbs/programmable_bootstrap_multibit.cu

@@ -220,6 +220,8 @@ void cuda_multi_bit_programmable_bootstrap_lwe_ciphertext_vector_64(
 
   if (base_log > 64)
     PANIC("Cuda error (multi-bit PBS): base log should be <= 64")
+  if ((glwe_dimension + 1) * level_count > 8)
+    PANIC("Cuda error (multi-bit PBS): (k + 1)*l should be <= 8")
 
   pbs_buffer<uint64_t, MULTI_BIT> *buffer =
       (pbs_buffer<uint64_t, MULTI_BIT> *)mem_ptr;
@@ -465,7 +467,7 @@ uint32_t get_lwe_chunk_size(uint32_t gpu_index, uint32_t max_num_pbs,
 #if CUDA_ARCH < 900
   // We pick a smaller divisor on GPUs other than H100, so 256-bit integer
   // multiplication can run
-  int log2_max_num_pbs = std::log2(max_num_pbs);
+  int log2_max_num_pbs = log2_int(max_num_pbs);
   if (log2_max_num_pbs > 13)
     ith_divisor = log2_max_num_pbs - 11;
 #endif

backends/tfhe-cuda-backend/cuda/src/polynomial/parameters.cuh

@@ -1,7 +1,24 @@
 #ifndef CUDA_PARAMETERS_CUH
 #define CUDA_PARAMETERS_CUH
 
-constexpr int log2(int n) { return (n <= 2) ? 1 : 1 + log2(n / 2); }
+#include "device.h"
+#include <cstdint>
+
+// If decide to support something else than 32 and 64 bits, this method will
+// need to be adjusted
+template <typename T> constexpr unsigned log2_int(T n) {
+  if (n == 0) {
+    PANIC("Cuda error (log2): log2 is undefined for 0");
+  }
+
+  if constexpr (sizeof(T) == 4) { // uint32_t
+    return (unsigned)(8 * sizeof(uint32_t) - __builtin_clz(n) - 1);
+  } else if constexpr (sizeof(T) == 8) { // uint64_t
+    return (unsigned)(8 * sizeof(uint64_t) - __builtin_clzll(n) - 1);
+  } else {
+    return (n <= 2) ? 1 : 1 + log2_int(n / 2);
+  }
+}
 
 constexpr int choose_opt_amortized(int degree) {
   if (degree <= 1024)
@@ -41,14 +58,14 @@ template <int N> class Degree {
 public:
   constexpr static int degree = N;
   constexpr static int opt = choose_opt(N);
-  constexpr static int log2_degree = log2(N);
+  constexpr static int log2_degree = log2_int(N);
 };
 
 template <int N> class AmortizedDegree {
 public:
   constexpr static int degree = N;
   constexpr static int opt = choose_opt_amortized(N);
-  constexpr static int log2_degree = log2(N);
+  constexpr static int log2_degree = log2_int(N);
 };
 enum sharedMemDegree { NOSM = 0, PARTIALSM = 1, FULLSM = 2 };
 

backends/tfhe-cuda-backend/src/cuda_bind.rs

@@ -9,6 +9,8 @@ extern "C" {
 
     pub fn cuda_synchronize_stream(stream: *mut c_void, gpu_index: u32);
 
+    pub fn cuda_is_available() -> u32;
+
     pub fn cuda_malloc(size: u64, gpu_index: u32) -> *mut c_void;
 
     pub fn cuda_malloc_async(size: u64, stream: *mut c_void, gpu_index: u32) -> *mut c_void;

ci/lattice_estimator.sage

@@ -91,6 +91,8 @@ if __name__ == "__main__":
         "boolean_parameters_lattice_estimator.sage",
         "shortint_classic_parameters_lattice_estimator.sage",
         "shortint_multi_bit_parameters_lattice_estimator.sage",
+        "shortint_cpke_parameters_lattice_estimator.sage",
+        "shortint_list_compression_parameters_lattice_estimator.sage",
     ):
         to_update, to_watch = check_security(params_filename)
         params_to_update.extend(to_update)

ci/webdriver.py

@@ -72,7 +72,7 @@ parser.add_argument(
 parser.add_argument(
     "--index-path",
     dest="index_path",
-    default="tfhe/web_wasm_parallel_tests/index.html",
+    default="crates/tfhe/tests/web_wasm_parallel/index.html",
     help="Path to HTML index file containing all the tests/benchmarks",
 )
 parser.add_argument(
@@ -376,7 +376,7 @@ def dump_benchmark_results(results, browser_kind):
             key.replace("mean", "_".join((browser_kind.name, "mean"))): val
             for key, val in results.items()
         }
-        pathlib.Path("tfhe/wasm_benchmark_results.json").write_text(json.dumps(results))
+        pathlib.Path("crates/tfhe/wasm_benchmark_results.json").write_text(json.dumps(results))
 
 
 def start_web_server(

crates/tfhe-core-crypto/Cargo.toml

@@ -0,0 +1,82 @@
+[package]
+name = "tfhe-core-crypto"
+version.workspace = true
+edition = "2021"
+license.workspace = true
+description = "Low level cryptographic primitives used in the TFHE-rs library."
+homepage = "https://zama.ai/"
+documentation.workspace = true
+repository.workspace = true
+keywords = ["fully", "homomorphic", "encryption", "fhe", "cryptography"]
+rust-version = "1.73"
+
+[dependencies]
+serde = { workspace = true, features = ["default", "derive"] }
+pulp = { workspace = true, features = ["default"] }
+aligned-vec = { workspace = true, features = ["default", "serde"] }
+dyn-stack = { workspace = true, features = ["default"] }
+# While we wait for repeat_n in rust standard library
+itertools = "0.11.0"
+rand_core = { version = "0.6.4", features = ["std"] }
+rayon = { version = "1.5.0" }
+bytemuck = { workspace = true }
+paste = "1.0.7"
+bincode = "1.3.3"
+fs2 = { version = "0.4.3", optional = true }
+lazy_static = { version = "1.4.0", optional = true }
+
+tfhe-csprng = { version = "0.4.1", path = "../tfhe-csprng", features = [
+	"generator_fallback",
+	"parallel",
+] }
+tfhe-fft = { version = "0.6.0", path = "../tfhe-fft", features = [
+	"serde",
+	"fft128",
+] }
+tfhe-ntt = { version = "0.3.0", path = "../tfhe-ntt" }
+tfhe-cuda-backend = { version = "0.6.0", path = "../../backends/tfhe-cuda-backend", optional = true }
+tfhe-zk-pok = { version = "0.3.1", path = "../tfhe-zk-pok", optional = true }
+tfhe-versionable = { version = "0.3.2", path = "../tfhe-versionable" }
+tfhe-safe-serialization = { version = "0.11.0", path = "../tfhe-safe-serialization" }
+
+getrandom = { version = "0.2.8", optional = true }
+
+[dev-dependencies]
+rand = "0.8.5"
+# For erf and normality test
+libm = "0.2.6"
+rand_distr = "0.4.3"
+
+[lints]
+workspace = true
+
+[features]
+nightly-avx512 = ["tfhe-fft/nightly", "tfhe-ntt/nightly", "pulp/nightly"]
+gpu = ["dep:tfhe-cuda-backend"]
+zk-pok = ["dep:tfhe-zk-pok"]
+
+internal-keycache = ["dep:lazy_static", "dep:fs2"]
+
+# Experimental section
+experimental = []
+experimental-force_fft_algo_dif4 = []
+
+# Private features
+__profiling = []
+
+__c_api = []
+__wasm_api = ["dep:getrandom", "getrandom/js"]
+
+# Make some internal mut getters pub for testing purpose
+__test_core_getters = []
+
+# Enable the x86_64 specific accelerated implementation of the random generator for the default
+# backend
+generator_x86_64_aesni = ["tfhe-csprng/generator_x86_64_aesni"]
+
+# Enable the aarch64 specific accelerated implementation of the random generator for the default
+# backend
+generator_aarch64_aes = ["tfhe-csprng/generator_aarch64_aes"]
+
+seeder_unix = ["tfhe-csprng/seeder_unix"]
+seeder_x86_64_rdseed = ["tfhe-csprng/seeder_x86_64_rdseed"]

crates/tfhe-core-crypto/src/algorithms/ggsw_conversion.rs

@@ -2,11 +2,11 @@
 //! [`standard GGSW ciphertexts`](`GgswCiphertext`) to various representations/numerical domains
 //! like the Fourier domain.
 
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft64::crypto::ggsw::fill_with_forward_fourier_scratch;
-use crate::core_crypto::fft_impl::fft64::math::fft::{Fft, FftView};
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft64::crypto::ggsw::fill_with_forward_fourier_scratch;
+use crate::fft_impl::fft64::math::fft::{Fft, FftView};
 use dyn_stack::{PodStack, SizeOverflow, StackReq};
 use tfhe_fft::c64;
 
@@ -43,7 +43,7 @@ pub fn convert_standard_ggsw_ciphertext_to_fourier<Scalar, InputCont, OutputCont
 
 /// Memory optimized version of [`convert_standard_ggsw_ciphertext_to_fourier`].
 ///
-/// See [`cmux_assign_mem_optimized`](`crate::core_crypto::algorithms::cmux_assign_mem_optimized`)
+/// See [`cmux_assign_mem_optimized`](`crate::algorithms::cmux_assign_mem_optimized`)
 /// for usage.
 pub fn convert_standard_ggsw_ciphertext_to_fourier_mem_optimized<Scalar, InputCont, OutputCont>(
     input_ggsw: &GgswCiphertext<InputCont>,

crates/tfhe-core-crypto/src/algorithms/ggsw_encryption.rs

@@ -1,18 +1,18 @@
 //! Module containing primitives pertaining to [`GGSW ciphertext
 //! encryption`](`GgswCiphertext#ggsw-encryption`).
 
-use crate::core_crypto::algorithms::misc::divide_round;
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::ciphertext_modulus::{CiphertextModulus, CiphertextModulusKind};
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::decomposition::{
+use crate::algorithms::misc::divide_round;
+use crate::algorithms::slice_algorithms::*;
+use crate::algorithms::*;
+use crate::commons::ciphertext_modulus::{CiphertextModulus, CiphertextModulusKind};
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::decomposition::{
     DecompositionLevel, DecompositionTerm, DecompositionTermNonNative, SignedDecomposer,
 };
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::{DecompositionBaseLog, PlaintextCount};
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::{DecompositionBaseLog, PlaintextCount};
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Compute the multiplicative factor for a GGSW encryption based on an input value and GGSW

crates/tfhe-core-crypto/src/algorithms/glwe_encryption.rs

@@ -1,16 +1,16 @@
 //! Module containing primitives pertaining to [`GLWE ciphertext
 //! encryption`](`GlweCiphertext#glwe-encryption`).
 
-use crate::core_crypto::algorithms::polynomial_algorithms::*;
-use crate::core_crypto::algorithms::slice_algorithms::{
+use crate::algorithms::polynomial_algorithms::*;
+use crate::algorithms::slice_algorithms::{
     slice_wrapping_scalar_div_assign, slice_wrapping_scalar_mul_assign,
 };
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulusKind;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::ciphertext_modulus::CiphertextModulusKind;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Convenience function to share the core logic of the GLWE assign encryption between all functions
 /// needing it.

crates/tfhe-core-crypto/src/algorithms/glwe_linear_algebra.rs

@@ -1,9 +1,9 @@
 //! Module containing primitives pertaining to [`GLWE ciphertext`](`GlweCiphertext`) linear algebra,
 //! like addition, multiplication, etc.
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Add the right-hand side [`GLWE ciphertext`](`GlweCiphertext`) to the left-hand side [`GLWE
 /// ciphertext`](`GlweCiphertext`) updating it in-place.

crates/tfhe-core-crypto/src/algorithms/glwe_sample_extraction.rs

@@ -2,10 +2,10 @@
 //! _sample extract_ in the literature. Allowing to extract a single
 //! [`LWE Ciphertext`](`LweCiphertext`) from a given [`GLWE ciphertext`](`GlweCiphertext`).
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Extract the nth coefficient from the body of a [`GLWE Ciphertext`](`GlweCiphertext`) as an

crates/tfhe-core-crypto/src/algorithms/glwe_secret_key_generation.rs

@@ -1,11 +1,11 @@
 //! Module containing primitives pertaining to the generation of
 //! [`GLWE secret keys`](`GlweSecretKey`).
 
-use crate::core_crypto::commons::generators::SecretRandomGenerator;
-use crate::core_crypto::commons::math::random::{RandomGenerable, UniformBinary};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::generators::SecretRandomGenerator;
+use crate::commons::math::random::{RandomGenerable, UniformBinary};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Allocate a new [`GLWE secret key`](`GlweSecretKey`) and fill it with uniformly random binary
 /// coefficients.

crates/tfhe-core-crypto/src/algorithms/lwe_bootstrap_key_conversion.rs

@@ -2,13 +2,13 @@
 //! [`standard LWE bootstrap keys`](`LweBootstrapKey`) to various representations/numerical domains
 //! like the Fourier domain.
 
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::math::ntt::ntt64::Ntt64;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft128::math::fft::Fft128;
-use crate::core_crypto::fft_impl::fft64::crypto::bootstrap::fill_with_forward_fourier_scratch;
-use crate::core_crypto::fft_impl::fft64::math::fft::{Fft, FftView};
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::math::ntt::ntt64::Ntt64;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft128::math::fft::Fft128;
+use crate::fft_impl::fft64::crypto::bootstrap::fill_with_forward_fourier_scratch;
+use crate::fft_impl::fft64::math::fft::{Fft, FftView};
 use dyn_stack::{PodStack, SizeOverflow, StackReq};
 use rayon::prelude::*;
 use tfhe_fft::c64;
@@ -16,7 +16,7 @@ use tfhe_fft::c64;
 /// Convert an [`LWE bootstrap key`](`LweBootstrapKey`) with standard coefficients to the Fourier
 /// domain.
 ///
-/// See [`programmable_bootstrap_lwe_ciphertext`](`crate::core_crypto::algorithms::programmable_bootstrap_lwe_ciphertext`) for usage.
+/// See [`programmable_bootstrap_lwe_ciphertext`](`crate::algorithms::programmable_bootstrap_lwe_ciphertext`) for usage.
 pub fn convert_standard_lwe_bootstrap_key_to_fourier<Scalar, InputCont, OutputCont>(
     input_bsk: &LweBootstrapKey<InputCont>,
     output_bsk: &mut FourierLweBootstrapKey<OutputCont>,
@@ -159,7 +159,7 @@ pub fn convert_standard_lwe_bootstrap_key_to_fourier_mem_optimized_requirement(
 /// Convert an [`LWE bootstrap key`](`LweBootstrapKey`) with standard coefficients to the Fourier
 /// domain.
 ///
-/// See [`programmable_bootstrap_f128_lwe_ciphertext`](`crate::core_crypto::algorithms::programmable_bootstrap_f128_lwe_ciphertext`) for usage.
+/// See [`programmable_bootstrap_f128_lwe_ciphertext`](`crate::algorithms::programmable_bootstrap_f128_lwe_ciphertext`) for usage.
 pub fn convert_standard_lwe_bootstrap_key_to_fourier_128<Scalar, InputCont, OutputCont>(
     input_bsk: &LweBootstrapKey<InputCont>,
     output_bsk: &mut Fourier128LweBootstrapKey<OutputCont>,
@@ -215,7 +215,7 @@ pub fn convert_standard_lwe_bootstrap_key_to_fourier_128<Scalar, InputCont, Outp
 /// Convert an [`LWE bootstrap key`](`LweBootstrapKey`) with standard coefficients to the NTT
 /// domain using a 64 bits NTT.
 ///
-/// See [`programmable_bootstrap_ntt64_lwe_ciphertext_mem_optimized`](`crate::core_crypto::algorithms::programmable_bootstrap_ntt64_lwe_ciphertext_mem_optimized`) for usage.
+/// See [`programmable_bootstrap_ntt64_lwe_ciphertext_mem_optimized`](`crate::algorithms::programmable_bootstrap_ntt64_lwe_ciphertext_mem_optimized`) for usage.
 pub fn convert_standard_lwe_bootstrap_key_to_ntt64<InputCont, OutputCont>(
     input_bsk: &LweBootstrapKey<InputCont>,
     output_bsk: &mut NttLweBootstrapKey<OutputCont>,

crates/tfhe-core-crypto/src/algorithms/lwe_bootstrap_key_generation.rs

@@ -2,12 +2,12 @@
 //! [`standard LWE bootstrap keys`](`LweBootstrapKey`) and [`seeded standard LWE bootstrap
 //! keys`](`SeededLweBootstrapKey`).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Fill an [`LWE bootstrap key`](`LweBootstrapKey`) with an actual bootstrapping key constructed

crates/tfhe-core-crypto/src/algorithms/lwe_compact_ciphertext_list_expansion.rs

@@ -1,9 +1,9 @@
 //! Module with primitives pertaining to [`LweCompactCiphertextList`] expansion.
 
-use crate::core_crypto::algorithms::polynomial_algorithms::polynomial_wrapping_monic_monomial_mul_assign;
-use crate::core_crypto::commons::parameters::MonomialDegree;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::polynomial_algorithms::polynomial_wrapping_monic_monomial_mul_assign;
+use crate::commons::parameters::MonomialDegree;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Expand an [`LweCompactCiphertextList`] into an [`LweCiphertextList`].

crates/tfhe-core-crypto/src/algorithms/lwe_compact_public_key_generation.rs

@@ -1,13 +1,13 @@
 //! Module containing primitives pertaining to [`LWE compact public key
 //! generation`](`LweCompactPublicKey`).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulus;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{Distribution, Uniform};
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::prelude::ActivatedRandomGenerator;
+use crate::algorithms::*;
+use crate::commons::ciphertext_modulus::CiphertextModulus;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{Distribution, Uniform};
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::prelude::ActivatedRandomGenerator;
 use slice_algorithms::*;
 
 /// Fill an [`LWE compact public key`](`LweCompactPublicKey`) with an actual public key constructed

crates/tfhe-core-crypto/src/algorithms/lwe_encryption.rs

@@ -1,22 +1,20 @@
 //! Module containing primitives pertaining to [`LWE ciphertext encryption and
 //! decryption`](`LweCiphertext#lwe-encryption`).
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulusKind;
-use crate::core_crypto::commons::generators::{EncryptionRandomGenerator, SecretRandomGenerator};
+use crate::algorithms::slice_algorithms::*;
+use crate::algorithms::*;
+use crate::commons::ciphertext_modulus::CiphertextModulusKind;
+use crate::commons::generators::{EncryptionRandomGenerator, SecretRandomGenerator};
 #[cfg(feature = "zk-pok")]
-use crate::core_crypto::commons::math::random::BoundedDistribution;
-use crate::core_crypto::commons::math::random::{
+use crate::commons::math::random::BoundedDistribution;
+use crate::commons::math::random::{
     ActivatedRandomGenerator, Distribution, RandomGenerable, RandomGenerator, Uniform,
     UniformBinary,
 };
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
-#[cfg(feature = "zk-pok")]
-use tfhe_zk_pok::proofs::pke::{commit, prove};
 
 /// Convenience function to share the core logic of the LWE encryption between all functions needing
 /// it.
@@ -1858,8 +1856,7 @@ where
     BodyDistribution: BoundedDistribution<Scalar::Signed>,
     KeyCont: Container<Element = Scalar>,
 {
-    let public_params = crs.public_params();
-    let exclusive_max = public_params.exclusive_max_noise();
+    let exclusive_max = crs.exclusive_max_noise();
     if Scalar::BITS < 64 && (1u64 << Scalar::BITS) >= exclusive_max {
         return Err(
             "The given random distribution would create random values out \
@@ -1893,28 +1890,23 @@ where
         return Err("Zero knowledge proof do not support moduli greater than 2**64".into());
     }
 
-    let expected_q = if Scalar::BITS == 64 {
-        0u64
-    } else {
-        164 << Scalar::BITS
-    };
-
-    if expected_q != public_params.q {
+    if ciphertext_modulus != crs.ciphertext_modulus() {
         return Err("Mismatched modulus between CRS and ciphertexts".into());
     }
 
-    if ciphertext_count.0 > public_params.k {
+    if ciphertext_count > crs.max_num_messages() {
         return Err(format!(
             "CRS allows at most {} ciphertexts to be proven at once, {} contained in the list",
-            public_params.k, ciphertext_count.0
+            crs.max_num_messages().0,
+            ciphertext_count.0
         )
         .into());
     }
 
-    if lwe_compact_public_key.lwe_dimension().0 > public_params.d {
+    if lwe_compact_public_key.lwe_dimension() > crs.lwe_dimension() {
         return Err(format!(
             "CRS allows a LweDimension of at most {}, current dimension: {}",
-            public_params.d,
+            crs.lwe_dimension().0,
             lwe_compact_public_key.lwe_dimension().0
         )
         .into());
@@ -1922,10 +1914,10 @@ where
 
     // 2**64 /delta == ((2**63) / delta) *2
     let plaintext_modulus = ((1u64 << (u64::BITS - 1) as usize) / u64::cast_from(delta)) * 2;
-    if plaintext_modulus != public_params.t {
+    if plaintext_modulus != crs.plaintext_modulus() {
         return Err(format!(
             "Mismatched plaintext modulus: CRS expects {}, requested modulus: {plaintext_modulus:?}",
-            public_params.t
+            crs.plaintext_modulus()
         ).into());
     }
 
@@ -2291,52 +2283,18 @@ where
         encryption_generator,
     );
 
-    let (c1, c2) = output.get_mask_and_body();
-
-    let (public_commit, private_commit) = commit(
-        lwe_compact_public_key
-            .get_mask()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        lwe_compact_public_key
-            .get_body()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        c1.as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        vec![i64::cast_from(*c2.data)],
-        binary_random_vector
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        mask_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        vec![i64::cast_from(message.0)],
-        body_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        crs.public_params(),
-        random_generator,
-    );
-
-    Ok(prove(
-        (crs.public_params(), &public_commit),
-        &private_commit,
+    Ok(crs.prove(
+        lwe_compact_public_key,
+        &vec![message.0],
+        &LweCompactCiphertextList::from_container(
+            output.as_ref(),
+            output.lwe_size(),
+            LweCiphertextCount(1),
+            output.ciphertext_modulus(),
+        ),
+        &binary_random_vector,
+        &mask_noise,
+        &body_noise,
         metadata,
         load,
         random_generator,
@@ -2807,61 +2765,13 @@ where
         encryption_generator,
     );
 
-    let (c1, c2) = output.get_mask_and_body_list();
-
-    let (public_commit, private_commit) = commit(
-        lwe_compact_public_key
-            .get_mask()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        lwe_compact_public_key
-            .get_body()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        c1.as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        c2.as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        binary_random_vector
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        mask_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        messages
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        body_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        crs.public_params(),
-        random_generator,
-    );
-
-    Ok(prove(
-        (crs.public_params(), &public_commit),
-        &private_commit,
+    Ok(crs.prove(
+        lwe_compact_public_key,
+        messages,
+        output,
+        &binary_random_vector,
+        &mask_noise,
+        &body_noise,
         metadata,
         load,
         random_generator,
@@ -3341,61 +3251,13 @@ where
         encryption_generator,
     );
 
-    let (c1, c2) = output.get_mask_and_body_list();
-
-    let (public_commit, private_commit) = commit(
-        lwe_compact_public_key
-            .get_mask()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        lwe_compact_public_key
-            .get_body()
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        c1.as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        c2.as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        binary_random_vector
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        mask_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        messages
-            .as_ref()
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        body_noise
-            .iter()
-            .copied()
-            .map(CastFrom::cast_from)
-            .collect::<Vec<_>>(),
-        crs.public_params(),
-        random_generator,
-    );
-
-    Ok(prove(
-        (crs.public_params(), &public_commit),
-        &private_commit,
+    Ok(crs.prove(
+        lwe_compact_public_key,
+        messages,
+        output,
+        &binary_random_vector,
+        &mask_noise,
+        &body_noise,
         metadata,
         load,
         random_generator,
@@ -3404,9 +3266,9 @@ where
 
 #[cfg(test)]
 mod test {
-    use crate::core_crypto::commons::generators::DeterministicSeeder;
-    use crate::core_crypto::commons::test_tools;
-    use crate::core_crypto::prelude::*;
+    use crate::commons::generators::DeterministicSeeder;
+    use crate::commons::test_tools;
+    use crate::prelude::*;
 
     #[test]
     fn test_compact_public_key_encryption() {

crates/tfhe-core-crypto/src/algorithms/lwe_keyswitch.rs

@@ -1,16 +1,12 @@
 //! Module containing primitives pertaining to [`LWE ciphertext
 //! keyswitch`](`LweKeyswitchKey#lwe-keyswitch`).
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulusKind;
-use crate::core_crypto::commons::math::decomposition::{
-    SignedDecomposer, SignedDecomposerNonNative,
-};
-use crate::core_crypto::commons::parameters::{
-    DecompositionBaseLog, DecompositionLevelCount, ThreadCount,
-};
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::ciphertext_modulus::CiphertextModulusKind;
+use crate::commons::math::decomposition::{SignedDecomposer, SignedDecomposerNonNative};
+use crate::commons::parameters::{DecompositionBaseLog, DecompositionLevelCount, ThreadCount};
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Keyswitch an [`LWE ciphertext`](`LweCiphertext`) encrypted under an
@@ -326,8 +322,8 @@ pub fn keyswitch_lwe_ciphertext_other_mod<Scalar, KSKCont, InputCont, OutputCont
 /// `input_bits` to a a smaller OutputScalar with `output_bits` and `output_bits` < `input_bits`.
 ///
 /// The product of the `lwe_keyswitch_key`'s
-/// [`DecompositionBaseLog`](`crate::core_crypto::commons::parameters::DecompositionBaseLog`) and
-/// [`DecompositionLevelCount`](`crate::core_crypto::commons::parameters::DecompositionLevelCount`)
+/// [`DecompositionBaseLog`](`crate::commons::parameters::DecompositionBaseLog`) and
+/// [`DecompositionLevelCount`](`crate::commons::parameters::DecompositionLevelCount`)
 /// needs to be smaller than `output_bits`.
 pub fn keyswitch_lwe_ciphertext_with_scalar_change<
     InputScalar,

crates/tfhe-core-crypto/src/algorithms/lwe_keyswitch_key_generation.rs

@@ -2,15 +2,15 @@
 //! generation`](`LweKeyswitchKey#key-switching-key`) and [`seeded LWE keyswitch keys
 //! generation`](`SeededLweKeyswitchKey`).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::decomposition::{
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::decomposition::{
     DecompositionLevel, DecompositionTerm, DecompositionTermNonNative,
 };
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Fill an [`LWE keyswitch key`](`LweKeyswitchKey`) with an actual keyswitching key constructed
 /// from an input and an output key [`LWE secret key`](`LweSecretKey`).

crates/tfhe-core-crypto/src/algorithms/lwe_linear_algebra.rs

@@ -1,10 +1,10 @@
 //! Module containing primitives pertaining to [`LWE ciphertext`](`LweCiphertext`) linear algebra,
 //! like addition, multiplication, etc.
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulusKind;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::ciphertext_modulus::CiphertextModulusKind;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Add the right-hand side [`LWE ciphertext`](`LweCiphertext`) to the left-hand side [`LWE
 /// ciphertext`](`LweCiphertext`) updating it in-place.

crates/tfhe-core-crypto/src/algorithms/lwe_multi_bit_bootstrap_key_conversion.rs

@@ -2,19 +2,17 @@
 //! [`standard LWE multi_bit bootstrap keys`](`LweMultiBitBootstrapKey`) to various
 //! representations/numerical domains like the Fourier domain.
 
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft64::math::fft::{
-    par_convert_polynomials_list_to_fourier, Fft, FftView,
-};
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft64::math::fft::{par_convert_polynomials_list_to_fourier, Fft, FftView};
 use dyn_stack::{PodStack, ReborrowMut, SizeOverflow, StackReq};
 use tfhe_fft::c64;
 
 /// Convert an [`LWE multi_bit bootstrap key`](`LweMultiBitBootstrapKey`) with standard
 /// coefficients to the Fourier domain.
 ///
-/// See [`multi_bit_programmable_bootstrap_lwe_ciphertext`](`crate::core_crypto::algorithms::multi_bit_programmable_bootstrap_lwe_ciphertext`) for usage.
+/// See [`multi_bit_programmable_bootstrap_lwe_ciphertext`](`crate::algorithms::multi_bit_programmable_bootstrap_lwe_ciphertext`) for usage.
 pub fn convert_standard_lwe_multi_bit_bootstrap_key_to_fourier<Scalar, InputCont, OutputCont>(
     input_bsk: &LweMultiBitBootstrapKey<InputCont>,
     output_bsk: &mut FourierLweMultiBitBootstrapKey<OutputCont>,

crates/tfhe-core-crypto/src/algorithms/lwe_multi_bit_bootstrap_key_generation.rs

@@ -1,12 +1,12 @@
 //! Module containing primitives pertaining to the generation of
 //! [`standard LWE multi_bit bootstrap keys`](`LweMultiBitBootstrapKey`).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// ```rust

crates/tfhe-core-crypto/src/algorithms/lwe_multi_bit_programmable_bootstrapping.rs

@@ -1,16 +1,16 @@
-use crate::core_crypto::algorithms::extract_lwe_sample_from_glwe_ciphertext;
-use crate::core_crypto::algorithms::polynomial_algorithms::*;
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::math::decomposition::SignedDecomposer;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::common::modulus_switch;
-use crate::core_crypto::fft_impl::fft64::crypto::ggsw::{
+use crate::algorithms::extract_lwe_sample_from_glwe_ciphertext;
+use crate::algorithms::polynomial_algorithms::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::math::decomposition::SignedDecomposer;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::common::modulus_switch;
+use crate::fft_impl::fft64::crypto::ggsw::{
     add_external_product_assign, add_external_product_assign_scratch, update_with_fmadd_factor,
 };
-use crate::core_crypto::fft_impl::fft64::math::fft::{Fft, FftView};
+use crate::fft_impl::fft64::math::fft::{Fft, FftView};
 use aligned_vec::ABox;
 use itertools::Itertools;
 use std::sync::atomic::{AtomicUsize, Ordering};

crates/tfhe-core-crypto/src/algorithms/lwe_packing_keyswitch.rs

@@ -1,13 +1,11 @@
-use crate::core_crypto::algorithms::polynomial_algorithms::polynomial_wrapping_monic_monomial_mul_assign;
-use crate::core_crypto::algorithms::slice_algorithms::{
+use crate::algorithms::polynomial_algorithms::polynomial_wrapping_monic_monomial_mul_assign;
+use crate::algorithms::slice_algorithms::{
     slice_wrapping_add_assign, slice_wrapping_sub_scalar_mul_assign,
 };
-use crate::core_crypto::commons::math::decomposition::SignedDecomposer;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::{
-    GlweCiphertext, LweCiphertext, LweCiphertextList, LwePackingKeyswitchKey,
-};
+use crate::commons::math::decomposition::SignedDecomposer;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::{GlweCiphertext, LweCiphertext, LweCiphertextList, LwePackingKeyswitchKey};
 use rayon::prelude::*;
 
 /// Apply a keyswitch on an input [`LWE ciphertext`](`LweCiphertext`) and

crates/tfhe-core-crypto/src/algorithms/lwe_packing_keyswitch_key_generation.rs

@@ -2,15 +2,15 @@
 //! generation`](`LwePackingKeyswitchKey`) and [`seeded LWE packing keyswitch keys
 //! generation`](`SeededLwePackingKeyswitchKey`).
 
-use crate::core_crypto::algorithms::{
+use crate::algorithms::{
     encrypt_glwe_ciphertext_list, encrypt_seeded_glwe_ciphertext_list_with_existing_generator,
 };
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::decomposition::{DecompositionLevel, DecompositionTerm};
-use crate::core_crypto::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::{
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::decomposition::{DecompositionLevel, DecompositionTerm};
+use crate::commons::math::random::{ActivatedRandomGenerator, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::{
     GlweSecretKey, LwePackingKeyswitchKey, LwePackingKeyswitchKeyOwned, LweSecretKey,
     PlaintextListOwned, SeededLwePackingKeyswitchKey, SeededLwePackingKeyswitchKeyOwned,
 };

crates/tfhe-core-crypto/src/algorithms/lwe_private_functional_packing_keyswitch.rs

@@ -7,12 +7,12 @@
 //!      J. Cryptol 33, 34–91 (2020). \
 //! &nbsp;&nbsp;&nbsp;&nbsp; <https://doi.org/10.1007/s00145-019-09319-x>
 
-use crate::core_crypto::algorithms::polynomial_algorithms::*;
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::math::decomposition::SignedDecomposer;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::polynomial_algorithms::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::math::decomposition::SignedDecomposer;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Apply a private functional keyswitch on an input [`LWE ciphertext`](`LweCiphertext`) and write

crates/tfhe-core-crypto/src/algorithms/lwe_private_functional_packing_keyswitch_key_generation.rs

@@ -1,14 +1,14 @@
 //! Module containing primitives pertaining to [`LWE private functional packing keyswitch key
 //! generation`](`LwePrivateFunctionalPackingKeyswitchKey`).
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::decomposition::{DecompositionLevel, DecompositionTerm};
-use crate::core_crypto::commons::math::random::{Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::decomposition::{DecompositionLevel, DecompositionTerm};
+use crate::commons::math::random::{Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Fill an [`LWE private functional packing keyswitch
@@ -249,9 +249,9 @@ pub fn par_generate_lwe_private_functional_packing_keyswitch_key<
 
 #[cfg(test)]
 mod test {
-    use crate::core_crypto::commons::generators::DeterministicSeeder;
-    use crate::core_crypto::commons::math::random::Seed;
-    use crate::core_crypto::prelude::*;
+    use crate::commons::generators::DeterministicSeeder;
+    use crate::commons::math::random::Seed;
+    use crate::prelude::*;
 
     const NB_TESTS: usize = 10;
 
@@ -261,19 +261,17 @@ mod test {
             // DISCLAIMER: these toy example parameters are not guaranteed to be secure or yield
             // correct computations
             let glwe_dimension =
-                GlweDimension(crate::core_crypto::commons::test_tools::random_usize_between(5..10));
-            let polynomial_size = PolynomialSize(
-                crate::core_crypto::commons::test_tools::random_usize_between(5..10),
-            );
-            let pfpksk_level_count = DecompositionLevelCount(
-                crate::core_crypto::commons::test_tools::random_usize_between(2..5),
-            );
-            let pfpksk_base_log = DecompositionBaseLog(
-                crate::core_crypto::commons::test_tools::random_usize_between(2..5),
-            );
+                GlweDimension(crate::commons::test_tools::random_usize_between(5..10));
+            let polynomial_size =
+                PolynomialSize(crate::commons::test_tools::random_usize_between(5..10));
+            let pfpksk_level_count =
+                DecompositionLevelCount(crate::commons::test_tools::random_usize_between(2..5));
+            let pfpksk_base_log =
+                DecompositionBaseLog(crate::commons::test_tools::random_usize_between(2..5));
 
-            let common_encryption_seed =
-                Seed(crate::core_crypto::commons::test_tools::random_uint_between(0..u128::MAX));
+            let common_encryption_seed = Seed(crate::commons::test_tools::random_uint_between(
+                0..u128::MAX,
+            ));
 
             let var_small = Variance::from_variance(2f64.powf(-80.0));
             let gaussian_small = Gaussian::from_dispersion_parameter(var_small, 0.0);

crates/tfhe-core-crypto/src/algorithms/lwe_programmable_bootstrapping/fft128.rs

@@ -1,13 +1,13 @@
 //! Module containing primitives pertaining to the [`LWE programmable
-//! bootstrap`](`crate::core_crypto::entities::LweBootstrapKey#programmable-bootstrapping`) using
+//! bootstrap`](`crate::entities::LweBootstrapKey#programmable-bootstrapping`) using
 //! 128 bits FFT for polynomial multiplication.
 
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft128::crypto::bootstrap::bootstrap_scratch as bootstrap_scratch_f128;
-use crate::core_crypto::fft_impl::fft128::math::fft::{Fft128, Fft128View};
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft128::crypto::bootstrap::bootstrap_scratch as bootstrap_scratch_f128;
+use crate::fft_impl::fft128::math::fft::{Fft128, Fft128View};
 use dyn_stack::{PodStack, SizeOverflow, StackReq};
 
 /// Perform a programmable bootstrap given an input [`LWE ciphertext`](`LweCiphertext`), a

crates/tfhe-core-crypto/src/algorithms/lwe_programmable_bootstrapping/fft64.rs

@@ -1,21 +1,21 @@
 //! Module containing primitives pertaining to the [`LWE programmable
-//! bootstrap`](`crate::core_crypto::entities::LweBootstrapKey#programmable-bootstrapping`) using 64
+//! bootstrap`](`crate::entities::LweBootstrapKey#programmable-bootstrapping`) using 64
 //! bits FFT for polynomial multiplication.
 
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::math::decomposition::SignedDecomposer;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft64::crypto::bootstrap::{
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::math::decomposition::SignedDecomposer;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft64::crypto::bootstrap::{
     batch_bootstrap_scratch, blind_rotate_assign_scratch, bootstrap_scratch,
 };
-use crate::core_crypto::fft_impl::fft64::crypto::ggsw::{
+use crate::fft_impl::fft64::crypto::ggsw::{
     add_external_product_assign as impl_add_external_product_assign,
     add_external_product_assign_scratch as impl_add_external_product_assign_scratch, cmux,
     cmux_scratch,
 };
-use crate::core_crypto::fft_impl::fft64::math::fft::{Fft, FftView};
+use crate::fft_impl::fft64::math::fft::{Fft, FftView};
 use dyn_stack::{PodStack, SizeOverflow, StackReq};
 use tfhe_fft::c64;
 

crates/tfhe-core-crypto/src/algorithms/lwe_programmable_bootstrapping/mod.rs

@@ -6,10 +6,10 @@ pub use fft128::*;
 pub use fft64::*;
 pub use ntt64::*;
 
-use crate::core_crypto::algorithms::glwe_encryption::allocate_and_trivially_encrypt_new_glwe_ciphertext;
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::glwe_encryption::allocate_and_trivially_encrypt_new_glwe_ciphertext;
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Helper function to generate an accumulator for a PBS
 ///

crates/tfhe-core-crypto/src/algorithms/lwe_programmable_bootstrapping/ntt64.rs

@@ -1,22 +1,22 @@
 //! Module containing primitives pertaining to the [`LWE programmable
-//! bootstrap`](`crate::core_crypto::entities::LweBootstrapKey#programmable-bootstrapping`) using 64
+//! bootstrap`](`crate::entities::LweBootstrapKey#programmable-bootstrapping`) using 64
 //! bits NTT for polynomial multiplication.
 
-use crate::core_crypto::algorithms::glwe_sample_extraction::extract_lwe_sample_from_glwe_ciphertext;
-use crate::core_crypto::algorithms::misc::divide_round;
-use crate::core_crypto::algorithms::polynomial_algorithms::{
+use crate::algorithms::glwe_sample_extraction::extract_lwe_sample_from_glwe_ciphertext;
+use crate::algorithms::misc::divide_round;
+use crate::algorithms::polynomial_algorithms::{
     polynomial_wrapping_monic_monomial_div_assign_custom_mod,
     polynomial_wrapping_monic_monomial_mul_assign_custom_mod,
 };
-use crate::core_crypto::commons::computation_buffers::ComputationBuffers;
-use crate::core_crypto::commons::math::decomposition::{
+use crate::commons::computation_buffers::ComputationBuffers;
+use crate::commons::math::decomposition::{
     SignedDecomposerNonNative, TensorSignedDecompositionLendingIterNonNative,
 };
-use crate::core_crypto::commons::math::ntt::ntt64::{Ntt64, Ntt64View};
-use crate::core_crypto::commons::parameters::{GlweSize, MonomialDegree, PolynomialSize};
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::commons::utils::izip;
-use crate::core_crypto::entities::*;
+use crate::commons::math::ntt::ntt64::{Ntt64, Ntt64View};
+use crate::commons::parameters::{GlweSize, MonomialDegree, PolynomialSize};
+use crate::commons::traits::*;
+use crate::commons::utils::izip;
+use crate::entities::*;
 use aligned_vec::CACHELINE_ALIGN;
 use dyn_stack::{PodStack, ReborrowMut, SizeOverflow, StackReq};
 

crates/tfhe-core-crypto/src/algorithms/lwe_public_key_generation.rs

@@ -2,12 +2,12 @@
 //! generation`](`LwePublicKey#lwe-public-key`) and [`seeded LWE public key
 //! generation`](`SeededLwePublicKey#lwe-public-key`).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{CompressionSeed, Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{CompressionSeed, Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Fill an [`LWE public key`](`LwePublicKey`) with an actual public key constructed from a private
 /// [`LWE secret key`](`LweSecretKey`).

crates/tfhe-core-crypto/src/algorithms/lwe_secret_key_generation.rs

@@ -1,11 +1,11 @@
 //! Module containing primitives pertaining to the generation of
 //! [`LWE secret keys`](`LweSecretKey`).
 
-use crate::core_crypto::commons::generators::SecretRandomGenerator;
-use crate::core_crypto::commons::math::random::{RandomGenerable, UniformBinary};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::commons::generators::SecretRandomGenerator;
+use crate::commons::math::random::{RandomGenerable, UniformBinary};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Allocate a new [`LWE secret key`](`LweSecretKey`) and fill it with uniformly random binary
 /// coefficients.

crates/tfhe-core-crypto/src/algorithms/lwe_wopbs.rs

@@ -1,16 +1,16 @@
 //! Module containing primitives pertaining to the Wopbs (WithOut padding PBS).
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::EncryptionRandomGenerator;
-use crate::core_crypto::commons::math::random::{Distribution, Uniform};
-use crate::core_crypto::commons::parameters::*;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
-use crate::core_crypto::fft_impl::fft64::crypto::wop_pbs::{
+use crate::algorithms::*;
+use crate::commons::generators::EncryptionRandomGenerator;
+use crate::commons::math::random::{Distribution, Uniform};
+use crate::commons::parameters::*;
+use crate::commons::traits::*;
+use crate::entities::*;
+use crate::fft_impl::fft64::crypto::wop_pbs::{
     circuit_bootstrap_boolean_vertical_packing, circuit_bootstrap_boolean_vertical_packing_scratch,
     extract_bits, extract_bits_scratch,
 };
-use crate::core_crypto::fft_impl::fft64::math::fft::FftView;
+use crate::fft_impl::fft64::math::fft::FftView;
 use dyn_stack::{PodStack, SizeOverflow, StackReq};
 use rayon::prelude::*;
 use tfhe_fft::c64;

crates/tfhe-core-crypto/src/algorithms/lwe_zero_knowledge_verification.rs

@@ -0,0 +1,48 @@
+use crate::entities::{LweCompactCiphertextList, LweCompactPublicKey};
+use crate::prelude::{CastFrom, Container, LweCiphertext, LweCiphertextCount, UnsignedInteger};
+use crate::zk::{CompactPkeCrs, CompactPkeProof, ZkVerificationOutcome};
+
+/// Verifies with the given proof that a [`LweCompactCiphertextList`]
+/// is valid.
+pub fn verify_lwe_compact_ciphertext_list<Scalar, ListCont, KeyCont>(
+    lwe_compact_list: &LweCompactCiphertextList<ListCont>,
+    compact_public_key: &LweCompactPublicKey<KeyCont>,
+    proof: &CompactPkeProof,
+    crs: &CompactPkeCrs,
+    metadata: &[u8],
+) -> ZkVerificationOutcome
+where
+    Scalar: UnsignedInteger,
+    i64: CastFrom<Scalar>,
+    ListCont: Container<Element = Scalar>,
+    KeyCont: Container<Element = Scalar>,
+{
+    crs.verify(lwe_compact_list, compact_public_key, proof, metadata)
+}
+
+/// Verifies with the given proof that a single [`LweCiphertext`] is valid.
+pub fn verify_lwe_ciphertext<Scalar, Cont, KeyCont>(
+    lwe_ciphertext: &LweCiphertext<Cont>,
+    compact_public_key: &LweCompactPublicKey<KeyCont>,
+    proof: &CompactPkeProof,
+    crs: &CompactPkeCrs,
+    metadata: &[u8],
+) -> ZkVerificationOutcome
+where
+    Scalar: UnsignedInteger,
+    i64: CastFrom<Scalar>,
+    Cont: Container<Element = Scalar>,
+    KeyCont: Container<Element = Scalar>,
+{
+    crs.verify(
+        &LweCompactCiphertextList::from_container(
+            lwe_ciphertext.as_ref(),
+            lwe_ciphertext.lwe_size(),
+            LweCiphertextCount(1),
+            lwe_ciphertext.ciphertext_modulus(),
+        ),
+        compact_public_key,
+        proof,
+        metadata,
+    )
+}

crates/tfhe-core-crypto/src/algorithms/misc.rs

@@ -1,6 +1,6 @@
 //! Miscellaneous algorithms.
 
-use crate::core_crypto::prelude::*;
+use crate::prelude::*;
 
 #[inline]
 pub fn divide_round<Scalar: UnsignedInteger>(numerator: Scalar, denominator: Scalar) -> Scalar {

crates/tfhe-core-crypto/src/algorithms/mod.rs

@@ -49,7 +49,7 @@ pub mod slice_algorithms;
 pub(crate) mod test;
 
 // No pub use for slice and polynomial algorithms which would not interest higher level users
-// They can still be used via `use crate::core_crypto::algorithms::slice_algorithms::*;`
+// They can still be used via `use crate::algorithms::slice_algorithms::*;`
 pub use ggsw_conversion::*;
 pub use ggsw_encryption::*;
 pub use glwe_encryption::*;

crates/tfhe-core-crypto/src/algorithms/polynomial_algorithms.rs

@@ -1,9 +1,9 @@
 //! Module providing algorithms to perform computations on polynomials modulo $X^{N} + 1$.
 
-use crate::core_crypto::algorithms::slice_algorithms::*;
-use crate::core_crypto::commons::parameters::MonomialDegree;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::*;
+use crate::commons::parameters::MonomialDegree;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Add a polynomial to the output polynomial.
 ///
@@ -1247,9 +1247,9 @@ fn induction_karatsuba_custom_mod<Scalar>(
 mod test {
     use rand::Rng;
 
-    use crate::core_crypto::algorithms::polynomial_algorithms::*;
-    use crate::core_crypto::commons::parameters::*;
-    use crate::core_crypto::commons::test_tools::*;
+    use crate::algorithms::polynomial_algorithms::*;
+    use crate::commons::parameters::*;
+    use crate::commons::test_tools::*;
 
     fn test_multiply_divide_unit_monomial<T: UnsignedTorus>() {
         //! tests if multiply_by_monomial and divide_by_monomial cancel each other

crates/tfhe-core-crypto/src/algorithms/seeded_ggsw_ciphertext_decompression.rs

@@ -1,10 +1,10 @@
 //! Module with primitives pertaining to [`SeededGgswCiphertext`] decompression.
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::math::random::Uniform;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::math::random::Uniform;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for

crates/tfhe-core-crypto/src/algorithms/seeded_ggsw_ciphertext_list_decompression.rs

@@ -1,10 +1,10 @@
 //! Module with primitives pertaining to [`SeededGgswCiphertextList`] decompression.
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::math::random::Uniform;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::math::random::Uniform;
+use crate::commons::traits::*;
+use crate::entities::*;
 use rayon::prelude::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for

crates/tfhe-core-crypto/src/algorithms/seeded_glwe_ciphertext_decompression.rs

@@ -1,9 +1,9 @@
 //! Module with primitives pertaining to [`SeededGlweCiphertext`] decompression.
 
-use crate::core_crypto::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for
 /// [`SeededGlweCiphertext`] between all functions needing it.

crates/tfhe-core-crypto/src/algorithms/seeded_glwe_ciphertext_list_decompression.rs

@@ -1,9 +1,9 @@
 //! Module with primitives pertaining to [`SeededGlweCiphertextList`] decompression.
 
-use crate::core_crypto::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for
 /// [`SeededGlweCiphertextList`] between all functions needing it.

crates/tfhe-core-crypto/src/algorithms/seeded_lwe_bootstrap_key_decompression.rs

@@ -1,9 +1,9 @@
 //! Module with primitives pertaining to [`SeededLweBootstrapKey`] decompression.
 
-use crate::core_crypto::algorithms::*;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::*;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for
 /// [`SeededLweBootstrapKey`] between all functions needing it.

crates/tfhe-core-crypto/src/algorithms/seeded_lwe_ciphertext_decompression.rs

@@ -1,10 +1,10 @@
 //! Module with primitives pertaining to [`SeededLweCiphertext`] decompression.
 
-use crate::core_crypto::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
-use crate::core_crypto::commons::ciphertext_modulus::CiphertextModulusKind;
-use crate::core_crypto::commons::generators::MaskRandomGenerator;
-use crate::core_crypto::commons::traits::*;
-use crate::core_crypto::entities::*;
+use crate::algorithms::slice_algorithms::slice_wrapping_scalar_mul_assign;
+use crate::commons::ciphertext_modulus::CiphertextModulusKind;
+use crate::commons::generators::MaskRandomGenerator;
+use crate::commons::traits::*;
+use crate::entities::*;
 
 /// Convenience function to share the core logic of the decompression algorithm for
 /// [`SeededLweCiphertext`] between all functions needing it.