chore(ci): set npm trusted publishing with node 24.12

Node is updated to be able to get npm version that
allow trusted publishing (>v11.5.1).

.github/workflows/aws_tfhe_backward_compat_tests.yml

@@ -66,14 +66,9 @@ jobs:
       - name: Checkout tfhe-rs
         uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
-          persist-credentials: 'false'
+          persist-credentials: 'true' # Needed to pull lfs data
           token: ${{ env.CHECKOUT_TOKEN }}
 
-      - name: Install latest stable
-        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
-        with:
-          toolchain: stable
-
       # Cache key is an aggregated hash of lfs files hashes
       - name: Get LFS data sha
         id: hash-lfs-data
@@ -83,7 +78,7 @@ jobs:
 
       - name: Retrieve data from cache
         id: retrieve-data-cache
-        uses: actions/cache/restore@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/restore@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         with:
           path: |
             utils/tfhe-backward-compat-data/**/*.cbor
@@ -95,6 +90,16 @@ jobs:
         run: |
           make pull_backward_compat_data
 
+      # Pull token was stored by action/checkout to be used by lfs, we don't need it anymore
+      - name: Remove git credentials
+        run: |
+          git config --local --unset-all http.https://github.com/.extraheader
+
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
+
       - name: Run backward compatibility tests
         run: |
           make test_backward_compatibility_ci
@@ -102,7 +107,7 @@ jobs:
       - name: Store data in cache
         if: steps.retrieve-data-cache.outputs.cache-hit != 'true'
         continue-on-error: true
-        uses: actions/cache/save@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/save@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         with:
           path: |
             utils/tfhe-backward-compat-data/**/*.cbor

.github/workflows/aws_tfhe_fast_tests.yml

@@ -217,7 +217,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/restore@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         with:
           path: |
             ~/.nvm
@@ -230,7 +230,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/save@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/aws_tfhe_integer_tests.yml

@@ -107,6 +107,7 @@ jobs:
       group: ${{ github.workflow_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
       cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
     runs-on: ${{ needs.setup-instance.outputs.runner-name }}
+    timeout-minutes: 480 # 8 hours
     steps:
       - name: Checkout tfhe-rs
         uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8

.github/workflows/aws_tfhe_wasm_tests.yml

@@ -78,7 +78,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/restore@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         with:
           path: |
             ~/.nvm
@@ -91,7 +91,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/save@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/benchmark_perf_regression.yml

@@ -21,20 +21,20 @@ env:
 permissions: { }
 
 jobs:
-  verify-actor:
+  verify-triggering-actor:
     name: benchmark_perf_regression/verify-actor
     if: (github.event_name == 'pull_request' &&
       (contains(github.event.label.name, 'bench-perfs-cpu') ||
       contains(github.event.label.name, 'bench-perfs-gpu'))) ||
       (github.event.issue.pull_request && startsWith(github.event.comment.body, '/bench'))
-    uses: ./.github/workflows/verify_commit_actor.yml
+    uses: ./.github/workflows/verify_triggering_actor.yml
     secrets:
       ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
       READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
 
   prepare-benchmarks:
     name: benchmark_perf_regression/prepare-benchmarks
-    needs: verify-actor
+    needs: verify-triggering-actor
     runs-on: ubuntu-latest
     outputs:
       commands: ${{ steps.set_commands.outputs.commands }}
@@ -44,7 +44,7 @@ jobs:
       custom-env: ${{ steps.get_custom_env.outputs.custom_env }}
     steps:
       - name: Checkout tfhe-rs repo
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
           persist-credentials: 'false'
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
@@ -132,7 +132,7 @@ jobs:
             gcc: 11
     steps:
       - name: Checkout tfhe-rs repo
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
           fetch-depth: 0
           persist-credentials: 'false'
@@ -159,7 +159,7 @@ jobs:
         command: ${{ fromJson(needs.prepare-benchmarks.outputs.commands) }}
     steps:
       - name: Checkout tfhe-rs repo with tags
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
           fetch-depth: 0
           persist-credentials: 'false'
@@ -213,7 +213,7 @@ jobs:
           toolchain: nightly
 
       - name: Checkout Slab repo
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
           repository: zama-ai/slab
           path: slab

.github/workflows/benchmark_wasm_client.yml

@@ -117,7 +117,7 @@ jobs:
 
       - name: Node cache restoration
         id: node-cache
-        uses: actions/cache/restore@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/restore@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         with:
           path: |
             ~/.nvm
@@ -130,7 +130,7 @@ jobs:
           make install_node
 
       - name: Node cache save
-        uses: actions/cache/save@0400d5f644dc74513175e3cd8d07132dd4860809 #v4.2.4
+        uses: actions/cache/save@0057852bfaa89a56745cba8c7296529d2fc39830 #v4.3.0
         if: steps.node-cache.outputs.cache-hit != 'true'
         with:
           path: |

.github/workflows/cargo_audit.yml

@@ -1,4 +1,6 @@
 # Run cargo audit
+name: cargo_audit
+
 on:
   workflow_dispatch:
   schedule:

.github/workflows/cargo_build.yml

@@ -18,17 +18,95 @@ permissions:
   contents: read
 
 jobs:
-  cargo-builds:
-    name: cargo_build/cargo-builds (bpr)
-    runs-on: ${{ matrix.os }}
+  prepare-parallel-pcc-matrix:
+    name: cargo_build/prepare-parallel-pcc-matrix
+    runs-on: ubuntu-latest
+    outputs:
+      matrix_command: ${{ steps.set-pcc-commands-matrix.outputs.commands }}
+    steps:
+      - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: "false"
+          token: ${{ env.CHECKOUT_TOKEN }}
 
+      # Fetch all the Make recipes that start with `pcc_batch_`
+      - name: Set pcc commands matrix
+        id: set-pcc-commands-matrix
+        run: |
+          COMMANDS=$(grep -oE '^pcc_batch_[^:]*:' Makefile | sed 's/:/\"/; s/^/\"/' | paste -sd,)
+          echo "commands=[${COMMANDS}]" >> "$GITHUB_OUTPUT"
+
+  parallel-pcc-cpu:
+    name: cargo_build/parallel-pcc-cpu
+    needs: prepare-parallel-pcc-matrix
+    runs-on: large_ubuntu_16
+    strategy:
+      matrix:
+        command: ${{fromJson(needs.prepare-parallel-pcc-matrix.outputs.matrix_command)}}
+      fail-fast: false
+    steps:
+      - name: Checkout tfhe-rs repo
+        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
+
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
+
+      - name: Run pcc checks batch
+        run: |
+          make "${COMMAND}"
+        env:
+          COMMAND: ${{ matrix.command }}
+
+  pcc-hpu:
+    name: cargo_build/pcc-hpu
+    runs-on: large_ubuntu_16
+    steps:
+      - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
+
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
+
+      - name: Run Hpu pcc checks
+        run: |
+          make pcc_hpu
+
+  build-tfhe-full:
+    name: cargo_build/build-tfhe-full
+    runs-on: ${{ matrix.os }}
     strategy:
       matrix:
         # GitHub macos-latest are now M1 macs, so use ours, we limit what runs so it will be fast
         # even with a few PRs
-        os: [large_ubuntu_16, macos-latest, windows-latest]
+        os: [large_ubuntu_16, macos-latest-xlarge, large_windows_16_latest]
       fail-fast: false
+    steps:
+      - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
 
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
+
+      - name: Build Release tfhe full
+        run: |
+          make build_tfhe_full
+
+  build:
+    name: cargo_build/build
+    runs-on: large_ubuntu_16
     steps:
       - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
         with:
@@ -41,7 +119,6 @@ jobs:
           toolchain: stable
 
       - name: Install and run newline linter checks
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           wget https://github.com/fernandrone/linelint/releases/download/0.0.6/linelint-linux-amd64
           echo "16b70fb7b471d6f95cbdc0b4e5dc2b0ac9e84ba9ecdc488f7bdf13df823aca4b linelint-linux-amd64" > checksum
@@ -50,60 +127,93 @@ jobs:
           mv linelint-linux-amd64 /usr/local/bin/linelint
           make check_newline
 
-      - name: Run pcc checks
-        if: ${{ contains(matrix.os, 'ubuntu') }}
-        run: |
-          make pcc
-
       - name: Build tfhe-csprng
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_tfhe_csprng
 
       - name: Build with MSRV
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_tfhe_msrv
 
+      - name: Build coverage tests
+        run: |
+          make build_tfhe_coverage
+
+  build-layers:
+    name: cargo_build/build-layers
+    runs-on: large_ubuntu_16
+    steps:
+      - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
+
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
+
       - name: Build Release core
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_core AVX512_SUPPORT=ON
           make build_core_experimental AVX512_SUPPORT=ON
 
       - name: Build Release boolean
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_boolean
 
       - name: Build Release shortint
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_shortint
 
       - name: Build Release integer
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_integer
 
-      - name: Build Release tfhe full
-        run: |
-          make build_tfhe_full
+  build-c-api:
+    name: cargo_build/build-c-api
+    runs-on: large_ubuntu_16
+    steps:
+      - uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
+
+      - name: Install latest stable
+        uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        with:
+          toolchain: stable
 
       - name: Build Release c_api
-        if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |
           make build_c_api
 
-      - name: Build coverage tests
-        if: ${{ contains(matrix.os, 'ubuntu') }}
-        run: |
-          make build_tfhe_coverage
-
-      - name: Run Hpu pcc checks
-        if: ${{ contains(matrix.os, 'ubuntu') }}
-        run: |
-          make pcc_hpu
-
       # The wasm build check is a bit annoying to set-up here and is done during the tests in
       # aws_tfhe_tests.yml
+
+  cargo-builds:
+    name: cargo_build/cargo-builds (bpr)
+    needs: [ parallel-pcc-cpu, pcc-hpu, build-tfhe-full, build, build-layers, build-c-api ]
+    if: ${{ always() }}
+    runs-on: ubuntu-latest
+    steps:
+      - name: Check all builds success
+        if: needs.parallel-pcc-cpu.result == 'success' &&
+          needs.pcc-hpu.result == 'success' &&
+          needs.build-tfhe-full.result == 'success' &&
+          needs.build.result == 'success' &&
+          needs.build-layers.result == 'success' &&
+          needs.build-c-api.result == 'success'
+        run: |
+          echo "All tfhe-rs build checks passed"
+
+      - name: Check builds failure
+        if: needs.parallel-pcc-cpu.result != 'success' ||
+          needs.pcc-hpu.result != 'success' ||
+          needs.build-tfhe-full.result != 'success' ||
+          needs.build.result != 'success' ||
+          needs.build-layers.result != 'success' ||
+          needs.build-c-api.result != 'success'
+        run: |
+          echo "Some tfhe-rs build checks failed"
+          exit 1

.github/workflows/ci_lint.yml

@@ -42,7 +42,7 @@ jobs:
           GH_TOKEN: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Ensure SHA pinned actions
-        uses: zgosalvez/github-actions-ensure-sha-pinned-actions@fc87bb5b5a97953d987372e74478de634726b3e5 # v3.0.25
+        uses: zgosalvez/github-actions-ensure-sha-pinned-actions@9e9574ef04ea69da568d6249bd69539ccc704e74 # v4.0.0
         with:
           allowlist: |
             slsa-framework/slsa-github-generator

.github/workflows/coprocessor-benchmark-gpu.yml

@@ -3,6 +3,22 @@ name: coprocessor-benchmark-gpu
 
 on:
   workflow_dispatch:
+    inputs:
+      profile:
+        description: "Instance type"
+        required: true
+        type: choice
+        options:
+          - "l40 (n3-L40x1)"
+          - "4-l40 (n3-L40x4)"
+          - "single-h100 (n3-H100x1)"
+          - "2-h100 (n3-H100x2)"
+          - "4-h100 (n3-H100x4)"
+          - "multi-h100 (n3-H100x8)"
+          - "multi-h100-nvlink (n3-H100x8-NVLink)"
+          - "multi-h100-sxm5 (n3-H100x8-SXM5)"
+          - "multi-h100-sxm5_fallback (n3-H100x8-SXM5)"
+
   schedule:
     # Weekly tests @ 1AM
     - cron: "0 1 * * 6"
@@ -17,7 +33,9 @@ env:
   RUST_BACKTRACE: "full"
   RUST_MIN_STACK: "8388608"
   CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
-  PROFILE: "multi-h100-sxm5 (n3-H100x8-SXM5)"
+  PROFILE_SCHEDULED_RUN: "multi-h100-sxm5 (n3-H100x8-SXM5)"
+  PROFILE_MANUAL_RUN: ${{ inputs.profile }}
+  IS_MANUAL_RUN: ${{ github.event_name == 'workflow_dispatch' }}
   BENCHMARK_TYPE: "ALL"
   OPTIMIZATION_TARGET: "throughput"
   BATCH_SIZE: "5000"
@@ -40,15 +58,25 @@ jobs:
       - name: Parse profile
         id: parse_profile
         run: |
+          if [[ ${IS_MANUAL_RUN} == true ]]; then
+            PROFILE_RAW="${PROFILE_MANUAL_RUN}"
+          else
+            PROFILE_RAW="${PROFILE_SCHEDULED_RUN}"
+          fi
           # shellcheck disable=SC2001
-          PROFILE_VAL=$(echo "${PROFILE}" | sed 's|\(.*\)[[:space:]](.*)|\1|')
+          PROFILE_VAL=$(echo "${PROFILE_RAW}" | sed 's|\(.*\)[[:space:]](.*)|\1|')
           echo "profile=$PROFILE_VAL" >> "${GITHUB_OUTPUT}"
 
       - name: Parse hardware name
         id: parse_hardware_name
         run: |
+          if [[ ${IS_MANUAL_RUN} == true ]]; then
+            PROFILE_RAW="${PROFILE_MANUAL_RUN}"
+          else
+            PROFILE_RAW="${PROFILE}"
+          fi
           # shellcheck disable=SC2001
-          PROFILE_VAL=$(echo "${PROFILE}" | sed 's|.*[[:space:]](\(.*\))|\1|')
+          PROFILE_VAL=$(echo "${PROFILE_RAW}" | sed 's|.*[[:space:]](\(.*\))|\1|')
           echo "name=$PROFILE_VAL" >> "${GITHUB_OUTPUT}"
 
   setup-instance:
@@ -130,6 +158,13 @@ jobs:
           } >> "${GITHUB_ENV}"
         working-directory: tfhe-rs/
 
+      - name: Setup Hyperstack dependencies
+        uses: ./tfhe-rs/.github/actions/gpu_setup
+        with:
+          cuda-version: ${{ matrix.cuda }}
+          gcc-version: ${{ matrix.gcc }}
+          github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
+
       - name: Check fhEVM and TFHE-rs repos
         run: |
           pwd
@@ -140,13 +175,6 @@ jobs:
         run: git lfs checkout
         working-directory: fhevm/
 
-      - name: Setup Hyperstack dependencies
-        uses: ./fhevm/.github/actions/gpu_setup
-        with:
-          cuda-version: ${{ matrix.cuda }}
-          gcc-version: ${{ matrix.gcc }}
-          github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
-
       - name: Install rust
         uses: dtolnay/rust-toolchain@e97e2d8cc328f1b50210efc529dca0028893a2d9 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
@@ -154,7 +182,7 @@ jobs:
 
       - name: Install cargo dependencies
         run: |
-          sudo apt-get install -y protobuf-compiler cmake pkg-config libssl-dev \
+          sudo apt-get install -y protobuf-compiler pkg-config libssl-dev \
                                   libclang-dev docker-compose-v2 docker.io acl
           sudo usermod -aG docker "$USER"
           newgrp docker
@@ -165,7 +193,7 @@ jobs:
         uses: foundry-rs/foundry-toolchain@82dee4ba654bd2146511f85f0d013af94670c4de
 
       - name: Cache cargo
-        uses: actions/cache@0400d5f644dc74513175e3cd8d07132dd4860809 # v4.2.4
+        uses: actions/cache@0057852bfaa89a56745cba8c7296529d2fc39830 # v4.3.0
         with:
           path: |
             ~/.cargo/registry
@@ -181,9 +209,16 @@ jobs:
           username: ${{ github.actor }}
           password: ${{ secrets.GITHUB_TOKEN }}
 
+      - name: Login to Chainguard Registry
+        uses: docker/login-action@184bdaa0721073962dff0199f1fb9940f07167d1 # v3.5.0
+        with:
+          registry: cgr.dev
+          username: ${{ secrets.CGR_USERNAME }}
+          password: ${{ secrets.CGR_PASSWORD }}
+
       - name: Init database
         run: make init_db
-        working-directory: fhevm/coprocessor/fhevm-engine/coprocessor
+        working-directory: fhevm/coprocessor/fhevm-engine/tfhe-worker
 
       - name: Use Node.js
         uses: actions/setup-node@a0853c24544627f65ddf259abe73b1d18a591444 # v5.0.0
@@ -203,8 +238,12 @@ jobs:
 
       - name: Profile erc20 no-cmux benchmark on GPU
         run: |
-          BENCHMARK_BATCH_SIZE="${BATCH_SIZE}" FHEVM_DF_SCHEDULE="${SCHEDULING_POLICY}" BENCHMARK_TYPE="LATENCY" OPTIMIZATION_TARGET="${OPTIMIZATION_TARGET}" make -e "profile_erc20_gpu"
-        working-directory: fhevm/coprocessor/fhevm-engine/coprocessor
+          BENCHMARK_BATCH_SIZE="${BATCH_SIZE}" \
+          FHEVM_DF_SCHEDULE="${SCHEDULING_POLICY}" \
+          BENCHMARK_TYPE="THROUGHPUT_200" \
+          OPTIMIZATION_TARGET="${OPTIMIZATION_TARGET}" \
+          make -e "profile_erc20_gpu"
+        working-directory: fhevm/coprocessor/fhevm-engine/tfhe-worker
 
       - name: Get nsys profile name
         id: nsys_profile_name
@@ -215,7 +254,7 @@ jobs:
           REPORT_NAME: ${{ steps.nsys_profile_name.outputs.profile }}
         run: |
           mv report1.nsys-rep ${{ env.REPORT_NAME }}
-        working-directory: fhevm/coprocessor/fhevm-engine/coprocessor
+        working-directory: fhevm/coprocessor/fhevm-engine/tfhe-worker
 
       - name: Upload profile artifact
         env:
@@ -223,17 +262,17 @@ jobs:
         uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02
         with:
           name: ${{ env.REPORT_NAME }}
-          path: fhevm/coprocessor/fhevm-engine/coprocessor/${{ env.REPORT_NAME }}
+          path: fhevm/coprocessor/fhevm-engine/tfhe-worker/${{ env.REPORT_NAME }}
 
       - name: Run latency benchmark on GPU
         run: |
           BENCHMARK_BATCH_SIZE="${BATCH_SIZE}" FHEVM_DF_SCHEDULE="${SCHEDULING_POLICY}" BENCHMARK_TYPE="LATENCY" OPTIMIZATION_TARGET="${OPTIMIZATION_TARGET}" make -e "benchmark_${BENCHMARKS}_gpu"
-        working-directory: fhevm/coprocessor/fhevm-engine/coprocessor
+        working-directory: fhevm/coprocessor/fhevm-engine/tfhe-worker
 
       - name: Run throughput benchmarks on GPU
         run: |
           BENCHMARK_BATCH_SIZE="${BATCH_SIZE}" FHEVM_DF_SCHEDULE="${SCHEDULING_POLICY}" BENCHMARK_TYPE="THROUGHPUT_200" OPTIMIZATION_TARGET="${OPTIMIZATION_TARGET}" make -e "benchmark_${BENCHMARKS}_gpu"
-        working-directory: fhevm/coprocessor/fhevm-engine/coprocessor
+        working-directory: fhevm/coprocessor/fhevm-engine/tfhe-worker
 
       - name: Parse results
         run: |
@@ -246,7 +285,7 @@ jobs:
           --commit-date "${COMMIT_DATE}" \
           --bench-date "${BENCH_DATE}" \
           --walk-subdirs \
-          --crate "coprocessor/fhevm-engine/coprocessor" \
+          --crate "coprocessor/fhevm-engine/tfhe-worker" \
           --name-suffix "operation_batch_size_${BATCH_SIZE}-schedule_${SCHEDULING_POLICY}-optimization_target_${OPTIMIZATION_TARGET}"
         working-directory: fhevm/
 

.github/workflows/gpu_full_multi_gpu_tests.yml

@@ -86,7 +86,7 @@ jobs:
           slab-url: ${{ secrets.SLAB_BASE_URL }}
           job-secret: ${{ secrets.JOB_SECRET }}
           backend: hyperstack
-          profile: multi-gpu-test
+          profile: 4-l40
 
       # This instance will be spawned especially for pull-request from forked repository
       - name: Start GitHub instance

.github/workflows/gpu_integer_long_run_tests.yml

@@ -43,7 +43,7 @@ jobs:
           slab-url: ${{ secrets.SLAB_BASE_URL }}
           job-secret: ${{ secrets.JOB_SECRET }}
           backend: hyperstack
-          profile: multi-gpu-test
+          profile: 4-l40
 
   cuda-tests:
     name: gpu_integer_long_run_tests/cuda-tests

.github/workflows/make_release_common.yml

@@ -1,33 +1,31 @@
-# Publish new release of tfhe-rs on various platform.
-name: make_release
+# Common workflow to make crate release
+name: make_release_common
 
 on:
-  workflow_dispatch:
+  workflow_call:
     inputs:
-      dry_run:
-        description: "Dry-run"
+      package-name:
+        type: string
+        required: true
+      dry-run:
         type: boolean
         default: true
-      push_to_crates:
-        description: "Push to crate"
-        type: boolean
-        default: true
-      push_web_package:
-        description: "Push web js package"
-        type: boolean
-        default: true
-      push_node_package:
-        description: "Push node js package"
-        type: boolean
-        default: true
-      npm_latest_tag:
-        description: "Set NPM tag as latest"
-        type: boolean
-        default: false
+    secrets:
+      REPO_CHECKOUT_TOKEN:
+        required: true
+      SLACK_CHANNEL:
+        required: true
+      BOT_USERNAME:
+        required: true
+      SLACK_WEBHOOK:
+        required: true
+      ALLOWED_TEAM:
+        required: true
+      READ_ORG_TOKEN:
+        required: true
 
 env:
   ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
-  NPM_TAG: ""
   SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
   SLACK_ICON: https://pbs.twimg.com/profile_images/1274014582265298945/OjBKP9kn_400x400.png
   SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
@@ -36,18 +34,18 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release/verify-tag
+  verify-triggering-actor:
+    name: make_release_common/verify-triggering-actor
     if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
+    uses: ./.github/workflows/verify_triggering_actor.yml
     secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      ALLOWED_TEAM: ${{ secrets.ALLOWED_TEAM }}
       READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
 
   package:
-    name: make_release/package
+    name: make_release_common/package
     runs-on: ubuntu-latest
-    needs: verify-tag
+    needs: verify-triggering-actor
     outputs:
       hash: ${{ steps.hash.outputs.hash }}
     steps:
@@ -58,20 +56,23 @@ jobs:
           persist-credentials: 'false'
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
       - name: Prepare package
+        env:
+          PACKAGE: ${{ inputs.package-name }}
         run: |
-          cargo package -p tfhe
+          cargo package -p "${PACKAGE}"
       - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
         with:
-          name: crate
+          name: crate-${{ inputs.package-name }}
           path: target/package/*.crate
       - name: generate hash
         id: hash
         run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
 
+
   provenance:
-    name: make_release/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [package]
+    name: make_release_common/provenance
+    if: ${{ !inputs.dry-run  }}
+    needs: package
     uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
     permissions:
       # Needed to detect the GitHub Actions environment
@@ -84,14 +85,14 @@ jobs:
       # SHA-256 hashes of the Crate package.
       base64-subjects: ${{ needs.package.outputs.hash }}
 
+
   publish_release:
-    name: make_release/publish_release
-    needs: [package] # for comparing hashes
+    name: make_release_common/publish-release
+    needs: package
     runs-on: ubuntu-latest
-    # For provenance of npmjs publish
     permissions:
-      contents: read
-      id-token: write # also needed for OIDC token exchange on crates.io
+      # Needed for OIDC token exchange on crates.io
+      id-token: write
     steps:
       - name: Checkout
         uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
@@ -99,28 +100,27 @@ jobs:
           fetch-depth: 0
           persist-credentials: 'false'
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Create NPM version tag
-        if: ${{ inputs.npm_latest_tag }}
-        run: |
-          echo "NPM_TAG=latest" >> "${GITHUB_ENV}"
+
       - name: Download artifact
         uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
         with:
-          name: crate
+          name: crate-${{ inputs.package-name }}
           path: target/package
+
       - name: Authenticate on registry
         uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
         id: auth
+
       - name: Publish crate.io package
-        if: ${{ inputs.push_to_crates }}
         env:
           CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
+          PACKAGE: ${{ inputs.package-name }}
+          DRY_RUN: ${{ inputs.dry-run && '--dry-run' || '' }}
         run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
+          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish
           # would fail. This is safe since DRY_RUN is handled in the env section above.
           # shellcheck disable=SC2086
-          cargo publish -p tfhe ${DRY_RUN}
+          cargo publish -p "${PACKAGE}" ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash
@@ -132,45 +132,12 @@ jobs:
         uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
         env:
           SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe crate - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-
-      - name: Build web package
-        if: ${{ inputs.push_web_package }}
-        run: |
-          make build_web_js_api_parallel
-
-      - name: Publish web package
-        if: ${{ inputs.push_web_package }}
-        uses: JS-DevTools/npm-publish@19c28f1ef146469e409470805ea4279d47c3d35c
-        with:
-          token: ${{ secrets.NPM_TOKEN }}
-          package: tfhe/pkg/package.json
-          dry-run: ${{ inputs.dry_run }}
-          tag: ${{ env.NPM_TAG }}
-          provenance: true
-
-      - name: Build Node package
-        if: ${{ inputs.push_node_package }}
-        run: |
-          rm -rf tfhe/pkg
-
-          make build_node_js_api
-          sed -i 's/"tfhe"/"node-tfhe"/g' tfhe/pkg/package.json
-
-      - name: Publish Node package
-        if: ${{ inputs.push_node_package }}
-        uses: JS-DevTools/npm-publish@19c28f1ef146469e409470805ea4279d47c3d35c
-        with:
-          token: ${{ secrets.NPM_TOKEN }}
-          package: tfhe/pkg/package.json
-          dry-run: ${{ inputs.dry_run }}
-          tag: ${{ env.NPM_TAG }}
-          provenance: true
+          SLACK_MESSAGE: "SLSA ${{ inputs.package-name }} - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
 
       - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
+        if: ${{ failure() }}
         continue-on-error: true
         uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
         env:
           SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe release failed: (${{ env.ACTION_RUN_URL }})"
+          SLACK_MESSAGE: "${{ inputs.package-name }} release finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"

.github/workflows/make_release_cuda.yml

@@ -18,17 +18,17 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_cuda/verify-tag
+  verify-triggering-actor:
+    name: make_release_cuda/verify-triggering-actor
     if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
+    uses: ./.github/workflows/verify_triggering_actor.yml
     secrets:
       ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
       READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
 
   setup-instance:
     name: make_release_cuda/setup-instance
-    needs: verify-tag
+    needs: verify-triggering-actor
     runs-on: ubuntu-latest
     outputs:
       runner-name: ${{ steps.start-instance.outputs.label }}
@@ -101,6 +101,12 @@ jobs:
       - name: Prepare package
         run: |
           cargo package -p tfhe-cuda-backend
+
+      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
+        with:
+          name: crate-tfhe-cuda-backend
+          path: target/package/*.crate
+
       - name: generate hash
         id: hash
         run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
@@ -169,6 +175,12 @@ jobs:
         env:
           GCC_VERSION: ${{ matrix.gcc }}
 
+      - name: Download artifact
+        uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
+        with:
+          name: crate-tfhe-cuda-backend
+          path: target/package
+
       - name: Authenticate on registry
         uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
         id: auth

.github/workflows/make_release_hpu.yml

@@ -18,43 +18,12 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_hpu/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
-    secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
-      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
-
-  package:
-    name: make_release_hpu/package
-    runs-on: ubuntu-latest
-    needs: verify-tag
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-hpu-backend
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance:
-    name: make_release_hpu/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [package]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_hpu/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-hpu-backend"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -62,55 +31,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-  publish_release:
-    name: make_release_hpu/publish-release
-    runs-on: ubuntu-latest
-    needs: [verify-tag, package] # for comparing hashes
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
-        run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
-          # would fail. This is safe since DRY_RUN is handled in the env section above.
-          # shellcheck disable=SC2086
-          cargo publish -p tfhe-hpu-backend ${DRY_RUN}
-
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-hpu-backend crate - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-hpu-backend release failed: (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/make_release_tfhe.yml

@@ -0,0 +1,123 @@
+# Publish new release of tfhe-rs on various platform.
+name: make_release_tfhe
+
+on:
+  workflow_dispatch:
+    inputs:
+      dry_run:
+        description: "Dry-run"
+        type: boolean
+        default: true
+      push_to_crates:
+        description: "Push to crate"
+        type: boolean
+        default: true
+      push_web_package:
+        description: "Push web js package"
+        type: boolean
+        default: true
+      push_node_package:
+        description: "Push node js package"
+        type: boolean
+        default: true
+      npm_latest_tag:
+        description: "Set NPM tag as latest"
+        type: boolean
+        default: false
+
+env:
+  ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
+  NPM_TAG: ""
+  SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+  SLACK_ICON: https://pbs.twimg.com/profile_images/1274014582265298945/OjBKP9kn_400x400.png
+  SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
+  SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+
+permissions: {}
+
+jobs:
+  make-release:
+    name: make_release_tfhe/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    if: ${{ inputs.push_to_crates }}
+    with:
+      package-name: "tfhe"
+      dry-run: ${{ inputs.dry_run }}
+    permissions:
+      actions: read # Needed to detect the GitHub Actions environment
+      id-token: write # Needed to create the provenance via GitHub OIDC
+      contents: write # Needed to upload assets/artifacts
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
+
+  make-release-js:
+    name: make_release_tfhe/make-release-js
+    needs: make-release
+    if: ${{ always() && needs.make-release.result != 'failure' }}
+    runs-on: ubuntu-latest
+    # For provenance of npmjs publish
+    permissions:
+      contents: read
+      id-token: write # also needed for OIDC token exchange on crates.io and npmjs.com
+    steps:
+      - name: Checkout
+        uses: actions/checkout@1af3b93b6815bc44a9784bd300feb67ff0d1eeb3 # v6.0.0
+        with:
+          fetch-depth: 0
+          persist-credentials: 'false'
+          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+
+      - name: Create NPM version tag
+        if: ${{ inputs.npm_latest_tag }}
+        run: |
+          echo "NPM_TAG=latest" >> "${GITHUB_ENV}"
+
+      - name: Build web package
+        if: ${{ inputs.push_web_package }}
+        run: |
+          make build_web_js_api_parallel
+
+      - name: Authenticate on NPM
+        uses: actions/setup-node@2028fbc5c25fe9cf00d9f06a71cc4710d4507903 # v6.0.0
+        with:
+          node-version: '24'
+          registry-url: 'https://registry.npmjs.org'
+
+      - name: Publish web package
+        if: ${{ inputs.push_web_package }}
+        uses: JS-DevTools/npm-publish@7f8fe47b3bea1be0c3aec2b717c5ec1f3e03410b
+        with:
+          package: tfhe/pkg/package.json
+          dry-run: ${{ inputs.dry_run }}
+          tag: ${{ env.NPM_TAG }}
+          provenance: true
+
+      - name: Build Node package
+        if: ${{ inputs.push_node_package }}
+        run: |
+          rm -rf tfhe/pkg
+
+          make build_node_js_api
+          sed -i 's/"tfhe"/"node-tfhe"/g' tfhe/pkg/package.json
+
+      - name: Publish Node package
+        if: ${{ inputs.push_node_package }}
+        uses: JS-DevTools/npm-publish@7f8fe47b3bea1be0c3aec2b717c5ec1f3e03410b
+        with:
+          package: tfhe/pkg/package.json
+          dry-run: ${{ inputs.dry_run }}
+          tag: ${{ env.NPM_TAG }}
+          provenance: true
+
+      - name: Slack Notification
+        if: ${{ failure() }}
+        continue-on-error: true
+        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
+        env:
+          SLACK_COLOR: ${{ job.status }}
+          SLACK_MESSAGE: "tfhe release finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"

.github/workflows/make_release_tfhe_csprng.yml

@@ -8,53 +8,15 @@ on:
         type: boolean
         default: true
 
-env:
-  ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
-  SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
-  SLACK_ICON: https://pbs.twimg.com/profile_images/1274014582265298945/OjBKP9kn_400x400.png
-  SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
-  SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
-
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_tfhe_csprng/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
-    secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
-      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
-
-  package:
-    name: make_release_tfhe_csprng/package
-    runs-on: ubuntu-latest
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-csprng
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate-tfhe-csprng
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-
-  provenance:
-    name: make_release_tfhe_csprng/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [package]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_tfhe_csprng/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-csprng"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -62,56 +24,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-
-  publish_release:
-    name: make_release_tfhe_csprng/publish-release
-    needs: [verify-tag, package]
-    runs-on: ubuntu-latest
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Download artifact
-        uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
-        with:
-          name: crate-tfhe-csprng
-          path: target/package
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
-        run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
-          # would fail. This is safe since DRY_RUN is handled in the env section above.
-          # shellcheck disable=SC2086
-          cargo publish -p tfhe-csprng ${DRY_RUN}
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-csprng - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-csprng release finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/make_release_tfhe_fft.yml

@@ -19,43 +19,12 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_tfhe_fft/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
-    secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
-      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
-
-  package:
-    name: make_release_tfhe_fft/package
-    runs-on: ubuntu-latest
-    needs: verify-tag
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-fft
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance:
-    name: make_release_tfhe_fft/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [package]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_tfhe_fft/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-fft"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -63,55 +32,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-  publish_release:
-    name: make_release_tfhe_fft/publish-release
-    runs-on: ubuntu-latest
-    needs: [verify-tag, package] # for comparing hashes
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
-        run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
-          # would fail. This is safe since DRY_RUN is handled in the env section above.
-          # shellcheck disable=SC2086
-          cargo publish -p tfhe-fft ${DRY_RUN}
-
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-fft crate - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-fft release failed: (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/make_release_tfhe_ntt.yml

@@ -19,43 +19,12 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_tfhe_ntt/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
-    secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
-      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
-
-  package:
-    name: make_release_tfhe_ntt/package
-    runs-on: ubuntu-latest
-    needs: verify-tag
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-ntt
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance:
-    name: make_release_tfhe_ntt/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [package]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_tfhe_ntt/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-ntt"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -63,55 +32,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-  publish_release:
-    name: make_release_tfhe_ntt/publish-release
-    runs-on: ubuntu-latest
-    needs: [verify-tag, package] # for comparing hashes
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
-        run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
-          # would fail. This is safe since DRY_RUN is handled in the env section above.
-          # shellcheck disable=SC2086
-          cargo publish -p tfhe-ntt ${DRY_RUN}
-
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-ntt crate - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-ntt release failed: (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/make_release_tfhe_versionable.yml

@@ -2,6 +2,11 @@ name: make_release_tfhe_versionable
 
 on:
   workflow_dispatch:
+    inputs:
+      dry_run:
+        description: "Dry-run"
+        type: boolean
+        default: true
 
 env:
   ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
@@ -13,41 +18,34 @@ env:
 permissions: {}
 
 jobs:
-  verify-tag:
-    name: make_release_tfhe_versionable/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
+  make-release-derive:
+    name: make_release_tfhe_versionable/make-release-derive
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-versionable-derive"
+      dry-run: ${{ inputs.dry_run }}
+    permissions:
+      # Needed to detect the GitHub Actions environment
+      actions: read
+      # Needed to create the provenance via GitHub OIDC
+      id-token: write
+      # Needed to upload assets/artifacts
+      contents: write
     secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
       ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
       READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
 
-  package-derive:
-    name: make_release_tfhe_versionable/package-derive
-    runs-on: ubuntu-latest
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-versionable-derive
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate-tfhe-versionable-derive
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance-derive:
-    name: make_release_tfhe_versionable/provenance-derive
-    needs: [package-derive]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_tfhe_versionable/make-release
+    needs: make-release-derive
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-versionable"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -55,132 +53,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package-derive.outputs.hash }}
-
-  publish_release-derive:
-    name: make_release_tfhe_versionable/publish_release_derive
-    needs: [ verify-tag, package-derive ] # for comparing hashes
-    runs-on: ubuntu-latest
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Download artifact
-        uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
-        with:
-          name: crate-tfhe-versionable-derive
-          path: target/package
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-        run: |
-          cargo publish -p tfhe-versionable-derive
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package-derive.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-versionable-derive - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-versionable-derive release finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"
-
-  package:
-    name: make_release_tfhe_versionable/package
-    needs: publish_release-derive
-    runs-on: ubuntu-latest
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-versionable
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate-tfhe-versionable
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance:
-    name: make_release_tfhe_versionable/provenance
-    needs: package
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
-    permissions:
-      # Needed to detect the GitHub Actions environment
-      actions: read
-      # Needed to create the provenance via GitHub OIDC
-      id-token: write
-      # Needed to upload assets/artifacts
-      contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-  publish_release:
-    name: make_release_tfhe_versionable/publish-release
-    needs: package # for comparing hashes
-    runs-on: ubuntu-latest
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Download artifact
-        uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
-        with:
-          name: crate-tfhe-versionable
-          path: target/package
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-        run: |
-          cargo publish -p tfhe-versionable
-      - name: Generate hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-versionable - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-versionable release finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/make_release_zk_pok.yml

@@ -18,43 +18,12 @@ env:
 permissions: { }
 
 jobs:
-  verify-tag:
-    name: make_release_zk_pok/verify-tag
-    if: startsWith(github.ref, 'refs/tags/')
-    uses: ./.github/workflows/verify_commit_actor.yml
-    secrets:
-      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
-      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}
-
-  package:
-    name: make_release_zk_pok/package
-    runs-on: ubuntu-latest
-    needs: verify-tag
-    outputs:
-      hash: ${{ steps.hash.outputs.hash }}
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Prepare package
-        run: |
-          cargo package -p tfhe-zk-pok
-      - uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02 # v4.6.2
-        with:
-          name: crate-zk-pok
-          path: target/package/*.crate
-      - name: generate hash
-        id: hash
-        run: cd target/package && echo "hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-
-  provenance:
-    name: make_release_zk_pok/provenance
-    if: ${{ !inputs.dry_run  }}
-    needs: [ package ]
-    uses: slsa-framework/slsa-github-generator/.github/workflows/generator_generic_slsa3.yml@v2.1.0
+  make-release:
+    name: make_release_zk_pok/make-release
+    uses: ./.github/workflows/make_release_common.yml
+    with:
+      package-name: "tfhe-zk-pok"
+      dry-run: ${{ inputs.dry_run }}
     permissions:
       # Needed to detect the GitHub Actions environment
       actions: read
@@ -62,55 +31,10 @@ jobs:
       id-token: write
       # Needed to upload assets/artifacts
       contents: write
-    with:
-      # SHA-256 hashes of the Crate package.
-      base64-subjects: ${{ needs.package.outputs.hash }}
-
-  publish_release:
-    name: make_release_zk_pok/publish-release
-    needs: [ verify-tag, package ] # for comparing hashes
-    runs-on: ubuntu-latest
-    permissions:
-      # Needed for OIDC token exchange on crates.io
-      id-token: write
-    steps:
-      - name: Checkout
-        uses: actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8 # v5.0.0
-        with:
-          fetch-depth: 0
-          persist-credentials: 'false'
-          token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      - name: Download artifact
-        uses: actions/download-artifact@634f93cb2916e3fdff6788551b99b062d0335ce0 # v5.0.0
-        with:
-          name: crate-zk-pok
-          path: target/package
-      - name: Authenticate on registry
-        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
-        id: auth
-      - name: Publish crate.io package
-        env:
-          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
-          DRY_RUN: ${{ inputs.dry_run && '--dry-run' || '' }}
-        run: |
-          # DRY_RUN expansion cannot be double quoted when variable contains empty string otherwise cargo publish 
-          # would fail. This is safe since DRY_RUN is handled in the env section above.
-          # shellcheck disable=SC2086
-          cargo publish -p tfhe-zk-pok ${DRY_RUN}
-      - name: Verify hash
-        id: published_hash
-        run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"
-      - name: Slack notification (hashes comparison)
-        if: ${{ needs.package.outputs.hash != steps.published_hash.outputs.pub_hash }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: failure
-          SLACK_MESSAGE: "SLSA tfhe-zk-pok crate - hash comparison failure: (${{ env.ACTION_RUN_URL }})"
-      - name: Slack Notification
-        if: ${{ failure() || (cancelled() && github.event_name != 'pull_request') }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661 # v2.3.3
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "tfhe-zk-pok release failed: (${{ env.ACTION_RUN_URL }})"
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      ALLOWED_TEAM: ${{ secrets.RELEASE_TEAM }}
+      READ_ORG_TOKEN: ${{ secrets.READ_ORG_TOKEN }}

.github/workflows/verify_triggering_actor.yml

@@ -1,5 +1,5 @@
-# Verify a commit actor
-name: verify_commit_actor
+# Verify a triggering actor
+name: verify_triggering_actor
 
 on:
   workflow_call:
@@ -13,7 +13,7 @@ permissions: {}
 
 jobs:
   check-actor:
-    name: verify_commit_actor/check-actor
+    name: verify_triggering_actor/check-actor
     runs-on: ubuntu-latest
     steps:
       # Check triggering actor membership

Cargo.toml

@@ -1,5 +1,5 @@
 [workspace]
-resolver = "2"
+resolver = "3"
 members = [
     "tfhe",
     "tfhe-benchmark",
@@ -24,7 +24,7 @@ exclude = [
 ]
 [workspace.dependencies]
 aligned-vec = { version = "0.6", default-features = false }
-bytemuck = "1.14.3"
+bytemuck = "<1.24"
 dyn-stack = { version = "0.11", default-features = false }
 itertools = "0.14"
 num-complex = "0.4"

Makefile

@@ -22,12 +22,13 @@ BENCH_TYPE?=latency
 BENCH_PARAM_TYPE?=classical
 BENCH_PARAMS_SET?=default
 BENCH_CUSTOM_COMMAND:=
-NODE_VERSION=22.6
+NODE_VERSION=24.12
 BACKWARD_COMPAT_DATA_DIR=utils/tfhe-backward-compat-data
 WASM_PACK_VERSION="0.13.1"
 WASM_BINDGEN_VERSION:=$(shell cargo tree --target wasm32-unknown-unknown -e all --prefix none | grep "wasm-bindgen v" | head -n 1 | cut -d 'v' -f2)
 WEB_RUNNER_DIR=web-test-runner
 WEB_SERVER_DIR=tfhe/web_wasm_parallel_tests
+TYPOS_VERSION=1.39.0
 # 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
@@ -165,9 +166,8 @@ install_cargo_audit: install_rs_build_toolchain
 
 .PHONY: install_typos_checker # Install typos checker
 install_typos_checker: install_rs_build_toolchain
-	@typos --version > /dev/null 2>&1 || \
-	cargo $(CARGO_RS_BUILD_TOOLCHAIN) install --locked typos-cli || \
-	( echo "Unable to install typos-cli, unknown error." && exit 1 )
+	@./scripts/install_typos.sh --rust-toolchain $(CARGO_RS_BUILD_TOOLCHAIN) \
+	--typos-version $(TYPOS_VERSION)
 
 .PHONY: install_zizmor # Install zizmor workflow security checker
 install_zizmor: install_rs_build_toolchain
@@ -1004,6 +1004,11 @@ test_list_gpu: install_rs_build_toolchain install_cargo_nextest
 		--features=integer,internal-keycache,gpu,zk-pok -p tfhe \
 		-E "test(/.*gpu.*/)"
 
+.PHONY: build_one_hl_api_test_gpu
+build_one_hl_api_test_gpu: install_rs_build_toolchain
+	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --no-run \
+		   --features=integer,gpu-debug -vv -p tfhe -- "$${TEST}" --test-threads=1 --nocapture
+
 test_high_level_api_hpu: install_rs_build_toolchain install_cargo_nextest
 ifeq ($(HPU_CONFIG), v80)
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) nextest run --cargo-profile $(CARGO_PROFILE) \
@@ -1184,6 +1189,8 @@ check_compile_tests: install_rs_build_toolchain
 		--features=experimental,boolean,shortint,integer,internal-keycache \
 		-p tfhe
 
+.PHONY: check_compile_tests_c_api # Build C API tests without running them
+check_compile_tests_c_api: install_rs_build_toolchain
 	@if [[ "$(OS)" == "Linux" || "$(OS)" == "Darwin" ]]; then \
 		"$(MAKE)" build_c_api && \
 		./scripts/c_api_tests.sh --build-only --cargo-profile "$(CARGO_PROFILE)"; \
@@ -1657,11 +1664,38 @@ sha256_bool: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_CHECK_TOOLCHAIN) run --profile $(CARGO_PROFILE) \
 	--example sha256_bool --features=boolean
 
-.PHONY: pcc # pcc stands for pre commit checks (except GPU)
-pcc: no_tfhe_typo no_dbg_log check_parameter_export_ok check_fmt check_typos lint_doc \
-check_md_docs_are_tested check_intra_md_links check_doc_paths_use_dash \
-clippy_all check_compile_tests test_tfhe_lints \
-tfhe_lints
+.PHONY: pcc # pcc stands for pre commit checks for CPU compilation
+pcc: pcc_batch_1 pcc_batch_2 pcc_batch_3 pcc_batch_4 pcc_batch_5 pcc_batch_6 pcc_batch_7
+
+#
+# PCC split into several batches to speed-up CI feedback.
+# Each batch have roughly the same execution time.
+# Durations are given from GitHub Ubuntu large runner with 16 CPU.
+#
+
+.PHONY: pcc_batch_1 # duration: 6'10''
+pcc_batch_1: no_tfhe_typo no_dbg_log check_parameter_export_ok check_fmt check_typos lint_doc \
+check_md_docs_are_tested check_intra_md_links check_doc_paths_use_dash test_tfhe_lints tfhe_lints \
+clippy_rustdoc
+
+.PHONY: pcc_batch_2 # duration: 6'10''
+pcc_batch_2: clippy clippy_all_targets
+
+.PHONY: pcc_batch_3 # duration: 6'50''
+pcc_batch_3: clippy_shortint clippy_integer
+
+.PHONY: pcc_batch_4 # duration: 7'40''
+pcc_batch_4: clippy_core clippy_js_wasm_api clippy_ws_tests clippy_bench
+
+.PHONY: pcc_batch_5 # duration: 7'20''
+pcc_batch_5: clippy_tfhe_lints check_compile_tests clippy_backward_compat_data
+
+.PHONY: pcc_batch_6  # duration: 4'50'' (shortest one, extend it with further checks)
+pcc_batch_6: clippy_boolean clippy_c_api clippy_tasks clippy_tfhe_csprng clippy_zk_pok \
+clippy_trivium clippy_versionable clippy_param_dedup
+
+.PHONY: pcc_batch_7 # duration: 7'50'' (currently PCC execution bottleneck)
+pcc_batch_7: check_compile_tests_c_api
 
 .PHONY: pcc_gpu # pcc stands for pre commit checks for GPU compilation
 pcc_gpu: check_rust_bindings_did_not_change clippy_rustdoc_gpu \

README.md

@@ -45,7 +45,7 @@ production-ready library for all the advanced features of TFHE.
 - **Short integer API** that enables exact, unbounded FHE integer arithmetics with up to 8 bits of message space
 - **Size-efficient public key encryption**
 - **Ciphertext and server key compression** for efficient data transfer
-- **Full Rust API, C bindings to the Rust High-Level API, and client-side Javascript API using WASM**.
+- **Full Rust API, C bindings to the Rust High-Level API, and client-side JavaScript API using WASM**.
 
 *Learn more about TFHE-rs features in the [documentation](https://docs.zama.ai/tfhe-rs/readme).*
 <br></br>
@@ -79,7 +79,7 @@ tfhe = { version = "*", features = ["boolean", "shortint", "integer"] }
 ```
 
 > [!Note]
-> Note: You need to use Rust version >= 1.84 to compile TFHE-rs.
+> Note: You need Rust version 1.84 or newer to compile TFHE-rs. You can check your version with `rustc --version`.
 
 > [!Note]
 > Note: AArch64-based machines are not supported for Windows as it's currently missing an entropy source to be able to seed the [CSPRNGs](https://en.wikipedia.org/wiki/Cryptographically_secure_pseudorandom_number_generator) used in TFHE-rs.
@@ -147,7 +147,7 @@ To run this code, use the following command:
 
 > [!Note]
 > Note that when running code that uses `TFHE-rs`, it is highly recommended
-to run in release mode with cargo's `--release` flag to have the best performances possible.
+to run in release mode with cargo's `--release` flag to have the best performance possible.
 
 *Find an example with more explanations in [this part of the documentation](https://docs.zama.ai/tfhe-rs/get-started/quick-start)*
 

_typos.toml

@@ -13,6 +13,7 @@ extend-ignore-identifiers-re = [
     # Example in trivium
     "C9217BA0D762ACA1",
     "0x[0-9a-fA-F]+",
+    "xrt_coreutil",
 ]
 
 [files]

backends/tfhe-cuda-backend/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "tfhe-cuda-backend"
-version = "0.11.0"
+version = "0.12.0-alpha.2"
 edition = "2021"
 authors = ["Zama team"]
 license = "BSD-3-Clause-Clear"

backends/tfhe-cuda-backend/build.rs

@@ -48,7 +48,6 @@ fn main() {
         // Conditionally pass the "USE_NVTOOLS" variable to CMake if the feature is enabled
         if cfg!(feature = "profile") {
             cmake_config.define("USE_NVTOOLS", "ON");
-            println!("cargo:rustc-link-lib=nvToolsExt");
         } else {
             cmake_config.define("USE_NVTOOLS", "OFF");
         }

backends/tfhe-cuda-backend/cuda/CMakeLists.txt

@@ -86,6 +86,7 @@ if(CMAKE_BUILD_TYPE_LOWERCASE STREQUAL "debug")
   message("Compiling in Debug mode")
   add_definitions(-DDEBUG)
   set(OPTIMIZATION_FLAGS "${OPTIMIZATION_FLAGS} -O0 -G -g")
+  set(USE_NVTOOLS 1)
 else()
   # Release mode
   message("Compiling in Release mode")

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

@@ -84,6 +84,11 @@ public:
     return CudaStreams(_streams, _gpu_indexes, 1);
   }
 
+  // Returns a subset containing only the ith stream
+  CudaStreams subset(int i) const {
+    return CudaStreams(&_streams[i], &_gpu_indexes[i], 1);
+  }
+
   // Synchronize all the streams in the set
   void synchronize() const {
     for (uint32_t i = 0; i < _gpu_count; i++) {

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

@@ -4,26 +4,6 @@
 #include "../../pbs/pbs_enums.h"
 #include "../integer.h"
 
-typedef struct {
-  void *ptr;
-  uint32_t num_radix_blocks;
-  uint32_t lwe_dimension;
-} CudaLweCiphertextListFFI;
-
-typedef struct {
-  void *ptr;
-  uint32_t storage_log_modulus;
-  uint32_t lwe_per_glwe;
-  // Input LWEs are grouped by groups of `lwe_per_glwe`(the last group may be
-  // smaller)
-  // Each group is then packed into one GLWE with `lwe_per_glwe` bodies (one for
-  // each LWE of the group). In the end the total number of bodies is equal to
-  // the number of input LWE
-  uint32_t total_lwe_bodies_count;
-  uint32_t glwe_dimension;
-  uint32_t polynomial_size;
-} CudaPackedGlweCiphertextListFFI;
-
 extern "C" {
 uint64_t scratch_cuda_integer_compress_radix_ciphertext_64(
     CudaStreamsFFI streams, int8_t **mem_ptr,

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

@@ -86,6 +86,26 @@ typedef struct {
   bool const divisor_has_more_bits_than_numerator;
 } CudaScalarDivisorFFI;
 
+typedef struct {
+  void *ptr;
+  uint32_t num_radix_blocks;
+  uint32_t lwe_dimension;
+} CudaLweCiphertextListFFI;
+
+typedef struct {
+  void *ptr;
+  uint32_t storage_log_modulus;
+  uint32_t lwe_per_glwe;
+  // Input LWEs are grouped by groups of `lwe_per_glwe`(the last group may be
+  // smaller)
+  // Each group is then packed into one GLWE with `lwe_per_glwe` bodies (one for
+  // each LWE of the group). In the end the total number of bodies is equal to
+  // the number of input LWE
+  uint32_t total_lwe_bodies_count;
+  uint32_t glwe_dimension;
+  uint32_t polynomial_size;
+} CudaPackedGlweCiphertextListFFI;
+
 uint64_t scratch_cuda_apply_univariate_lut_kb_64(
     CudaStreamsFFI streams, int8_t **mem_ptr, void const *input_lut,
     uint32_t lwe_dimension, uint32_t glwe_dimension, uint32_t polynomial_size,
@@ -333,8 +353,8 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
     uint32_t small_lwe_dimension, uint32_t ks_level, uint32_t ks_base_log,
     uint32_t pbs_level, uint32_t pbs_base_log, uint32_t grouping_factor,
     uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, uint32_t requested_flag, uint32_t uses_carry,
-    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type);
+    PBS_TYPE pbs_type, uint32_t requested_flag, bool allocate_gpu_memory,
+    PBS_MS_REDUCTION_T noise_reduction_type);
 
 uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     CudaStreamsFFI streams, int8_t **mem_ptr, uint32_t glwe_dimension,
@@ -342,8 +362,8 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     uint32_t small_lwe_dimension, uint32_t ks_level, uint32_t ks_base_log,
     uint32_t pbs_level, uint32_t pbs_base_log, uint32_t grouping_factor,
     uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, uint32_t requested_flag, uint32_t uses_carry,
-    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type);
+    PBS_TYPE pbs_type, uint32_t requested_flag, bool allocate_gpu_memory,
+    PBS_MS_REDUCTION_T noise_reduction_type);
 
 void cuda_propagate_single_carry_kb_64_inplace(
     CudaStreamsFFI streams, CudaRadixCiphertextFFI *lwe_array,
@@ -692,10 +712,9 @@ uint64_t scratch_cuda_integer_grouped_oprf_64(
     uint32_t polynomial_size, uint32_t lwe_dimension, uint32_t ks_level,
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks_to_process,
-    uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, bool allocate_gpu_memory,
-    uint32_t message_bits_per_block, uint32_t total_random_bits,
-    PBS_MS_REDUCTION_T noise_reduction_type);
+    uint32_t message_modulus, uint32_t carry_modulus, PBS_TYPE pbs_type,
+    bool allocate_gpu_memory, uint32_t message_bits_per_block,
+    uint32_t total_random_bits, PBS_MS_REDUCTION_T noise_reduction_type);
 
 void cuda_integer_grouped_oprf_async_64(
     CudaStreamsFFI streams, CudaRadixCiphertextFFI *radix_lwe_out,

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

@@ -36,7 +36,7 @@ __device__ Torus *get_ith_block(Torus *ksk, int i, int level,
  *
  */
 // Each thread in x are used to calculate one output.
-// threads in y are used to paralelize the lwe_dimension_in loop.
+// threads in y are used to parallelize the lwe_dimension_in loop.
 // shared memory is used to store intermediate results of the reduction.
 // Note: To reduce register pressure we have slightly changed the algorithm,
 // the idea consists in calculating the negate value of the output. So, instead

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

@@ -2,6 +2,9 @@
 #include <cstdint>
 #include <cuda_runtime.h>
 #include <mutex>
+#ifdef USE_NVTOOLS
+#include <cuda_profiler_api.h>
+#endif
 
 uint32_t cuda_get_device() {
   int device;
@@ -83,6 +86,9 @@ void cuda_set_device(uint32_t gpu_index) {
   check_cuda_error(cudaSetDevice(gpu_index));
   // Mempools are initialized only once in all the GPUS available
   cuda_setup_mempool(gpu_index);
+#ifdef USE_NVTOOLS
+  check_cuda_error(cudaProfilerStart());
+#endif
 }
 
 cudaEvent_t cuda_create_event(uint32_t gpu_index) {

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

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

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

@@ -4,6 +4,7 @@
 #include "crypto/keyswitch.cuh"
 #include "device.h"
 #include "integer/abs.cuh"
+#include "integer/cast.cuh"
 #include "integer/comparison.cuh"
 #include "integer/integer.cuh"
 #include "integer/integer_utilities.h"
@@ -31,6 +32,482 @@ __host__ uint64_t scratch_cuda_integer_div_rem_kb(
   return size_tracker;
 }
 
+template <typename Torus>
+__host__ void host_unsigned_integer_div_rem_kb_block_by_block_2_2(
+    CudaStreams streams, CudaRadixCiphertextFFI *quotient,
+    CudaRadixCiphertextFFI *remainder, CudaRadixCiphertextFFI const *numerator,
+    CudaRadixCiphertextFFI const *divisor, void *const *bsks,
+    uint64_t *const *ksks,
+    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key,
+    unsigned_int_div_rem_2_2_memory<uint64_t> *mem_ptr) {
+
+  if (streams.count() < 4) {
+    PANIC("GPU count should be greater than 4 when using div_rem_2_2");
+  }
+  if (mem_ptr->params.message_modulus != 4 ||
+      mem_ptr->params.carry_modulus != 4) {
+    PANIC("Only message_modulus == 4 && carry_modulus == 4 parameters are "
+          "supported");
+  }
+
+  // alias
+  auto radix_params = mem_ptr->params;
+  auto num_blocks = quotient->num_radix_blocks;
+  auto remainder_gpu_0 = remainder;
+  auto remainder_gpu_1 = mem_ptr->remainder_gpu_1;
+  auto remainder_gpu_2 = mem_ptr->remainder_gpu_2;
+  auto remainder_gpu_3 = mem_ptr->remainder_gpu_3;
+  auto divisor_gpu_0 = divisor;
+  auto divisor_gpu_1 = mem_ptr->divisor_gpu_1;
+  auto divisor_gpu_2 = mem_ptr->divisor_gpu_2;
+
+  auto make_view = [](CudaModulusSwitchNoiseReductionKeyFFI const *src,
+                      size_t i) {
+    CudaModulusSwitchNoiseReductionKeyFFI v;
+    v.ptr = (src == nullptr)        ? nullptr
+            : (src->ptr == nullptr) ? nullptr
+                                    : src->ptr + i;
+    v.num_zeros = (src == nullptr) ? 0 : src->num_zeros;
+    v.ms_bound = (src == nullptr) ? 0. : src->ms_bound;
+    v.ms_r_sigma = (src == nullptr) ? 0. : src->ms_r_sigma;
+    v.ms_input_variance = (src == nullptr) ? 0. : src->ms_input_variance;
+    return v;
+  };
+
+  CudaModulusSwitchNoiseReductionKeyFFI nrk0 =
+      make_view(ms_noise_reduction_key, 0);
+  CudaModulusSwitchNoiseReductionKeyFFI nrk1 =
+      make_view(ms_noise_reduction_key, 1);
+  CudaModulusSwitchNoiseReductionKeyFFI nrk2 =
+      make_view(ms_noise_reduction_key, 2);
+  CudaModulusSwitchNoiseReductionKeyFFI nrk3 =
+      make_view(ms_noise_reduction_key, 3);
+
+  CudaModulusSwitchNoiseReductionKeyFFI *ms_noise_reduction_keys[4] = {
+      &nrk0, &nrk1, &nrk2, &nrk3};
+
+  // gpu[0] -> gpu[0]
+  copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                     remainder_gpu_0, numerator);
+
+  // gpu[0] -> gpu[1]
+  copy_radix_ciphertext_async<Torus>(streams.stream(1), streams.gpu_index(1),
+                                     remainder_gpu_1, numerator);
+  // gpu[0] -> gpu[1]
+  copy_radix_ciphertext_async<Torus>(streams.stream(1), streams.gpu_index(1),
+                                     divisor_gpu_1, divisor);
+  // gpu[0] -> gpu[2]
+  copy_radix_ciphertext_async<Torus>(streams.stream(2), streams.gpu_index(2),
+                                     remainder_gpu_2, numerator);
+  // gpu[0] -> gpu[3]
+  copy_radix_ciphertext_async<Torus>(streams.stream(3), streams.gpu_index(3),
+                                     remainder_gpu_3, numerator);
+  // gpu[0] -> gpu[2]
+  copy_radix_ciphertext_async<Torus>(streams.stream(2), streams.gpu_index(2),
+                                     divisor_gpu_2, divisor);
+
+  // gpu[0]
+  set_zero_radix_ciphertext_slice_async<Torus>(
+      streams.stream(0), streams.gpu_index(0), quotient, 0, num_blocks);
+  quotient->num_radix_blocks = 0;
+
+  // Copy divisor_gpu_2 into d1 gpu[2] -> gpu[2]
+  mem_ptr->d1->num_radix_blocks = divisor_gpu_2->num_radix_blocks;
+  copy_radix_ciphertext_async<Torus>(streams.stream(2), streams.gpu_index(2),
+                                     mem_ptr->d1, divisor_gpu_2);
+
+  // Computes 2*d by extending and shifting on gpu[1]
+  host_extend_radix_with_trivial_zero_blocks_msb<Torus>(
+      mem_ptr->d2, divisor_gpu_1, streams.subset(1));
+  host_integer_radix_logical_scalar_shift_kb_inplace<Torus>(
+      streams.subset(1), mem_ptr->d2, 1, mem_ptr->shift_mem, &bsks[1], &ksks[1],
+      ms_noise_reduction_keys[1], mem_ptr->d2->num_radix_blocks);
+
+  // Computes 3*d = 4*d - d using block shift and subtraction on gpu[0]
+  host_extend_radix_with_trivial_zero_blocks_msb<Torus>(
+      mem_ptr->tmp_gpu_0, divisor_gpu_0, streams.subset(0));
+  host_radix_blocks_rotate_right<Torus>(streams.subset(0), mem_ptr->d3,
+                                        mem_ptr->tmp_gpu_0, 1,
+                                        mem_ptr->tmp_gpu_0->num_radix_blocks);
+  set_zero_radix_ciphertext_slice_async<Torus>(
+      streams.stream(0), streams.gpu_index(0), mem_ptr->d3, 0, 1);
+  host_sub_and_propagate_single_carry(
+      streams.subset(0), mem_ptr->d3, mem_ptr->tmp_gpu_0, nullptr, nullptr,
+      mem_ptr->sub_and_propagate_mem, &bsks[0], &ksks[0],
+      ms_noise_reduction_keys[0], outputFlag::FLAG_NONE, 0);
+
+  // +-----------------+-----------------+-----------------+-----------------+
+  // |     GPU[0]      |     GPU[1]      |     GPU[2]      |     GPU[3]      |
+  // +-----------------+-----------------+-----------------+-----------------+
+  // | d3              | d2              | d1              | -               |
+  // | low3            | low2            | low1            | -               |
+  // | rem3            | rem2            | rem1            | rem0            |
+  // | sub_result_1    | sub_result_2    | sub_result_3    | -               |
+  // | s_1_overflowed  | s_2_overflowed  | s_3_overflowed  | -               |
+  // | cmp_1           | cmp_2           | cmp_3           | -               |
+  // | r3              | r2              | r1              | -               |
+  // | o3              | o2              | o1              | -               |
+  // | c3 = !o3        | c2 = !o2 + o3   | c1 = !o1 + o2   | c0 = o1         |
+  // | z_o_not_1_lut_1 | z_o_not_2_lut_1 | z_o_not_2_lut_2 | z_o_not_1_lut_2 |
+  // +-----------------+-----------------+-----------------+-----------------+
+  for (int block_index = num_blocks - 1; block_index >= 0; block_index--) {
+
+    uint32_t slice_len = num_blocks - block_index;
+
+    auto init_low_rem_f =
+        [&](CudaRadixCiphertextFFI *low, CudaRadixCiphertextFFI *xd,
+            CudaRadixCiphertextFFI *rem, CudaRadixCiphertextFFI *cur_remainder,
+            size_t gpu_index, bool init_low) {
+          rem->num_radix_blocks = slice_len;
+          if (init_low) {
+            low->num_radix_blocks = slice_len;
+            copy_radix_ciphertext_slice_async<Torus>(
+                streams.stream(gpu_index), streams.gpu_index(gpu_index), low, 0,
+                slice_len, xd, 0, slice_len);
+          }
+          copy_radix_ciphertext_slice_async<Torus>(
+              streams.stream(gpu_index), streams.gpu_index(gpu_index), rem, 0,
+              slice_len, cur_remainder, block_index, num_blocks);
+        };
+
+    init_low_rem_f(nullptr, nullptr, mem_ptr->rem0, remainder_gpu_3, 3, false);
+    init_low_rem_f(mem_ptr->low1, mem_ptr->d1, mem_ptr->rem1, remainder_gpu_2,
+                   2, true);
+    init_low_rem_f(mem_ptr->low2, mem_ptr->d2, mem_ptr->rem2, remainder_gpu_1,
+                   1, true);
+    init_low_rem_f(mem_ptr->low3, mem_ptr->d3, mem_ptr->rem3, remainder_gpu_0,
+                   0, true);
+
+    auto sub_result_f = [&](CudaStreams streams, size_t gpu_index,
+                            CudaRadixCiphertextFFI *sub_result,
+                            CudaRadixCiphertextFFI *sub_overflowed,
+                            int_borrow_prop_memory<Torus> *overflow_sub_mem,
+                            CudaRadixCiphertextFFI *low,
+                            CudaRadixCiphertextFFI *rem, Torus *first_indexes,
+                            Torus *second_indexes, Torus *scalar_indexes) {
+      uint32_t compute_overflow = 1;
+      uint32_t uses_input_borrow = 0;
+      sub_result->num_radix_blocks = low->num_radix_blocks;
+      overflow_sub_mem->update_lut_indexes(
+          streams.subset(gpu_index), first_indexes, second_indexes,
+          scalar_indexes, rem->num_radix_blocks);
+      host_integer_overflowing_sub<uint64_t>(
+          streams.subset(gpu_index), sub_result, rem, low, sub_overflowed,
+          (const CudaRadixCiphertextFFI *)nullptr, overflow_sub_mem,
+          &bsks[gpu_index], &ksks[gpu_index],
+          ms_noise_reduction_keys[gpu_index], compute_overflow,
+          uses_input_borrow);
+    };
+
+    auto cmp_f = [&](CudaStreams streams, size_t gpu_index,
+                     CudaRadixCiphertextFFI *out_boolean_block,
+                     CudaRadixCiphertextFFI *comparison_blocks,
+                     CudaRadixCiphertextFFI *d,
+                     int_comparison_buffer<Torus> *comparison_buffer) {
+      CudaRadixCiphertextFFI *d_msb = new CudaRadixCiphertextFFI;
+      uint32_t slice_start = num_blocks - block_index;
+      uint32_t slice_end = d->num_radix_blocks;
+      as_radix_ciphertext_slice<Torus>(d_msb, d, slice_start, slice_end);
+      comparison_blocks->num_radix_blocks = d_msb->num_radix_blocks;
+      if (d_msb->num_radix_blocks == 0) {
+        cuda_memset_async(
+            (Torus *)out_boolean_block->ptr, 0,
+            sizeof(Torus) * (out_boolean_block->lwe_dimension + 1),
+            streams.stream(gpu_index), streams.gpu_index(gpu_index));
+      } else {
+        host_compare_blocks_with_zero<Torus>(
+            streams.subset(gpu_index), comparison_blocks, d_msb,
+            comparison_buffer, &bsks[gpu_index], &ksks[gpu_index],
+            ms_noise_reduction_keys[gpu_index], d_msb->num_radix_blocks,
+            comparison_buffer->is_zero_lut);
+        are_all_comparisons_block_true(
+            streams.subset(gpu_index), out_boolean_block, comparison_blocks,
+            comparison_buffer, &bsks[gpu_index], &ksks[gpu_index],
+            ms_noise_reduction_keys[gpu_index],
+            comparison_blocks->num_radix_blocks);
+
+        host_negation<Torus>(
+            streams.stream(gpu_index), streams.gpu_index(gpu_index),
+            (Torus *)out_boolean_block->ptr, (Torus *)out_boolean_block->ptr,
+            radix_params.big_lwe_dimension, 1);
+
+        // we calculate encoding because this block works only for
+        // message_modulus = 4 and carry_modulus = 4.
+        const Torus encoded_scalar = 1ULL << (sizeof(Torus) * 8 - 5);
+        host_addition_plaintext_scalar<Torus>(
+            streams.stream(gpu_index), streams.gpu_index(gpu_index),
+            (Torus *)out_boolean_block->ptr, (Torus *)out_boolean_block->ptr,
+            encoded_scalar, radix_params.big_lwe_dimension, 1);
+      }
+      delete d_msb;
+    };
+
+    for (uint j = 0; j < 3; j++) {
+      cuda_synchronize_stream(streams.stream(j), streams.gpu_index(j));
+    }
+
+    size_t indexes_id = mem_ptr->rem3->num_radix_blocks - 1;
+    sub_result_f(streams, 0, mem_ptr->sub_result_1, mem_ptr->sub_1_overflowed,
+                 mem_ptr->overflow_sub_mem_1, mem_ptr->low3, mem_ptr->rem3,
+                 mem_ptr->first_indexes_for_overflow_sub_gpu_0[indexes_id],
+                 mem_ptr->second_indexes_for_overflow_sub_gpu_0[indexes_id],
+                 mem_ptr->scalars_for_overflow_sub_gpu_0[indexes_id]);
+    sub_result_f(streams, 1, mem_ptr->sub_result_2, mem_ptr->sub_2_overflowed,
+                 mem_ptr->overflow_sub_mem_2, mem_ptr->low2, mem_ptr->rem2,
+                 mem_ptr->first_indexes_for_overflow_sub_gpu_1[indexes_id],
+                 mem_ptr->second_indexes_for_overflow_sub_gpu_1[indexes_id],
+                 mem_ptr->scalars_for_overflow_sub_gpu_1[indexes_id]);
+    sub_result_f(streams, 2, mem_ptr->sub_result_3, mem_ptr->sub_3_overflowed,
+                 mem_ptr->overflow_sub_mem_3, mem_ptr->low1, mem_ptr->rem1,
+                 mem_ptr->first_indexes_for_overflow_sub_gpu_2[indexes_id],
+                 mem_ptr->second_indexes_for_overflow_sub_gpu_2[indexes_id],
+                 mem_ptr->scalars_for_overflow_sub_gpu_2[indexes_id]);
+
+    cmp_f(mem_ptr->sub_streams_1, 0, mem_ptr->cmp_1,
+          mem_ptr->comparison_blocks_1, mem_ptr->d3,
+          mem_ptr->comparison_buffer_1);
+    cmp_f(mem_ptr->sub_streams_1, 1, mem_ptr->cmp_2,
+          mem_ptr->comparison_blocks_2, mem_ptr->d2,
+          mem_ptr->comparison_buffer_2);
+    cmp_f(mem_ptr->sub_streams_1, 2, mem_ptr->cmp_3,
+          mem_ptr->comparison_blocks_3, mem_ptr->d1,
+          mem_ptr->comparison_buffer_3);
+
+    for (uint j = 0; j < 3; j++) {
+      cuda_synchronize_stream(streams.stream(j), streams.gpu_index(j));
+      cuda_synchronize_stream(mem_ptr->sub_streams_1.stream(j),
+                              mem_ptr->sub_streams_1.gpu_index(j));
+    }
+
+    auto r1 = mem_ptr->sub_result_3;
+    auto r2 = mem_ptr->sub_result_2;
+    auto r3 = mem_ptr->sub_result_1;
+    auto o1 = mem_ptr->sub_3_overflowed;
+    auto o2 = mem_ptr->sub_2_overflowed;
+    auto o3 = mem_ptr->sub_1_overflowed;
+
+    // used as a bitor
+    host_integer_radix_bitop_kb(streams.subset(0), o3, o3, mem_ptr->cmp_1,
+                                mem_ptr->bitor_mem_1, &bsks[0], &ksks[0],
+                                ms_noise_reduction_keys[0]);
+    // used as a bitor
+    host_integer_radix_bitop_kb(streams.subset(1), o2, o2, mem_ptr->cmp_2,
+                                mem_ptr->bitor_mem_2, &bsks[1], &ksks[1],
+                                ms_noise_reduction_keys[1]);
+    // used as a bitor
+    host_integer_radix_bitop_kb(streams.subset(2), o1, o1, mem_ptr->cmp_3,
+                                mem_ptr->bitor_mem_3, &bsks[2], &ksks[2],
+                                ms_noise_reduction_keys[2]);
+
+    // cmp_1, cmp_2, cmp_3 are not needed anymore, we can reuse them as c3,
+    // c2, c1. c0 is allocated on gpu[3], we take it from mem_ptr.
+    auto c3 = mem_ptr->cmp_1;
+    auto c2 = mem_ptr->cmp_2;
+    auto c1 = mem_ptr->cmp_3;
+    auto c0 = mem_ptr->c0;
+
+    // move all `o` so that each gpu has required `o` for calculating `c`
+    auto o3_gpu_1 = mem_ptr->tmp_gpu_1;
+    auto o2_gpu_2 = mem_ptr->tmp_gpu_2;
+    auto o1_gpu_3 = mem_ptr->tmp_gpu_3;
+
+    o3_gpu_1->num_radix_blocks = o3->num_radix_blocks;
+    o2_gpu_2->num_radix_blocks = o2->num_radix_blocks;
+    o1_gpu_3->num_radix_blocks = o1->num_radix_blocks;
+
+    for (uint j = 0; j < 4; j++) {
+      cuda_synchronize_stream(streams.stream(j), streams.gpu_index(j));
+    }
+
+    copy_radix_ciphertext_async<Torus>(streams.stream(1), streams.gpu_index(1),
+                                       o3_gpu_1, o3);
+    copy_radix_ciphertext_async<Torus>(streams.stream(2), streams.gpu_index(2),
+                                       o2_gpu_2, o2);
+    copy_radix_ciphertext_async<Torus>(streams.stream(3), streams.gpu_index(3),
+                                       o1_gpu_3, o1);
+
+    // c3 = !o3
+    copy_radix_ciphertext_slice_async<Torus>(
+        streams.stream(0), streams.gpu_index(0), c3, 0, 1, o3, 0, 1);
+    host_negation<Torus>(streams.stream(0), streams.gpu_index(0),
+                         (Torus *)c3->ptr, (Torus *)c3->ptr,
+                         radix_params.big_lwe_dimension, 1);
+    const Torus encoded_scalar = 1ULL << (sizeof(Torus) * 8 - 5);
+    host_addition_plaintext_scalar<Torus>(
+        streams.stream(0), streams.gpu_index(0), (Torus *)c3->ptr,
+        (Torus *)c3->ptr, encoded_scalar, radix_params.big_lwe_dimension, 1);
+
+    // c2 = !o2 + o3
+    copy_radix_ciphertext_slice_async<Torus>(
+        streams.stream(1), streams.gpu_index(1), c2, 0, 1, o2, 0, 1);
+    host_negation<Torus>(streams.stream(1), streams.gpu_index(1),
+                         (Torus *)c2->ptr, (Torus *)c2->ptr,
+                         radix_params.big_lwe_dimension, 1);
+    host_addition_plaintext_scalar<Torus>(
+        streams.stream(1), streams.gpu_index(1), (Torus *)c2->ptr,
+        (Torus *)c2->ptr, encoded_scalar, radix_params.big_lwe_dimension, 1);
+    host_addition<Torus>(streams.stream(1), streams.gpu_index(1), c2, c2,
+                         o3_gpu_1, 1, 4, 4);
+
+    // c1 = !o1 + o2
+    copy_radix_ciphertext_slice_async<Torus>(
+        streams.stream(2), streams.gpu_index(2), c1, 0, 1, o1, 0, 1);
+    host_negation<Torus>(streams.stream(2), streams.gpu_index(2),
+                         (Torus *)c1->ptr, (Torus *)c1->ptr,
+                         radix_params.big_lwe_dimension, 1);
+    host_addition_plaintext_scalar<Torus>(
+        streams.stream(2), streams.gpu_index(2), (Torus *)c1->ptr,
+        (Torus *)c1->ptr, encoded_scalar, radix_params.big_lwe_dimension, 1);
+    host_addition<Torus>(streams.stream(2), streams.gpu_index(2), c1, c1,
+                         o2_gpu_2, 1, 4, 4);
+
+    // c0 = o1 (direct copy)
+    copy_radix_ciphertext_slice_async<Torus>(streams.stream(3),
+                                             streams.gpu_index(3), mem_ptr->c0,
+                                             0, 1, o1_gpu_3, 0, 1);
+
+    auto conditional_update = [&](CudaStreams streams, size_t gpu_index,
+                                  CudaRadixCiphertextFFI *cx,
+                                  CudaRadixCiphertextFFI *rx,
+                                  int_radix_lut<Torus> *lut, Torus factor) {
+      auto rx_list = to_lwe_ciphertext_list(rx);
+      host_cleartext_multiplication<Torus>(streams.stream(gpu_index),
+                                           streams.gpu_index(gpu_index),
+                                           (Torus *)rx->ptr, &rx_list, factor);
+      host_add_the_same_block_to_all_blocks<Torus>(streams.stream(gpu_index),
+                                                   streams.gpu_index(gpu_index),
+                                                   rx, rx, cx, 4, 4);
+      integer_radix_apply_univariate_lookup_table_kb<Torus>(
+          streams.subset(gpu_index), rx, rx, &bsks[gpu_index], &ksks[gpu_index],
+          ms_noise_reduction_keys[gpu_index], lut, rx->num_radix_blocks);
+    };
+
+    for (uint j = 0; j < 4; j++) {
+      cuda_synchronize_stream(streams.stream(j), streams.gpu_index(j));
+      cuda_synchronize_stream(mem_ptr->sub_streams_1.stream(j),
+                              mem_ptr->sub_streams_1.gpu_index(j));
+    }
+
+    conditional_update(streams, 0, c3, r3, mem_ptr->zero_out_if_not_1_lut_1, 2);
+    conditional_update(streams, 1, c2, r2, mem_ptr->zero_out_if_not_2_lut_1, 3);
+    conditional_update(streams, 2, c1, r1, mem_ptr->zero_out_if_not_2_lut_2, 3);
+    conditional_update(streams, 3, c0, mem_ptr->rem0,
+                       mem_ptr->zero_out_if_not_1_lut_2, 2);
+
+    // calculate quotient bits GPU[2]
+    integer_radix_apply_univariate_lookup_table_kb<Torus>(
+        mem_ptr->sub_streams_1.subset(2), mem_ptr->q1, c1, &bsks[2], &ksks[2],
+        ms_noise_reduction_keys[2], mem_ptr->quotient_lut_1, 1);
+    // calculate quotient bits GPU[1]
+    integer_radix_apply_univariate_lookup_table_kb<Torus>(
+        mem_ptr->sub_streams_1.subset(1), mem_ptr->q2, c2, &bsks[1], &ksks[1],
+        ms_noise_reduction_keys[1], mem_ptr->quotient_lut_2, 1);
+    // calculate quotient bits GPU[0]
+    integer_radix_apply_univariate_lookup_table_kb<Torus>(
+        mem_ptr->sub_streams_1.subset(0), mem_ptr->q3, c3, &bsks[0], &ksks[0],
+        ms_noise_reduction_keys[0], mem_ptr->quotient_lut_3, 1);
+
+    for (uint j = 0; j < 4; j++) {
+      cuda_synchronize_stream(streams.stream(j), streams.gpu_index(j));
+      cuda_synchronize_stream(mem_ptr->sub_streams_1.stream(j),
+                              mem_ptr->sub_streams_1.gpu_index(j));
+    }
+
+    // We need to accumulate rem, r1, r2, and r3, but each buffer currently
+    // lives on a different GPU. To gather them on GPU[0], we’ll **reuse**
+    // buffers already allocated on GPU[0]. At this point, the contents of rem3,
+    // tmp_gpu_0, and low3 are no longer needed, so it’s safe to repurpose them.
+    // Aliases for the GPU[0] destinations:
+    auto r3_gpu_0 = r3;                 // reuse: destination for r3 on GPU[0]
+    auto r2_gpu_0 = mem_ptr->tmp_gpu_0; // reuse: destination for r2 on GPU[0]
+    auto r1_gpu_0 = mem_ptr->low3;      // reuse: destination for r1 on GPU[0]
+    auto rem_gpu_0 = mem_ptr->rem3;     // reuse: destination for rem on GPU[0]
+
+    r2_gpu_0->num_radix_blocks = r2->num_radix_blocks;
+    // r3 is already on GPU 0, so no need to copy it.
+
+    // Copy r2 from GPU[1] to GPU[0]
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       r2_gpu_0, r2);
+
+    // Copy r1 from GPU[2] to GPU[0]
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       r1_gpu_0, r1);
+
+    // Copy rem from GPU[3] to GPU[0]
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       rem_gpu_0, mem_ptr->rem0);
+
+    // We do the same to accumulate quotient bits q1, q2 and q3. q3 is already
+    // on GPU[0]. To copy q1 and q2 we will reuse buffers allocated on GPU[0]:
+    // sub_1_overflowed and cmp_1.
+    auto q3_gpu_0 = mem_ptr->q3; // q3 is already on GPU[0]
+    auto q2_gpu_0 =
+        mem_ptr->sub_1_overflowed;  // reuse: destination for q2 on GPU[0]
+    auto q1_gpu_0 = mem_ptr->cmp_1; // reuse: destination for q1 on GPU[0]
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       q2_gpu_0, mem_ptr->q2);
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       q1_gpu_0, mem_ptr->q1);
+
+    host_addition<Torus>(streams.stream(0), streams.gpu_index(0), rem_gpu_0,
+                         rem_gpu_0, r3_gpu_0, rem_gpu_0->num_radix_blocks, 4,
+                         4);
+    host_addition<Torus>(streams.stream(0), streams.gpu_index(0), rem_gpu_0,
+                         rem_gpu_0, r2_gpu_0, rem_gpu_0->num_radix_blocks, 4,
+                         4);
+    host_addition<Torus>(streams.stream(0), streams.gpu_index(0), rem_gpu_0,
+                         rem_gpu_0, r1_gpu_0, rem_gpu_0->num_radix_blocks, 4,
+                         4);
+
+    host_addition<Torus>(streams.stream(0), streams.gpu_index(0), q3_gpu_0,
+                         q3_gpu_0, q2_gpu_0, 1, 4, 4);
+    host_addition<Torus>(streams.stream(0), streams.gpu_index(0), q3_gpu_0,
+                         q3_gpu_0, q1_gpu_0, 1, 4, 4);
+
+    streams.synchronize();
+
+    integer_radix_apply_univariate_lookup_table_kb<Torus>(
+        streams, rem_gpu_0, rem_gpu_0, bsks, ksks, ms_noise_reduction_key,
+        mem_ptr->message_extract_lut_1, rem_gpu_0->num_radix_blocks);
+    integer_radix_apply_univariate_lookup_table_kb<Torus>(
+        mem_ptr->sub_streams_1, q3_gpu_0, q3_gpu_0, bsks, ksks,
+        ms_noise_reduction_key, mem_ptr->message_extract_lut_2, 1);
+    streams.synchronize();
+    mem_ptr->sub_streams_1.synchronize();
+
+    copy_radix_ciphertext_slice_async<Torus>(
+        streams.stream(0), streams.gpu_index(0), remainder_gpu_0, block_index,
+        remainder_gpu_0->num_radix_blocks, rem_gpu_0, 0,
+        rem_gpu_0->num_radix_blocks);
+    insert_block_in_radix_ciphertext_async<Torus>(
+        streams.stream(0), streams.gpu_index(0), q3_gpu_0, quotient, 0);
+
+    // Copy remainder_gpu_0 to all other GPUs
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       remainder_gpu_1, remainder_gpu_0);
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       remainder_gpu_2, remainder_gpu_0);
+    copy_radix_ciphertext_async<Torus>(streams.stream(0), streams.gpu_index(0),
+                                       remainder_gpu_3, remainder_gpu_0);
+
+    // non boolean blocks
+    for (int block_id = 0; block_id < slice_len; block_id++) {
+      mem_ptr->sub_result_1->degrees[block_id] =
+          radix_params.message_modulus - 1;
+      mem_ptr->rem0->degrees[block_id] = radix_params.message_modulus - 1;
+    }
+
+    // boolean blocks
+    mem_ptr->cmp_3->degrees[0] = 0;
+    mem_ptr->cmp_2->degrees[0] = 0;
+    mem_ptr->cmp_1->degrees[0] = 0;
+    mem_ptr->cmp_3->noise_levels[0] = 0;
+
+    streams.synchronize();
+  }
+}
+
 template <typename Torus>
 __host__ void host_unsigned_integer_div_rem_kb(
     CudaStreams streams, CudaRadixCiphertextFFI *quotient,
@@ -48,6 +525,14 @@ __host__ void host_unsigned_integer_div_rem_kb(
       remainder->lwe_dimension != divisor->lwe_dimension ||
       remainder->lwe_dimension != quotient->lwe_dimension)
     PANIC("Cuda error: input and output lwe dimension must be equal")
+
+  if (mem_ptr->params.message_modulus == 4 &&
+      mem_ptr->params.carry_modulus == 4 && streams.count() >= 4) {
+    host_unsigned_integer_div_rem_kb_block_by_block_2_2<Torus>(
+        streams, quotient, remainder, numerator, divisor, bsks, ksks,
+        ms_noise_reduction_key, mem_ptr->div_rem_2_2_mem);
+    return;
+  }
   auto radix_params = mem_ptr->params;
   auto num_blocks = quotient->num_radix_blocks;
 

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

@@ -51,8 +51,8 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
     uint32_t small_lwe_dimension, uint32_t ks_level, uint32_t ks_base_log,
     uint32_t pbs_level, uint32_t pbs_base_log, uint32_t grouping_factor,
     uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, uint32_t requested_flag, uint32_t uses_carry,
-    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type) {
+    PBS_TYPE pbs_type, uint32_t requested_flag, bool allocate_gpu_memory,
+    PBS_MS_REDUCTION_T noise_reduction_type) {
   int_radix_params params(pbs_type, glwe_dimension, polynomial_size,
                           big_lwe_dimension, small_lwe_dimension, ks_level,
                           ks_base_log, pbs_level, pbs_base_log, grouping_factor,
@@ -60,7 +60,7 @@ uint64_t scratch_cuda_propagate_single_carry_kb_64_inplace(
 
   return scratch_cuda_propagate_single_carry_kb_inplace<uint64_t>(
       CudaStreams(streams), (int_sc_prop_memory<uint64_t> **)mem_ptr,
-      num_blocks, params, requested_flag, uses_carry, allocate_gpu_memory);
+      num_blocks, params, requested_flag, allocate_gpu_memory);
 }
 
 uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
@@ -69,8 +69,8 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
     uint32_t small_lwe_dimension, uint32_t ks_level, uint32_t ks_base_log,
     uint32_t pbs_level, uint32_t pbs_base_log, uint32_t grouping_factor,
     uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, uint32_t requested_flag, uint32_t uses_carry,
-    bool allocate_gpu_memory, PBS_MS_REDUCTION_T noise_reduction_type) {
+    PBS_TYPE pbs_type, uint32_t requested_flag, bool allocate_gpu_memory,
+    PBS_MS_REDUCTION_T noise_reduction_type) {
   int_radix_params params(pbs_type, glwe_dimension, polynomial_size,
                           big_lwe_dimension, small_lwe_dimension, ks_level,
                           ks_base_log, pbs_level, pbs_base_log, grouping_factor,
@@ -78,7 +78,7 @@ uint64_t scratch_cuda_add_and_propagate_single_carry_kb_64_inplace(
 
   return scratch_cuda_propagate_single_carry_kb_inplace<uint64_t>(
       CudaStreams(streams), (int_sc_prop_memory<uint64_t> **)mem_ptr,
-      num_blocks, params, requested_flag, uses_carry, allocate_gpu_memory);
+      num_blocks, params, requested_flag, allocate_gpu_memory);
 }
 
 uint64_t scratch_cuda_integer_overflowing_sub_kb_64_inplace(

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

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

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

@@ -5,10 +5,9 @@ uint64_t scratch_cuda_integer_grouped_oprf_64(
     uint32_t polynomial_size, uint32_t lwe_dimension, uint32_t ks_level,
     uint32_t ks_base_log, uint32_t pbs_level, uint32_t pbs_base_log,
     uint32_t grouping_factor, uint32_t num_blocks_to_process,
-    uint32_t num_blocks, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, bool allocate_gpu_memory,
-    uint32_t message_bits_per_block, uint32_t total_random_bits,
-    PBS_MS_REDUCTION_T noise_reduction_type) {
+    uint32_t message_modulus, uint32_t carry_modulus, PBS_TYPE pbs_type,
+    bool allocate_gpu_memory, uint32_t message_bits_per_block,
+    uint32_t total_random_bits, PBS_MS_REDUCTION_T noise_reduction_type) {
 
   int_radix_params params(pbs_type, glwe_dimension, polynomial_size,
                           glwe_dimension * polynomial_size, lwe_dimension,
@@ -18,8 +17,8 @@ uint64_t scratch_cuda_integer_grouped_oprf_64(
 
   return scratch_cuda_integer_grouped_oprf<uint64_t>(
       CudaStreams(streams), (int_grouped_oprf_memory<uint64_t> **)mem_ptr,
-      params, num_blocks_to_process, num_blocks, message_bits_per_block,
-      total_random_bits, allocate_gpu_memory);
+      params, num_blocks_to_process, message_bits_per_block, total_random_bits,
+      allocate_gpu_memory);
 }
 
 void cuda_integer_grouped_oprf_async_64(

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

@@ -8,14 +8,13 @@ template <typename Torus>
 uint64_t scratch_cuda_integer_grouped_oprf(
     CudaStreams streams, int_grouped_oprf_memory<Torus> **mem_ptr,
     int_radix_params params, uint32_t num_blocks_to_process,
-    uint32_t num_blocks, uint32_t message_bits_per_block,
-    uint64_t total_random_bits, bool allocate_gpu_memory) {
+    uint32_t message_bits_per_block, uint64_t total_random_bits,
+    bool allocate_gpu_memory) {
   uint64_t size_tracker = 0;
 
   *mem_ptr = new int_grouped_oprf_memory<Torus>(
-      streams, params, num_blocks_to_process, num_blocks,
-      message_bits_per_block, total_random_bits, allocate_gpu_memory,
-      size_tracker);
+      streams, params, num_blocks_to_process, message_bits_per_block,
+      total_random_bits, allocate_gpu_memory, size_tracker);
 
   return size_tracker;
 }
@@ -32,8 +31,8 @@ void host_integer_grouped_oprf(
 
   if (active_streams.count() == 1) {
     execute_pbs_async<Torus, Torus>(
-        streams, (Torus *)(radix_lwe_out->ptr), lut->lwe_indexes_out,
-        lut->lut_vec, lut->lut_indexes_vec,
+        streams.subset_first_gpu(), (Torus *)(radix_lwe_out->ptr),
+        lut->lwe_indexes_out, lut->lut_vec, lut->lut_indexes_vec,
         const_cast<Torus *>(seeded_lwe_input), lut->lwe_indexes_in, bsks,
         ms_noise_reduction_key, lut->buffer, mem_ptr->params.glwe_dimension,
         mem_ptr->params.small_lwe_dimension, mem_ptr->params.polynomial_size,
@@ -52,15 +51,13 @@ void host_integer_grouped_oprf(
                              streams.gpu_index(j));
     }
 
-    if (!lut->using_trivial_lwe_indexes) {
-      PANIC("lut->using_trivial_lwe_indexes should be true");
-    }
-
+    PUSH_RANGE("scatter")
     multi_gpu_scatter_lwe_async<Torus>(
         active_streams, lwe_array_in_vec, seeded_lwe_input, lut->lwe_indexes_in,
         lut->using_trivial_lwe_indexes, lut->lwe_aligned_vec,
         active_streams.count(), num_blocks_to_process,
         mem_ptr->params.small_lwe_dimension + 1);
+    POP_RANGE()
 
     execute_pbs_async<Torus, Torus>(
         active_streams, lwe_after_pbs_vec, lwe_trivial_indexes_vec,
@@ -71,12 +68,13 @@ void host_integer_grouped_oprf(
         mem_ptr->params.pbs_level, mem_ptr->params.grouping_factor,
         num_blocks_to_process, mem_ptr->params.pbs_type, 1, 0);
 
+    PUSH_RANGE("gather")
     multi_gpu_gather_lwe_async<Torus>(
         active_streams, (Torus *)radix_lwe_out->ptr, lwe_after_pbs_vec,
         lut->lwe_indexes_out, lut->using_trivial_lwe_indexes,
         lut->lwe_aligned_vec, num_blocks_to_process,
         mem_ptr->params.big_lwe_dimension + 1);
-
+    POP_RANGE()
     // other gpus record their events
     for (int j = 1; j < active_streams.count(); j++) {
       cuda_event_record(lut->event_scatter_out[j], streams.stream(j),

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

@@ -7,6 +7,13 @@
 #include "utils/helper_profile.cuh"
 #include "utils/kernel_dimensions.cuh"
 
+inline CudaLweCiphertextListFFI
+to_lwe_ciphertext_list(CudaRadixCiphertextFFI *radix) {
+  return {.ptr = radix->ptr,
+          .num_radix_blocks = radix->num_radix_blocks,
+          .lwe_dimension = radix->lwe_dimension};
+}
+
 template <typename Torus>
 void create_zero_radix_ciphertext_async(cudaStream_t const stream,
                                         uint32_t const gpu_index,

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

@@ -281,6 +281,55 @@ const _: () = {
         divisor_has_more_bits_than_numerator
     ) - 60usize];
 };
+#[repr(C)]
+#[derive(Debug, Copy, Clone)]
+pub struct CudaLweCiphertextListFFI {
+    pub ptr: *mut ffi::c_void,
+    pub num_radix_blocks: u32,
+    pub lwe_dimension: u32,
+}
+#[allow(clippy::unnecessary_operation, clippy::identity_op)]
+const _: () = {
+    ["Size of CudaLweCiphertextListFFI"]
+        [::std::mem::size_of::<CudaLweCiphertextListFFI>() - 16usize];
+    ["Alignment of CudaLweCiphertextListFFI"]
+        [::std::mem::align_of::<CudaLweCiphertextListFFI>() - 8usize];
+    ["Offset of field: CudaLweCiphertextListFFI::ptr"]
+        [::std::mem::offset_of!(CudaLweCiphertextListFFI, ptr) - 0usize];
+    ["Offset of field: CudaLweCiphertextListFFI::num_radix_blocks"]
+        [::std::mem::offset_of!(CudaLweCiphertextListFFI, num_radix_blocks) - 8usize];
+    ["Offset of field: CudaLweCiphertextListFFI::lwe_dimension"]
+        [::std::mem::offset_of!(CudaLweCiphertextListFFI, lwe_dimension) - 12usize];
+};
+#[repr(C)]
+#[derive(Debug, Copy, Clone)]
+pub struct CudaPackedGlweCiphertextListFFI {
+    pub ptr: *mut ffi::c_void,
+    pub storage_log_modulus: u32,
+    pub lwe_per_glwe: u32,
+    pub total_lwe_bodies_count: u32,
+    pub glwe_dimension: u32,
+    pub polynomial_size: u32,
+}
+#[allow(clippy::unnecessary_operation, clippy::identity_op)]
+const _: () = {
+    ["Size of CudaPackedGlweCiphertextListFFI"]
+        [::std::mem::size_of::<CudaPackedGlweCiphertextListFFI>() - 32usize];
+    ["Alignment of CudaPackedGlweCiphertextListFFI"]
+        [::std::mem::align_of::<CudaPackedGlweCiphertextListFFI>() - 8usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::ptr"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, ptr) - 0usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::storage_log_modulus"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, storage_log_modulus) - 8usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::lwe_per_glwe"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, lwe_per_glwe) - 12usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::total_lwe_bodies_count"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, total_lwe_bodies_count) - 16usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::glwe_dimension"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, glwe_dimension) - 20usize];
+    ["Offset of field: CudaPackedGlweCiphertextListFFI::polynomial_size"]
+        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, polynomial_size) - 24usize];
+};
 unsafe extern "C" {
     pub fn scratch_cuda_apply_univariate_lut_kb_64(
         streams: CudaStreamsFFI,
@@ -808,7 +857,6 @@ unsafe extern "C" {
         carry_modulus: u32,
         pbs_type: PBS_TYPE,
         requested_flag: u32,
-        uses_carry: u32,
         allocate_gpu_memory: bool,
         noise_reduction_type: PBS_MS_REDUCTION_T,
     ) -> u64;
@@ -831,7 +879,6 @@ unsafe extern "C" {
         carry_modulus: u32,
         pbs_type: PBS_TYPE,
         requested_flag: u32,
-        uses_carry: u32,
         allocate_gpu_memory: bool,
         noise_reduction_type: PBS_MS_REDUCTION_T,
     ) -> u64;
@@ -1553,7 +1600,6 @@ unsafe extern "C" {
         pbs_base_log: u32,
         grouping_factor: u32,
         num_blocks_to_process: u32,
-        num_blocks: u32,
         message_modulus: u32,
         carry_modulus: u32,
         pbs_type: PBS_TYPE,
@@ -1619,55 +1665,6 @@ unsafe extern "C" {
 unsafe extern "C" {
     pub fn cleanup_cuda_integer_ilog2_kb_64(streams: CudaStreamsFFI, mem_ptr_void: *mut *mut i8);
 }
-#[repr(C)]
-#[derive(Debug, Copy, Clone)]
-pub struct CudaLweCiphertextListFFI {
-    pub ptr: *mut ffi::c_void,
-    pub num_radix_blocks: u32,
-    pub lwe_dimension: u32,
-}
-#[allow(clippy::unnecessary_operation, clippy::identity_op)]
-const _: () = {
-    ["Size of CudaLweCiphertextListFFI"]
-        [::std::mem::size_of::<CudaLweCiphertextListFFI>() - 16usize];
-    ["Alignment of CudaLweCiphertextListFFI"]
-        [::std::mem::align_of::<CudaLweCiphertextListFFI>() - 8usize];
-    ["Offset of field: CudaLweCiphertextListFFI::ptr"]
-        [::std::mem::offset_of!(CudaLweCiphertextListFFI, ptr) - 0usize];
-    ["Offset of field: CudaLweCiphertextListFFI::num_radix_blocks"]
-        [::std::mem::offset_of!(CudaLweCiphertextListFFI, num_radix_blocks) - 8usize];
-    ["Offset of field: CudaLweCiphertextListFFI::lwe_dimension"]
-        [::std::mem::offset_of!(CudaLweCiphertextListFFI, lwe_dimension) - 12usize];
-};
-#[repr(C)]
-#[derive(Debug, Copy, Clone)]
-pub struct CudaPackedGlweCiphertextListFFI {
-    pub ptr: *mut ffi::c_void,
-    pub storage_log_modulus: u32,
-    pub lwe_per_glwe: u32,
-    pub total_lwe_bodies_count: u32,
-    pub glwe_dimension: u32,
-    pub polynomial_size: u32,
-}
-#[allow(clippy::unnecessary_operation, clippy::identity_op)]
-const _: () = {
-    ["Size of CudaPackedGlweCiphertextListFFI"]
-        [::std::mem::size_of::<CudaPackedGlweCiphertextListFFI>() - 32usize];
-    ["Alignment of CudaPackedGlweCiphertextListFFI"]
-        [::std::mem::align_of::<CudaPackedGlweCiphertextListFFI>() - 8usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::ptr"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, ptr) - 0usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::storage_log_modulus"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, storage_log_modulus) - 8usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::lwe_per_glwe"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, lwe_per_glwe) - 12usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::total_lwe_bodies_count"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, total_lwe_bodies_count) - 16usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::glwe_dimension"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, glwe_dimension) - 20usize];
-    ["Offset of field: CudaPackedGlweCiphertextListFFI::polynomial_size"]
-        [::std::mem::offset_of!(CudaPackedGlweCiphertextListFFI, polynomial_size) - 24usize];
-};
 unsafe extern "C" {
     pub fn scratch_cuda_integer_compress_radix_ciphertext_64(
         streams: CudaStreamsFFI,

backends/tfhe-hpu-backend/Cargo.toml

@@ -33,7 +33,7 @@ serde = { version = "1", features = ["derive"] }
 toml = { version = "0.8", features = [] }
 paste = "1.0.15"
 thiserror = "1.0.61"
-bytemuck = "1.16.0"
+bytemuck = { workspace = true }
 anyhow = "1.0.82"
 lazy_static = "1.4.0"
 rand = "0.8.5"
@@ -61,8 +61,12 @@ bitvec = { version = "1.0", optional = true }
 serde_json = { version = "1.0", optional = true }
 
 # Dependencies used for v80 pdi handling
-bincode ={ version = "1.3", optional = true}
-serde_derive ={ version = "1.0", optional = true}
+bincode = { version = "1.3", optional = true }
+serde_derive = { version = "1.0", optional = true }
+
+# Fix build for toolchain < 1.85, newer versions are edition 2024 which requires rust 1.85
+ignore = { version = "<0.4.24" }
+globset = { version = "<0.4.17" }
 
 # Binary for manual debugging
 # Enable to access Hpu register and drive some custom sequence by hand

backends/tfhe-hpu-backend/config_store/sim/hpu_config.toml

@@ -26,7 +26,7 @@
   lut_pc = {Hbm={pc=34}}
 
   fw_size= 16777215 # i.e. 16 MiB
-  fw_pc = {Ddr= {offset= 0x3900_0000}} # NB: Allocation must take place in the Discret DDR
+  fw_pc = {Ddr= {offset= 0x3900_0000}} # NB: Allocation must take place in the Discrete DDR
 
   bsk_pc = [
     {Hbm={pc=8}},

backends/tfhe-hpu-backend/config_store/v80/hpu_config.toml

@@ -31,7 +31,7 @@
   lut_pc = {Hbm={pc=34}}
 
   fw_size= 16777216 # i.e. 16 MiB
-  fw_pc = {Ddr= {offset= 0x3900_0000}} # NB: Allocation must take place in the Discret DDR
+  fw_pc = {Ddr= {offset= 0x3900_0000}} # NB: Allocation must take place in the Discrete DDR
 
   bsk_pc = [
     {Hbm={pc=8}},

backends/tfhe-hpu-backend/src/asm/dop/field.rs

@@ -81,7 +81,7 @@ pub enum ImmId {
 }
 
 impl ImmId {
-    /// Create new immediat template
+    /// Create new immediate template
     pub fn new_var(tid: u8, bid: u8) -> Self {
         Self::Var { tid, bid }
     }
@@ -132,7 +132,7 @@ pub struct PeArithInsn {
 }
 
 /// PeaMsg instructions
-/// Arithmetic operation that use one destination register, one source register and an immediat
+/// Arithmetic operation that use one destination register, one source register and an immediate
 /// value
 #[derive(Clone, Copy, Debug, PartialEq, Eq, serde::Serialize, serde::Deserialize)]
 pub struct PeArithMsgInsn {

backends/tfhe-hpu-backend/src/asm/dop/fmt.rs

@@ -69,7 +69,7 @@ impl From<&PeArithHex> for PeArithInsn {
 }
 
 /// PeaMsg instructions
-/// Arithmetic operation that use one destination register, one source register and an immediat
+/// Arithmetic operation that use one destination register, one source register and an immediate
 /// value
 #[bitfield(u32)]
 pub struct PeArithMsgHex {

backends/tfhe-hpu-backend/src/asm/iop/arg.rs

@@ -4,7 +4,7 @@
 
 use super::*;
 use field::{
-    FwMode, IOpHeader, IOpcode, ImmBundle, Immediat, Operand, OperandBlock, OperandBundle,
+    FwMode, IOpHeader, IOpcode, ImmBundle, Immediate, Operand, OperandBlock, OperandBundle,
 };
 use lazy_static::lazy_static;
 
@@ -393,7 +393,7 @@ impl std::str::FromStr for ImmBundle {
                         .map_err(|err| ParsingError::InvalidArg(err.to_string()))?
                 };
 
-                Ok(Immediat::from_cst(imm))
+                Ok(Immediate::from_cst(imm))
             })
             .collect::<Result<Vec<_>, ParsingError>>()?;
 
@@ -475,7 +475,7 @@ impl std::fmt::Display for IOp {
         // Source operands list
         write!(f, " <{}>", self.src)?;
 
-        // Immediat operands list [Optional]
+        // Immediate operands list [Optional]
         if self.header.has_imm {
             write!(f, " <{}>", self.imm)?;
         }

backends/tfhe-hpu-backend/src/asm/iop/field.rs

@@ -75,7 +75,7 @@ impl Operand {
     }
 }
 
-/// Create a dedicated type for a collection of Immediat
+/// Create a dedicated type for a collection of Immediate
 /// This is to enable trait implementation on it (c.f arg)
 #[derive(Debug, Clone)]
 pub struct OperandBundle(Vec<Operand>);
@@ -142,32 +142,32 @@ impl OperandBundle {
 
 // Immediate operands
 // ------------------------------------------------------------------------------------------------
-/// Immediat Size
-/// => Number of valid digit in following immediat
+/// Immediate Size
+/// => Number of valid digit in following immediate
 /// To obtain the number of valid bits, user should multiply by the msg_width
 #[derive(Debug, Clone, Copy, PartialEq)]
 pub struct ImmBlock(pub u16);
 
-/// Immediat header
+/// Immediate header
 /// Use to implement top-level parser manually
 #[derive(Debug, Clone, Copy, PartialEq)]
-pub struct ImmediatHeader {
+pub struct ImmediateHeader {
     pub(super) lsb_msg: u16,
     pub(super) block: ImmBlock,
     pub(super) is_last: bool,
     pub(super) kind: OperandKind,
 }
 
-/// Full Immediat representation (i.e. header + data)
+/// Full Immediate representation (i.e. header + data)
 #[derive(Debug, Clone, PartialEq)]
-pub struct Immediat {
+pub struct Immediate {
     pub(super) kind: OperandKind,
     pub(super) is_last: bool,
     pub(super) block: ImmBlock,
     pub(super) msg: Vec<u16>,
 }
 
-impl Immediat {
+impl Immediate {
     /// Access imm msg for template patching
     /// Extract the correct block (i.e. MSG_WIDTH chunk)
     pub fn msg_block(&self, bid: u8) -> u16 {
@@ -217,7 +217,7 @@ impl Immediat {
     }
 }
 
-impl Immediat {
+impl Immediate {
     #[tracing::instrument(level = "trace", ret)]
     pub fn from_words(stream: &[IOpWordRepr]) -> Result<(Self, usize), HexParsingError> {
         // Keep track of the current peak index
@@ -226,7 +226,7 @@ impl Immediat {
         // 1. Parse header
         let header = if let Some(header_word) = stream.get(peak_words) {
             peak_words += 1;
-            ImmediatHeader::from(&fmt::ImmediatHeaderHex::from_bits(*header_word))
+            ImmediateHeader::from(&fmt::ImmediateHeaderHex::from_bits(*header_word))
         } else {
             return Err(HexParsingError::EmptyStream);
         };
@@ -276,13 +276,13 @@ impl Immediat {
 
     pub fn to_words(&self) -> Vec<IOpWordRepr> {
         let mut words = Vec::new();
-        let header = ImmediatHeader {
+        let header = ImmediateHeader {
             lsb_msg: *self.msg.first().unwrap_or(&0),
             block: self.block,
             is_last: self.is_last,
             kind: self.kind,
         };
-        words.push(fmt::ImmediatHeaderHex::from(&header).into_bits());
+        words.push(fmt::ImmediateHeaderHex::from(&header).into_bits());
 
         if self.msg.len() > 1 {
             for imm in self.msg[1..]
@@ -302,10 +302,10 @@ impl Immediat {
     }
 }
 
-/// Create a dedicated type for a collection of Immediat
+/// Create a dedicated type for a collection of Immediate
 /// This is to enable trait implementation on it (c.f arg)
 #[derive(Debug, Clone)]
-pub struct ImmBundle(Vec<Immediat>);
+pub struct ImmBundle(Vec<Immediate>);
 
 impl ImmBundle {
     #[tracing::instrument(level = "trace", ret)]
@@ -315,7 +315,7 @@ impl ImmBundle {
 
         let mut imm_list = Vec::new();
         loop {
-            let (imm, peaked) = Immediat::from_words(&stream[peak_words..])?;
+            let (imm, peaked) = Immediate::from_words(&stream[peak_words..])?;
             peak_words += peaked;
 
             let is_last = imm.is_last;
@@ -335,9 +335,9 @@ impl ImmBundle {
     }
 }
 
-impl From<Vec<Immediat>> for ImmBundle {
+impl From<Vec<Immediate>> for ImmBundle {
     #[tracing::instrument(level = "trace", ret)]
-    fn from(inner: Vec<Immediat>) -> Self {
+    fn from(inner: Vec<Immediate>) -> Self {
         let mut inner = inner;
         // Enforce correct is_last handling
         inner.iter_mut().for_each(|op| op.is_last = false);
@@ -349,7 +349,7 @@ impl From<Vec<Immediat>> for ImmBundle {
 }
 
 impl std::ops::Deref for ImmBundle {
-    type Target = Vec<Immediat>;
+    type Target = Vec<Immediate>;
 
     fn deref(&self) -> &Self::Target {
         &self.0
@@ -392,7 +392,7 @@ use std::collections::VecDeque;
 /// Implement construction
 /// Used to construct IOp from Backend HpuVar
 impl IOp {
-    pub fn new(opcode: IOpcode, dst: Vec<Operand>, src: Vec<Operand>, imm: Vec<Immediat>) -> Self {
+    pub fn new(opcode: IOpcode, dst: Vec<Operand>, src: Vec<Operand>, imm: Vec<Immediate>) -> Self {
         let dst_align = dst.iter().map(|x| x.block).max().unwrap();
         let src_align = src.iter().map(|x| x.block).max().unwrap();
         let has_imm = !imm.is_empty();
@@ -504,7 +504,7 @@ impl IOp {
             src
         };
 
-        // 4. Parse Immediat [Optional]
+        // 4. Parse Immediate [Optional]
         let (imm, peaked) = if header.has_imm {
             ImmBundle::from_words(&stream.as_slices().0[peak_words..])?
         } else {
@@ -533,7 +533,7 @@ impl IOp {
         words.extend(self.dst.to_words());
         // 3. Sources
         words.extend(self.src.to_words());
-        // 4. Immediat
+        // 4. Immediate
         words.extend(self.imm.to_words());
         words
     }

backends/tfhe-hpu-backend/src/asm/iop/fmt.rs

@@ -60,7 +60,7 @@ impl From<&Operand> for OperandHex {
 }
 
 #[bitfield(u32)]
-pub struct ImmediatHeaderHex {
+pub struct ImmediateHeaderHex {
     #[bits(16)]
     lsb_msg: u16,
     #[bits(12)]
@@ -73,8 +73,8 @@ pub struct ImmediatHeaderHex {
     kind: u8,
 }
 
-impl From<&ImmediatHeaderHex> for field::ImmediatHeader {
-    fn from(value: &ImmediatHeaderHex) -> Self {
+impl From<&ImmediateHeaderHex> for field::ImmediateHeader {
+    fn from(value: &ImmediateHeaderHex) -> Self {
         let kind = if value.kind() == OperandKind::Src as u8 {
             OperandKind::Src
         } else if value.kind() == OperandKind::Dst as u8 {
@@ -94,8 +94,8 @@ impl From<&ImmediatHeaderHex> for field::ImmediatHeader {
     }
 }
 
-impl From<&field::ImmediatHeader> for ImmediatHeaderHex {
-    fn from(value: &field::ImmediatHeader) -> Self {
+impl From<&field::ImmediateHeader> for ImmediateHeaderHex {
+    fn from(value: &field::ImmediateHeader) -> Self {
         Self::new()
             .with_lsb_msg(value.lsb_msg)
             .with_block(value.block.0)

backends/tfhe-hpu-backend/src/asm/iop/mod.rs

@@ -2,7 +2,7 @@
 //! IOp definition
 
 mod field;
-pub use field::{HexParsingError, IOp, IOpcode, Immediat, Operand, OperandKind};
+pub use field::{HexParsingError, IOp, IOpcode, Immediate, Operand, OperandKind};
 mod fmt;
 pub use fmt::{IOpRepr, IOpWordRepr};
 mod iop_macro;

backends/tfhe-hpu-backend/src/asm/mod.rs

@@ -127,11 +127,7 @@ where
                     Ok(op) => asm_ops.push(AsmOp::Stmt(op)),
                     Err(err) => {
                         tracing::warn!("ReadAsm failed @{file}:{}", line + 1);
-                        anyhow::bail!(
-                            "ReadAsm failed @{file}:{} with {}",
-                            line + 1,
-                            err.to_string()
-                        );
+                        anyhow::bail!("ReadAsm failed @{file}:{} with {}", line + 1, err);
                     }
                 }
             }

backends/tfhe-hpu-backend/src/asm/tests/dop.asm

@@ -20,7 +20,7 @@ SUB R2 R1 R3
 MUL R2 R1 R3
 MAC R2 R1 R3 4
 
-; Test ArithMsg operation with various immediat template format
+; Test ArithMsg operation with various immediate template format
 ADDS R2 R1 10
 SUBS R2 R1 TI[4].0
 SSUB R2 R1 TI[2].4

backends/tfhe-hpu-backend/src/asm/tests/iop.asm

@@ -2,7 +2,7 @@
 ; with the != available arguments modes
 
 ; Simple Mode:
-; 1 destination, 2 sources, no immediat
+; 1 destination, 2 sources, no immediate
 ; With raw opcode -> 0x35
 IOP[0x35] <I8 I8> <I8@0x08> <I8@0x0 I8@0x4>
 ; With raw opcode -> 40 and dynamic Fw generation
@@ -11,7 +11,7 @@ IOP[0x35] <dyn I8 I8> <I8@0x08> <I8@0x0 I8@0x4>
 ; With opcode alias -> MUL
 MUL <I8 I8> <I8@0x08> <I8@0x0 I4@0x4>
 
-; Simple Mode with immediat
+; Simple Mode with immediate
 ; Source operands are defined through vector mode
 MULS <I8 I8> <I8@0x8> <I8[2]@0x0> <0xaf>
 
@@ -26,6 +26,6 @@ IOP[0x60] <dyn I8 I8> <I8@0x10 I8@0x14> <I8@0x0 I8@0x4>
 ; Previous operation could be defined with vector format.
 IOP[0x40] <dyn I8 I8> <I8[2]@0x10> <I8[2]@0x0>
 
-; With multiple immediat
+; With multiple immediate
 ; Example this operation could compute D <- A*4 + B*8  
 IOP[0x0] <I16 I16> <I16@16> <I16@0x0 I8@0x8> <0xdeadc0de>

backends/tfhe-hpu-backend/src/entities/parameters.rs

@@ -148,10 +148,10 @@ pub struct HpuPcParameters {
     pub ksk_pc: usize,
     pub ksk_bytes_w: usize,
     pub bsk_pc: usize,
+    pub glwe_pc: usize, // Currently hardcoded to 1
     pub bsk_bytes_w: usize,
     pub pem_pc: usize,
     pub pem_bytes_w: usize,
-    // pub glwe_pc: usize, // Currently hardcoded to 1
     pub glwe_bytes_w: usize,
 }
 

backends/tfhe-hpu-backend/src/ffi/v80/pdi/mod.rs

@@ -84,7 +84,7 @@ impl HpuV80Pdi {
     }
 
     #[allow(unused)]
-    /// Deconstruct HpuV80Pdi into a folder with discret files
+    /// Deconstruct HpuV80Pdi into a folder with discrete files
     pub fn to_folder(&self, folder_path: &str) -> Result<(), Box<dyn Error>> {
         let metadata_path = Path::new(folder_path).join("metadata.toml");
         self.metadata.to_toml(

backends/tfhe-hpu-backend/src/fw/fw_impl/ilp.rs

@@ -100,11 +100,11 @@ pub fn iop_adds(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic addx function
     iop_addx(prog, &mut dst, None, &src_a, &src_b);
 }
@@ -134,18 +134,18 @@ pub fn iop_overflow_adds(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic addx function
     iop_addx(prog, &mut dst, Some(&mut flag[0]), &src_a, &src_b);
 }
 
 /// Generic Add operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_addx(
     prog: &mut Program,
@@ -191,7 +191,7 @@ pub fn iop_sub(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -207,11 +207,11 @@ pub fn iop_subs(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic subx function
     iop_subx(prog, &mut dst, None, &src_a, &src_b);
 }
@@ -224,7 +224,7 @@ pub fn iop_overflow_sub(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -241,18 +241,18 @@ pub fn iop_overflow_subs(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic subx function
     iop_subx(prog, &mut dst, Some(&mut flag[0]), &src_a, &src_b);
 }
 
 /// Generic sub operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_subx(
     prog: &mut Program,
@@ -324,18 +324,18 @@ pub fn iop_ssub(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SSUB Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SSUB Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic subx function
     iop_ssubx(prog, &mut dst, None, &src_a, &src_b);
 }
 
 /// Generic SSUB operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_ssubx(
     prog: &mut Program,
@@ -406,11 +406,11 @@ pub fn iop_overflow_ssub(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic ssubx function
     iop_ssubx(prog, &mut dst, Some(&mut flag[0]), &src_a, &src_b);
 }
@@ -422,7 +422,7 @@ pub fn iop_mul(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -437,11 +437,11 @@ pub fn iop_muls(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic mulx function
     iop_mulx(prog, &mut dst, None, &src_a, &src_b);
 }
@@ -454,7 +454,7 @@ pub fn iop_overflow_mul(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -471,18 +471,18 @@ pub fn iop_overflow_muls(prog: &mut Program) {
     let mut flag = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic mulx function
     iop_mulx(prog, &mut dst, Some(&mut flag[0]), &src_a, &src_b);
 }
 
 /// Generic mul operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_mulx(
     prog: &mut Program,
@@ -994,11 +994,11 @@ pub fn iop_shift_scalar_right(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // Src -> Operand
     let src = prog.iop_template_var(OperandKind::Src, 0);
-    // Amount-> Immediat
+    // Amount-> Immediate
     let amount = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SHIFTS_R Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SHIFTS_R Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic rotx function
     iop_scalar_shiftrotx(prog, ShiftKind::ShiftRight, &mut dst, &src, &amount);
 }
@@ -1010,11 +1010,11 @@ pub fn iop_shift_scalar_left(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // Src -> Operand
     let src = prog.iop_template_var(OperandKind::Src, 0);
-    // Amount-> Immediat
+    // Amount-> Immediate
     let amount = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SHIFTS_L Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SHIFTS_L Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic rotx function
     iop_scalar_shiftrotx(prog, ShiftKind::ShiftLeft, &mut dst, &src, &amount);
 }
@@ -1026,11 +1026,11 @@ pub fn iop_rotate_scalar_right(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // Src -> Operand
     let src = prog.iop_template_var(OperandKind::Src, 0);
-    // Amount-> Immediat
+    // Amount-> Immediate
     let amount = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("ROTS_R Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("ROTS_R Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic rotx function
     iop_scalar_shiftrotx(prog, ShiftKind::RotRight, &mut dst, &src, &amount);
 }
@@ -1041,11 +1041,11 @@ pub fn iop_rotate_scalar_left(prog: &mut Program) {
     let mut dst = prog.iop_template_var(OperandKind::Dst, 0);
     // Src -> Operand
     let src = prog.iop_template_var(OperandKind::Src, 0);
-    // Amount-> Immediat
+    // Amount-> Immediate
     let amount = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("ROTS_L Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("ROTS_L Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic rotx function
     iop_scalar_shiftrotx(prog, ShiftKind::RotRight, &mut dst, &src, &amount);
 }
@@ -1118,7 +1118,7 @@ pub fn iop_cmp(prog: &mut Program, cmp_op: Pbs) {
 
 /// Generic Cmp operation
 /// One destination block and two sources operands
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_cmpx(
     prog: &mut Program,

backends/tfhe-hpu-backend/src/fw/fw_impl/ilp_div.rs

@@ -21,7 +21,7 @@ pub fn iop_div(prog: &mut Program) {
     let mut dst_remain = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -37,11 +37,11 @@ pub fn iop_divs(prog: &mut Program) {
     let mut dst_remain = prog.iop_template_var(OperandKind::Dst, 1);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("DIVS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("DIVS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic divx function
     // TODO: do computation on immediate directly for more efficiency.
     // Workaround: transform immediate into ct.
@@ -52,7 +52,7 @@ pub fn iop_divs(prog: &mut Program) {
 
 /// Generic div operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 pub fn iop_divx(
     prog: &mut Program,
     dst_quotient: &mut [metavar::MetaVarCell],
@@ -80,7 +80,7 @@ pub fn iop_mod(prog: &mut Program) {
     let mut dst_remain = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -95,11 +95,11 @@ pub fn iop_mods(prog: &mut Program) {
     let mut dst_remain = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("MODS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("MODS Operand::Dst Operand::Src Operand::Immediate".to_string());
     // Deferred implementation to generic modx function
     // TODO: do computation on immediate directly for more efficiency.
     // Workaround: transform immediate into ct.
@@ -110,7 +110,7 @@ pub fn iop_mods(prog: &mut Program) {
 
 /// Generic mod operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 pub fn iop_modx(
     prog: &mut Program,
     dst_remain: &mut [metavar::MetaVarCell],

backends/tfhe-hpu-backend/src/fw/fw_impl/llt/mod.rs

@@ -98,7 +98,7 @@ pub fn iop_add(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -123,11 +123,11 @@ pub fn iop_adds(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("ADDS Operand::Dst Operand::Src Operand::Immediate".to_string());
     iop_addx(prog, dst, src_a, src_b);
 }
 
@@ -138,7 +138,7 @@ pub fn iop_sub(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -152,11 +152,11 @@ pub fn iop_subs(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SUBS Operand::Dst Operand::Src Operand::Immediate".to_string());
     iop_subx(prog, dst, src_a, src_b);
 }
 
@@ -166,11 +166,11 @@ pub fn iop_ssub(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Imm, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 0);
 
     // Add Comment header
-    prog.push_comment("SSUB Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("SSUB Operand::Dst Operand::Src Operand::Immediate".to_string());
     iop_subx(prog, dst, src_a, src_b);
 }
 
@@ -199,7 +199,7 @@ pub fn iop_mul(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Src, 1);
 
     // Add Comment header
@@ -215,11 +215,11 @@ pub fn iop_muls(prog: &mut Program) {
     let dst = prog.iop_template_var(OperandKind::Dst, 0);
     // SrcA -> Operand
     let src_a = prog.iop_template_var(OperandKind::Src, 0);
-    // SrcB -> Immediat
+    // SrcB -> Immediate
     let src_b = prog.iop_template_var(OperandKind::Imm, 0);
 
     // Add Comment header
-    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediat".to_string());
+    prog.push_comment("MULS Operand::Dst Operand::Src Operand::Immediate".to_string());
 
     iop_mulx(prog, dst, src_a, src_b).add_to_prog(prog);
 }
@@ -534,7 +534,7 @@ pub fn iop_mulx(
 
 /// Generic mul operation
 /// One destination and two sources operation
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_mulx_ser(
     prog: &mut Program,
@@ -654,7 +654,7 @@ pub fn iop_mulx_ser(
 
 /// Generic Cmp operation
 /// One destination block and two sources operands
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_cmpx(
     prog: &mut Program,
@@ -673,7 +673,7 @@ pub fn iop_cmpx(
 
 /// Generic Cmp operation
 /// One destination block and two sources operands
-/// Source could be Operand or Immediat
+/// Source could be Operand or Immediate
 #[instrument(level = "trace", skip(prog))]
 pub fn iop_cmpx_rtl(
     prog: &mut Program,

backends/tfhe-hpu-backend/src/fw/mod.rs

@@ -3,7 +3,7 @@
 //!
 //! Provide two concrete implementation of those traits
 //! * DigitOperations (DOp)
-//! * IntegerOperarions (IOp)
+//! * IntegerOperations (IOp)
 
 pub mod fw_impl;
 pub mod isc_sim;
@@ -72,7 +72,7 @@ impl From<FwParameters> for asm::DigitParameters {
 }
 
 /// Fw trait abstraction
-/// Use to handle Fw implemantion in an abstract way
+/// Use to handle Fw implementation in an abstract way
 #[enum_dispatch]
 pub trait Fw {
     /// Expand a program of IOp into a program of DOp

backends/tfhe-hpu-backend/src/fw/program.rs

@@ -382,9 +382,9 @@ impl Program {
     }
 
     /// Create templated arguments
-    /// kind is used to specify if it's bind to src/dst or immediat template
+    /// kind is used to specify if it's bind to src/dst or immediate template
     /// pos_id is used to bind the template to an IOp operand position
-    // TODO pass the associated operand or immediat to obtain the inner blk properties instead of
+    // TODO pass the associated operand or immediate to obtain the inner blk properties instead of
     // using the global one
     pub fn iop_template_var(&mut self, kind: asm::OperandKind, pos_id: u8) -> Vec<MetaVarCell> {
         let nb_blk = self.params().blk_w() as u8;

backends/tfhe-hpu-backend/src/interface/cmd.rs

@@ -2,11 +2,11 @@
 /// Help with IOp management over HPU
 /// Track IOp status and handle backward update of associated HpuVariable
 use super::*;
-use crate::asm::iop::{Immediat, Operand, OperandKind};
+use crate::asm::iop::{Immediate, Operand, OperandKind};
 use crate::asm::{IOp, IOpcode};
 use variable::HpuVarWrapped;
 
-/// Underlying type used for Immediat value;
+/// Underlying type used for Immediate value;
 pub type HpuImm = u128;
 
 /// Structure that hold an IOp with there associated operands
@@ -15,7 +15,7 @@ pub struct HpuCmd {
     pub(crate) op: IOp,
     pub(crate) dst: Vec<HpuVarWrapped>,
     pub(crate) src: Vec<HpuVarWrapped>,
-    // NB: No need to track Immediat lifetime. It's simply constant completely held by the IOp
+    // NB: No need to track Immediate lifetime. It's simply constant completely held by the IOp
     // definition
 }
 
@@ -75,7 +75,7 @@ impl HpuCmd {
             .collect::<Vec<_>>();
         let imm_op = imm
             .iter()
-            .map(|var| Immediat::from_cst(*var))
+            .map(|var| Immediate::from_cst(*var))
             .collect::<Vec<_>>();
 
         let op = IOp::new(opcode, dst_op, src_op, imm_op);

backends/tfhe-hpu-backend/src/interface/rtl/params.rs

@@ -186,6 +186,7 @@ impl FromRtl for HpuPcParameters {
         let ksk_pc = *hbm_pc_fields.get("ksk_pc").expect("Unknown field") as usize;
         let bsk_pc = *hbm_pc_fields.get("bsk_pc").expect("Unknown field") as usize;
         let pem_pc = *hbm_pc_fields.get("pem_pc").expect("Unknown field") as usize;
+        let glwe_pc = *hbm_pc_fields.get("glwe_pc").expect("Unknown field") as usize;
 
         // Extract bus width for each channel
         let ksk_bytes_w = {
@@ -229,6 +230,7 @@ impl FromRtl for HpuPcParameters {
             ksk_pc,
             bsk_pc,
             pem_pc,
+            glwe_pc,
             ksk_bytes_w,
             bsk_bytes_w,
             pem_bytes_w,

backends/tfhe-hpu-backend/src/utils/fw.rs

@@ -84,6 +84,11 @@ pub struct Args {
     #[arg(long, default_value = "Patient")]
     flush_behaviour: FlushBehaviour,
 
+    /// Force parallel IOp implementations or not. By default this is derived
+    /// from the batch size.
+    #[arg(long, default_value = None)]
+    parallel: Option<bool>,
+
     /// Integer bit width
     #[arg(long, default_value_t = 8)]
     integer_w: usize,
@@ -168,6 +173,7 @@ fn main() -> Result<(), anyhow::Error> {
             use_tiers: args.use_tiers,
             flush: args.flush,
             flush_behaviour: args.flush_behaviour,
+            parallel: None,
         }),
         cur_op_cfg: OpCfg::default(),
         pe_cfg,

ci/slab.toml

@@ -83,18 +83,6 @@ image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
 flavor_name = "n3-A100x8-NVLink"
 user = "ubuntu"
 
-[backend.hyperstack.multi-gpu-test]
-environment_name = "canada"
-image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
-flavor_name = "n3-L40x4"
-user = "ubuntu"
-
-[backend.hyperstack.multi-gpu-test_fallback]
-environment_name = "canada"
-image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
-flavor_name = "n3-RTX-A6000x2"
-user = "ubuntu"
-
 [backend.hyperstack.l40]
 environment_name = "canada"
 image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
@@ -106,3 +94,9 @@ environment_name = "canada"
 image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
 flavor_name = "n3-RTX-A6000x1"
 user = "ubuntu"
+
+[backend.hyperstack.4-l40]
+environment_name = "canada"
+image_name = "Ubuntu Server 22.04 LTS R570 CUDA 12.8"
+flavor_name = "n3-L40x4"
+user = "ubuntu"

mockups/tfhe-hpu-mockup/Cargo.toml

@@ -18,7 +18,7 @@ ipc-channel = "0.18.3"
 
 strum = { version = "0.26.2", features = ["derive"] }
 strum_macros = "0.26.2"
-bytemuck = "1.16.0"
+bytemuck = { workspace = true }
 
 clap = { version = "4.4.4", features = ["derive"] }
 clap-num = "*"

mockups/tfhe-hpu-mockup/params/gaussian_44b.toml

@@ -37,6 +37,7 @@
     bsk_bytes_w= 64
     pem_pc= 2
     pem_bytes_w= 64
+    glwe_pc= 1
     glwe_bytes_w= 64
 [regf_params]
     reg_nb= 64

mockups/tfhe-hpu-mockup/params/gaussian_44b_fast.toml

@@ -37,6 +37,7 @@
     bsk_bytes_w= 64
     pem_pc= 2
     pem_bytes_w= 64
+    glwe_pc= 1
     glwe_bytes_w= 64
 [regf_params]
     reg_nb= 64

mockups/tfhe-hpu-mockup/params/gaussian_64b.toml

@@ -39,6 +39,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 
 [regf_params]

mockups/tfhe-hpu-mockup/params/gaussian_64b_fast.toml

@@ -37,6 +37,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 [regf_params]
     reg_nb= 64

mockups/tfhe-hpu-mockup/params/gaussian_64b_pfail64.toml

@@ -39,6 +39,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 
 [regf_params]

mockups/tfhe-hpu-mockup/params/gaussian_64b_pfail64_psi64.toml

@@ -40,6 +40,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 
 [regf_params]

mockups/tfhe-hpu-mockup/params/tuniform_64b_fast.toml

@@ -37,6 +37,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 [regf_params]
     reg_nb= 64

mockups/tfhe-hpu-mockup/params/tuniform_64b_pfail128_psi64.toml

@@ -40,6 +40,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 
 [regf_params]

mockups/tfhe-hpu-mockup/params/tuniform_64b_pfail64_psi64.toml

@@ -40,6 +40,7 @@
     bsk_bytes_w= 32
     pem_pc= 2
     pem_bytes_w= 32
+    glwe_pc= 1
     glwe_bytes_w= 32
 
 [regf_params]

mockups/tfhe-hpu-mockup/src/modules/regmap.rs

@@ -1,4 +1,4 @@
-use std::collections::VecDeque;
+use std::collections::{HashMap, VecDeque};
 use std::sync::atomic::{AtomicBool, Ordering};
 
 use super::*;
@@ -101,10 +101,21 @@ impl RegisterMap {
     /// Return register value from parameter value
     pub fn read_reg(&mut self, addr: u64) -> u32 {
         let register_name = self.get_register_name(addr);
+        let cur_reg = self
+            .regmap
+            .register()
+            .get(register_name)
+            .expect("Unknown register, check regmap definition");
+
         match register_name {
             "info::ntt_structure" => {
                 let ntt_p = &self.rtl_params.ntt_params;
-                (ntt_p.radix + (ntt_p.psi << 8) /*+(ntt_p.div << 16)*/ + (ntt_p.delta << 24)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("radix", ntt_p.radix as u32),
+                    ("psi", ntt_p.psi as u32),
+                    ("delta", ntt_p.delta as u32),
+                    //("div", ntt_p.div as u32),
+                ]))
             }
             "info::ntt_rdx_cut" => {
                 let ntt_p = &self.rtl_params.ntt_params;
@@ -112,10 +123,11 @@ impl RegisterMap {
                     HpuNttCoreArch::GF64(cut_w) => cut_w,
                     _ => &vec![ntt_p.delta as u8],
                 };
-                cut_w
-                    .iter()
-                    .enumerate()
-                    .fold(0, |acc, (id, val)| acc + ((*val as u32) << (id * 4)))
+                let mut hash_cut: HashMap<&str, u32> = HashMap::new();
+                for (f, val) in cur_reg.field().iter().zip(cut_w) {
+                    hash_cut.insert(f.name(), *val as u32);
+                }
+                cur_reg.from_field(hash_cut)
             }
             "info::ntt_architecture" => match self.rtl_params.ntt_params.core_arch {
                 HpuNttCoreArch::WmmCompactPcg => NTT_CORE_ARCH_OFS + 4,
@@ -124,7 +136,10 @@ impl RegisterMap {
             },
             "info::ntt_pbs" => {
                 let ntt_p = &self.rtl_params.ntt_params;
-                (ntt_p.batch_pbs_nb + (ntt_p.total_pbs_nb << 8)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("batch_pbs_nb", ntt_p.batch_pbs_nb as u32),
+                    ("total_pbs_nb", ntt_p.total_pbs_nb as u32),
+                ]))
             }
             "info::ntt_modulo" => {
                 MOD_NTT_NAME_OFS + (self.rtl_params.ntt_params.prime_modulus as u8) as u32
@@ -159,17 +174,29 @@ impl RegisterMap {
             }
             "info::ks_structure" => {
                 let ks_p = &self.rtl_params.ks_params;
-                (ks_p.lbx + (ks_p.lby << 8) + (ks_p.lbz << 16)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("x", ks_p.lbx as u32),
+                    ("y", ks_p.lby as u32),
+                    ("z", ks_p.lbz as u32),
+                ]))
             }
             "info::ks_crypto_param" => {
                 let ks_p = &self.rtl_params.ks_params;
                 let pbs_p = &self.rtl_params.pbs_params;
-                (ks_p.width + (pbs_p.ks_level << 8) + (pbs_p.ks_base_log << 16)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("mod_ksk_w", ks_p.width as u32),
+                    ("ks_l", pbs_p.ks_level as u32),
+                    ("ks_b", pbs_p.ks_base_log as u32),
+                ]))
             }
             "info::hbm_axi4_nb" => {
                 let pc_p = &self.rtl_params.pc_params;
-                // TODO: Cut number currently not reverted
-                (pc_p.bsk_pc + (pc_p.ksk_pc << 8) + (pc_p.pem_pc << 16)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("ksk_pc", pc_p.ksk_pc as u32),
+                    ("bsk_pc", pc_p.bsk_pc as u32),
+                    ("pem_pc", pc_p.pem_pc as u32),
+                    ("glwe_pc", pc_p.glwe_pc as u32),
+                ]))
             }
             "info::hbm_axi4_dataw_ksk" => {
                 let bytes_w = &self.rtl_params.pc_params.ksk_bytes_w;
@@ -190,11 +217,17 @@ impl RegisterMap {
 
             "info::regf_structure" => {
                 let regf_p = &self.rtl_params.regf_params;
-                (regf_p.reg_nb + (regf_p.coef_nb << 8)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("reg_nb", regf_p.reg_nb as u32),
+                    ("coef_nb", regf_p.coef_nb as u32),
+                ]))
             }
             "info::isc_structure" => {
                 let isc_p = &self.rtl_params.isc_params;
-                (isc_p.depth + (isc_p.min_iop_size << 8)) as u32
+                cur_reg.from_field(HashMap::from([
+                    ("depth", isc_p.depth as u32),
+                    ("min_iop_size", isc_p.min_iop_size as u32),
+                ]))
             }
 
             "bsk_avail::avail" => self.bsk.avail.load(Ordering::SeqCst) as u32,
@@ -217,9 +250,10 @@ impl RegisterMap {
             }
 
             // Bpip configuration registers
-            "bpip::use" => {
-                ((self.bpip.used as u8) + ((self.bpip.use_opportunism as u8) << 1)) as u32
-            }
+            "bpip::use" => cur_reg.from_field(HashMap::from([
+                ("use_bpip", self.bpip.used as u32),
+                ("use_opportunism", self.bpip.use_opportunism as u32),
+            ])),
             "bpip::timeout" => self.bpip.timeout,
 
             // Add offset configuration registers
@@ -318,13 +352,23 @@ impl RegisterMap {
 
     pub fn write_reg(&mut self, addr: u64, value: u32) -> RegisterEvent {
         let register_name = self.get_register_name(addr);
+        let cur_reg = self
+            .regmap
+            .register()
+            .get(register_name)
+            .expect("Unknown register, check regmap definition");
+        let hash_val = cur_reg.as_field(value);
+
         match register_name {
             "bsk_avail::avail" => {
-                self.bsk.avail.store((value & 0x1) == 0x1, Ordering::SeqCst);
+                let avail = hash_val.get("avail");
+
+                self.bsk.avail.store(avail == Some(&0x1), Ordering::SeqCst);
                 RegisterEvent::None
             }
             "bsk_avail::reset" => {
-                if (value & 0x1) == 0x1 {
+                let req = hash_val.get("request");
+                if req == Some(&0x1) {
                     self.bsk.rst_pdg.store(true, Ordering::SeqCst);
                     self.bsk.avail.store(false, Ordering::SeqCst);
                     RegisterEvent::KeyReset
@@ -333,11 +377,14 @@ impl RegisterMap {
                 }
             }
             "ksk_avail::avail" => {
-                self.ksk.avail.store((value & 0x1) == 0x1, Ordering::SeqCst);
+                let avail = hash_val.get("avail");
+
+                self.ksk.avail.store(avail == Some(&0x1), Ordering::SeqCst);
                 RegisterEvent::None
             }
             "ksk_avail::reset" => {
-                if (value & 0x1) == 0x1 {
+                let req = hash_val.get("request");
+                if req == Some(&0x1) {
                     self.ksk.rst_pdg.store(true, Ordering::SeqCst);
                     self.ksk.avail.store(false, Ordering::SeqCst);
                     RegisterEvent::KeyReset
@@ -348,8 +395,10 @@ impl RegisterMap {
 
             // Bpip configuration registers
             "bpip::use" => {
-                self.bpip.used = (value & 0x1) == 0x1;
-                self.bpip.use_opportunism = (value & 0x2) == 0x2;
+                let use_bpip = hash_val.get("use_bpip");
+                let use_opportunism = hash_val.get("use_opportunism");
+                self.bpip.used = use_bpip == Some(&0x1);
+                self.bpip.use_opportunism = use_opportunism == Some(&0x1);
                 RegisterEvent::None
             }
             "bpip::timeout" => {

mockups/tfhe-hpu-mockup/src/modules/ucore.rs

@@ -113,7 +113,7 @@ impl UCore {
                         patch_imm(iop, imm);
                         dop_patch
                     }
-                    // TODO Patch immediat
+                    // TODO Patch immediate
                     _ => dop_patch,
                 }
             })
@@ -126,7 +126,7 @@ impl UCore {
     }
 }
 
-/// Utility function to patch immediat argument
+/// Utility function to patch immediate argument
 fn patch_imm(iop: &hpu_asm::IOp, imm: &mut hpu_asm::ImmId) {
     *imm = match imm {
         hpu_asm::ImmId::Cst(val) => hpu_asm::ImmId::Cst(*val),

scripts/install_typos.sh

@@ -0,0 +1,64 @@
+#!/usr/bin/env bash
+
+set -e
+
+rust_toolchain=
+required_typos_version=""
+
+function usage() {
+    echo "$0: install typos-cli"
+    echo
+    echo "--help                    Print this message"
+    echo "--rust-toolchain          The toolchain to use"
+    echo "--typos-version           Version of typos-cli to install"
+    echo
+}
+
+while [ -n "$1" ]
+do
+    case "$1" in
+        "--rust-toolchain" )
+            shift
+            rust_toolchain="$1"
+            ;;
+
+        "--typos-version" )
+            shift
+            required_typos_version="$1"
+            ;;
+
+        *)
+            echo "Unknown param : $1"
+            exit 1
+            ;;
+    esac
+    shift
+done
+
+if [[ "${rust_toolchain::1}" != "+" ]]; then
+    rust_toolchain=${rust_toolchain:"+${rust_toolchain}"}
+fi
+
+if ! which typos ; then
+    cargo ${rust_toolchain:+"$rust_toolchain"} install --locked typos-cli --version ~"${required_typos_version}" || \
+    ( echo "Unable to install typos-cli, unknown error." && exit 1 )
+
+    exit 0
+fi
+
+ver_string="$(typos --version | cut -d ' ' -f 2)"
+
+ver_major="$(echo "${ver_string}" | cut -d '.' -f 1)"
+ver_minor="$(echo "${ver_string}" | cut -d '.' -f 2)"
+
+min_ver_major="$(echo "${required_typos_version}" | cut -d '.' -f 1)"
+min_ver_minor="$(echo "${required_typos_version}" | cut -d '.' -f 2)"
+
+if [[ "${ver_major}" -gt "${min_ver_major}" ]]; then
+    exit 0
+elif [[ "${ver_major}" -eq "${min_ver_major}" ]] && [[ "${ver_minor}" -ge "${min_ver_minor}" ]]; then
+    exit 0
+else
+    cargo ${rust_toolchain:+"$rust_toolchain"} install --locked typos-cli --version ~"${required_typos_version}" || \
+    ( echo "Unable to install typos-cli, unknown error." && exit 1 )
+fi

tests/Cargo.toml

@@ -10,6 +10,7 @@ tfhe-versionable = { path = "../utils/tfhe-versionable" }
 tfhe-backward-compat-data = { path = "../utils/tfhe-backward-compat-data", default-features = false, features = [
     "load",
 ] }
+rand = { workspace = true }
 cargo_toml = "0.22"
 
 

tests/backward_compatibility/high_level_api.rs

@@ -3,7 +3,9 @@ use crate::{load_and_unversionize, TestedModule};
 use std::path::Path;
 #[cfg(feature = "zk-pok")]
 use tfhe::integer::parameters::DynamicDistribution;
-use tfhe::prelude::{CiphertextList, FheDecrypt, FheEncrypt, ParameterSetConformant, SquashNoise};
+use tfhe::prelude::{
+    CiphertextList, FheDecrypt, FheEncrypt, ParameterSetConformant, ReRandomize, SquashNoise,
+};
 #[cfg(feature = "zk-pok")]
 use tfhe::shortint::parameters::{
     CompactCiphertextListExpansionKind, CompactPublicKeyEncryptionParameters,
@@ -17,7 +19,8 @@ use tfhe::{
     set_server_key, ClientKey, CompactCiphertextList, CompressedCiphertextList,
     CompressedCompactPublicKey, CompressedFheBool, CompressedFheInt8, CompressedFheUint8,
     CompressedPublicKey, CompressedServerKey, CompressedSquashedNoiseCiphertextList, FheBool,
-    FheInt8, FheUint8, SquashedNoiseFheBool, SquashedNoiseFheInt, SquashedNoiseFheUint,
+    FheInt8, FheUint8, ReRandomizationContext, SquashedNoiseFheBool, SquashedNoiseFheInt,
+    SquashedNoiseFheUint,
 };
 #[cfg(feature = "zk-pok")]
 use tfhe::{CompactPublicKey, ProvenCompactCiphertextList};
@@ -347,6 +350,7 @@ pub fn test_hl_pubkey(
         };
         FheUint8::encrypt(value, &public_key)
     };
+
     let decrypted: u8 = encrypted.decrypt(&client_key);
 
     if decrypted != value {
@@ -375,7 +379,7 @@ pub fn test_hl_serverkey(
     .map_err(|e| test.failure(e, format))?;
 
     let v1 = 73u8;
-    let ct1 = FheUint8::encrypt(v1, &client_key);
+    let mut ct1 = FheUint8::encrypt(v1, &client_key);
     let v2 = 102u8;
     let ct2 = FheUint8::encrypt(v2, &client_key);
 
@@ -386,7 +390,8 @@ pub fn test_hl_serverkey(
         load_and_unversionize(dir, test, format)?
     };
 
-    let has_noise_squashing = key.noise_squashing_key().is_some();
+    let has_noise_squashing = key.supports_noise_squashing();
+    let has_rerand = key.supports_ciphertext_re_randomization();
     set_server_key(key);
 
     if has_noise_squashing {
@@ -402,6 +407,47 @@ pub fn test_hl_serverkey(
         }
     }
 
+    if let Some(rerand_cpk_filename) = test.rerand_cpk_filename.as_ref() {
+        if has_rerand {
+            let rerand_cpk_file = dir.join(rerand_cpk_filename.to_string());
+            let public_key = CompressedCompactPublicKey::unversionize(
+                load_versioned_auxiliary(rerand_cpk_file).map_err(|e| test.failure(e, format))?,
+            )
+            .map_err(|e| test.failure(e, format))?
+            .decompress();
+
+            let nonce: [u8; 256 / 8] = rand::random();
+            let mut re_rand_context = ReRandomizationContext::new(
+                *b"TFHE_Rrd",
+                [b"FheUint8".as_slice(), nonce.as_slice()],
+                *b"TFHE_Enc",
+            );
+
+            re_rand_context.add_ciphertext(&ct1);
+            let mut seed_gen = re_rand_context.finalize();
+
+            ct1.re_randomize(&public_key, seed_gen.next_seed().unwrap())
+                .unwrap();
+
+            #[allow(clippy::eq_op)]
+            let rrd = &ct1 & &ct1;
+            let res: u8 = rrd.decrypt(&client_key);
+            if res != v1 {
+                return Err(test.failure(
+                    format!(
+                    "Invalid result for rerand using loaded server key, expected {v1} got {res}",
+                ),
+                    format,
+                ));
+            }
+        } else {
+            return Err(test.failure(
+                "Test requires rerand key but server key does not have it".to_string(),
+                format,
+            ));
+        }
+    }
+
     let ct_sum = ct1 + ct2;
     let sum: u8 = ct_sum.decrypt(&client_key);
 

tfhe-benchmark/benches/high_level_api/dex.rs

@@ -945,10 +945,10 @@ fn bench_swap_claim_throughput<FheType, F1, F2>(
         let bench_id =
             format!("{bench_name}::throughput::{fn_name}::{type_name}::{num_elems}_elems");
         group.bench_with_input(&bench_id, &num_elems, |b, &num_elems| {
-            let pendings_0_in = (0..num_elems)
+            let pending_0_in = (0..num_elems)
                 .map(|_| FheType::encrypt(rng.gen::<u64>(), client_key))
                 .collect::<Vec<_>>();
-            let pendings_1_in = (0..num_elems)
+            let pending_1_in = (0..num_elems)
                 .map(|_| FheType::encrypt(rng.gen::<u64>(), client_key))
                 .collect::<Vec<_>>();
             let total_tokens_0_in = (0..num_elems).map(|_| rng.gen::<u64>()).collect::<Vec<_>>();
@@ -969,9 +969,9 @@ fn bench_swap_claim_throughput<FheType, F1, F2>(
                 .collect::<Vec<_>>();
 
             b.iter(|| {
-                let (amounts_0_out, amounts_1_out): (Vec<_>, Vec<_>) = pendings_0_in
+                let (amounts_0_out, amounts_1_out): (Vec<_>, Vec<_>) = pending_0_in
                     .par_iter()
-                    .zip(pendings_1_in.par_iter())
+                    .zip(pending_1_in.par_iter())
                     .zip(total_tokens_0_in.par_iter())
                     .zip(total_tokens_1_in.par_iter())
                     .zip(total_tokens_0_out.par_iter())
@@ -1084,10 +1084,10 @@ fn cuda_bench_swap_claim_throughput<FheType, F1, F2>(
         let bench_id =
             format!("{bench_name}::throughput::{fn_name}::{type_name}::{num_elems}_elems");
         group.bench_with_input(&bench_id, &num_elems, |b, &num_elems| {
-            let pendings_0_in = (0..num_elems)
+            let pending_0_in = (0..num_elems)
                 .map(|_| FheType::encrypt(rng.gen::<u64>(), client_key))
                 .collect::<Vec<_>>();
-            let pendings_1_in = (0..num_elems)
+            let pending_1_in = (0..num_elems)
                 .map(|_| FheType::encrypt(rng.gen::<u64>(), client_key))
                 .collect::<Vec<_>>();
             let total_tokens_0_in = (0..num_elems).map(|_| rng.gen::<u64>()).collect::<Vec<_>>();
@@ -1116,9 +1116,9 @@ fn cuda_bench_swap_claim_throughput<FheType, F1, F2>(
             let num_streams_per_gpu = 2.min(num_elems / num_gpus);
             let chunk_size = (num_elems / num_gpus) as usize;
             b.iter(|| {
-                pendings_0_in
+                pending_0_in
                     .par_chunks(chunk_size)
-                    .zip(pendings_1_in.par_chunks(chunk_size))
+                    .zip(pending_1_in.par_chunks(chunk_size))
                     .zip(total_tokens_0_in.par_chunks(chunk_size))
                     .zip(total_tokens_1_in.par_chunks(chunk_size))
                     .zip(total_tokens_0_out.par_chunks(chunk_size))
@@ -1131,7 +1131,7 @@ fn cuda_bench_swap_claim_throughput<FheType, F1, F2>(
                                 (
                                     (
                                         (
-                                            (pendings_0_in_gpu_i, pendings_1_in_gpu_i),
+                                            (pending_0_in_gpu_i, pending_1_in_gpu_i),
                                             total_tokens_0_in_gpu_i,
                                         ),
                                         total_tokens_1_in_gpu_i,
@@ -1142,10 +1142,10 @@ fn cuda_bench_swap_claim_throughput<FheType, F1, F2>(
                             ),
                         )| {
                             let stream_chunk_size =
-                                pendings_0_in_gpu_i.len() / num_streams_per_gpu as usize;
-                            pendings_0_in_gpu_i
+                                pending_0_in_gpu_i.len() / num_streams_per_gpu as usize;
+                            pending_0_in_gpu_i
                                 .par_chunks(stream_chunk_size)
-                                .zip(pendings_1_in_gpu_i.par_chunks(stream_chunk_size))
+                                .zip(pending_1_in_gpu_i.par_chunks(stream_chunk_size))
                                 .zip(total_tokens_0_in_gpu_i.par_chunks(stream_chunk_size))
                                 .zip(total_tokens_1_in_gpu_i.par_chunks(stream_chunk_size))
                                 .zip(total_tokens_0_out_gpu_i.par_chunks(stream_chunk_size))

tfhe-benchmark/src/bin/shortint_key_sizes.rs

@@ -6,15 +6,16 @@ use std::io::Write;
 use std::path::Path;
 use tfhe::keycache::NamedParam;
 use tfhe::shortint::atomic_pattern::compressed::CompressedAtomicPatternServerKey;
+use tfhe::shortint::atomic_pattern::AtomicPatternServerKey;
 use tfhe::shortint::keycache::KEY_CACHE;
 use tfhe::shortint::list_compression::{
     NoiseSquashingCompressionKey, NoiseSquashingCompressionPrivateKey,
 };
 use tfhe::shortint::noise_squashing::{NoiseSquashingKey, NoiseSquashingPrivateKey};
-use tfhe::shortint::server_key::{StandardServerKey, StandardServerKeyView};
+use tfhe::shortint::server_key::StandardServerKeyView;
 use tfhe::shortint::{
     ClientKey, CompactPrivateKey, CompressedCompactPublicKey, CompressedKeySwitchingKey,
-    CompressedServerKey, PBSParameters,
+    CompressedServerKey, PBSParameters, ServerKey,
 };
 
 fn write_result(file: &mut File, name: &str, value: usize) {
@@ -189,50 +190,86 @@ fn tuniform_key_set_sizes(results_file: &Path) {
         let param_fhe_name = compute_param.name();
         let cks = ClientKey::new(compute_param);
         let compressed_sks = CompressedServerKey::new(&cks);
-        let sks = StandardServerKey::try_from(compressed_sks.decompress()).unwrap();
+        let sks = ServerKey::try_from(compressed_sks.decompress()).unwrap();
 
-        let std_compressed_ap_key = match &compressed_sks.compressed_ap_server_key {
-            CompressedAtomicPatternServerKey::Standard(
-                compressed_standard_atomic_pattern_server_key,
-            ) => compressed_standard_atomic_pattern_server_key,
-            CompressedAtomicPatternServerKey::KeySwitch32(_) => {
-                panic!("KS32 is unsupported to measure key sizes at the moment")
+        match &sks.atomic_pattern {
+            AtomicPatternServerKey::Standard(ap) => {
+                measure_serialized_size(
+                    &ap.key_switching_key,
+                    compute_param,
+                    &param_fhe_name,
+                    "ksk",
+                    "KSK",
+                    &mut file,
+                );
+                measure_serialized_size(
+                    &ap.bootstrapping_key,
+                    compute_param,
+                    &param_fhe_name,
+                    "bsk",
+                    "BSK",
+                    &mut file,
+                );
             }
-        };
+            AtomicPatternServerKey::KeySwitch32(ap) => {
+                measure_serialized_size(
+                    &ap.key_switching_key,
+                    compute_param,
+                    &param_fhe_name,
+                    "ksk",
+                    "KSK",
+                    &mut file,
+                );
+                measure_serialized_size(
+                    &ap.bootstrapping_key,
+                    compute_param,
+                    &param_fhe_name,
+                    "bsk",
+                    "BSK",
+                    &mut file,
+                );
+            }
+            AtomicPatternServerKey::Dynamic(_) => panic!("Dynamic atomic pattern not supported"),
+        }
 
-        measure_serialized_size(
-            &sks.atomic_pattern.key_switching_key,
-            compute_param,
-            &param_fhe_name,
-            "ksk",
-            "KSK",
-            &mut file,
-        );
-        measure_serialized_size(
-            std_compressed_ap_key.key_switching_key(),
-            compute_param,
-            &param_fhe_name,
-            "ksk_compressed",
-            "KSK",
-            &mut file,
-        );
-
-        measure_serialized_size(
-            &sks.atomic_pattern.bootstrapping_key,
-            compute_param,
-            &param_fhe_name,
-            "bsk",
-            "BSK",
-            &mut file,
-        );
-        measure_serialized_size(
-            &std_compressed_ap_key.bootstrapping_key(),
-            compute_param,
-            &param_fhe_name,
-            "bsk_compressed",
-            "BSK",
-            &mut file,
-        );
+        match &compressed_sks.compressed_ap_server_key {
+            CompressedAtomicPatternServerKey::Standard(comp_ap) => {
+                measure_serialized_size(
+                    comp_ap.key_switching_key(),
+                    compute_param,
+                    &param_fhe_name,
+                    "ksk_compressed",
+                    "KSK",
+                    &mut file,
+                );
+                measure_serialized_size(
+                    &comp_ap.bootstrapping_key(),
+                    compute_param,
+                    &param_fhe_name,
+                    "bsk_compressed",
+                    "BSK",
+                    &mut file,
+                );
+            }
+            CompressedAtomicPatternServerKey::KeySwitch32(comp_ap) => {
+                measure_serialized_size(
+                    comp_ap.key_switching_key(),
+                    compute_param,
+                    &param_fhe_name,
+                    "ksk_compressed",
+                    "KSK",
+                    &mut file,
+                );
+                measure_serialized_size(
+                    &comp_ap.bootstrapping_key(),
+                    compute_param,
+                    &param_fhe_name,
+                    "bsk_compressed",
+                    "BSK",
+                    &mut file,
+                );
+            }
+        }
 
         if let Some(dedicated_pke_params) = meta_params.dedicated_compact_public_key_parameters {
             let pke_param = dedicated_pke_params.pke_params;

tfhe-benchmark/src/utilities.rs

@@ -421,23 +421,32 @@ pub fn throughput_num_threads(num_block: usize, op_pbs_count: u64) -> u64 {
     let block_multiplicator = (ref_block_count as f64 / num_block as f64).ceil().min(1.0);
     // Some operations with a high serial workload (e.g. division) would yield an operation
     // loading value so low that the number of elements in the end wouldn't be meaningful.
-    let minimum_loading = if num_block < 64 { 0.2 } else { 0.01 };
+    let minimum_loading = if num_block < 64 { 1.0 } else { 0.015 };
 
     #[cfg(feature = "gpu")]
     {
         let num_sms_per_gpu = get_number_of_sms();
         let total_num_sm = num_sms_per_gpu * get_number_of_gpus();
 
-        let total_blocks_per_sm = 4u32; // Assume each SM can handle 4 blocks concurrently
-        let total_num_sm = total_blocks_per_sm * total_num_sm;
+        let total_blocks_per_sm = 4u64; // Assume each SM can handle 4 blocks concurrently
         let min_num_waves = 4u64; //Enforce at least 4 waves in the GPU
-        let elements_per_wave = total_num_sm as u64 / (num_block as u64);
-
+        let block_factor = ((2.0f64 * num_block as f64) / 4.0f64).ceil() as u64;
+        let elements_per_wave = total_blocks_per_sm * total_num_sm as u64 / block_factor;
+        // We need to enable the new load for pbs benches and for sizes larger than 16 blocks in
+        // demanding operations for the rest of operations we maintain a minimum of 200
+        // elements
+        let min_elements = if op_pbs_count == 1
+            || (op_pbs_count > (num_block * num_block) as u64 && num_block >= 16)
+        {
+            elements_per_wave * min_num_waves
+        } else {
+            200u64
+        };
         let operation_loading = ((total_num_sm as u64 / op_pbs_count) as f64).max(minimum_loading);
         let elements = (total_num_sm as f64 * block_multiplicator * operation_loading) as u64;
-        elements.min(elements_per_wave * min_num_waves) // This threshold is useful for operation
-                                                        // with both a small number of
-                                                        // block and low PBs count.
+        elements.min(min_elements) // This threshold is useful for operation
+                                   // with both a small number of
+                                   // block and low PBs count.
     }
     #[cfg(feature = "hpu")]
     {

tfhe-csprng/Cargo.toml

@@ -15,6 +15,8 @@ rust-version = "1.72"
 aes = "0.8.2"
 rayon = { workspace = true, optional = true }
 getrandom = { workspace = true }
+serde = "1.0.226"
+tfhe-versionable = { version = "0.6.2", path = "../utils/tfhe-versionable" }
 
 [target.'cfg(target_os = "macos")'.dependencies]
 libc = "0.2.133"

tfhe-csprng/src/seeders/backward_compatibility/mod.rs

@@ -0,0 +1,8 @@
+use tfhe_versionable::VersionsDispatch;
+
+use crate::seeders::XofSeed;
+
+#[derive(VersionsDispatch)]
+pub enum XofSeedVersions {
+    V0(XofSeed),
+}

tfhe-csprng/src/seeders/mod.rs

@@ -9,14 +9,15 @@
 #[derive(Debug, Copy, Clone, PartialEq, Eq)]
 pub struct Seed(pub u128);
 
-/// A Seed as described in the [NIST document]
+/// A Seed as described in the [Threshold (Fully) Homomorphic Encryption]
 ///
 /// This seed contains 2 information:
 /// * The domain separator bytes (ASCII string)
 /// * The seed bytes
 ///
-/// [NIST document]: https://eprint.iacr.org/2025/699
-#[derive(Debug, Clone, PartialEq, Eq)]
+/// [Threshold (Fully) Homomorphic Encryption]: https://eprint.iacr.org/2025/699
+#[derive(Debug, Clone, PartialEq, Eq, serde::Deserialize, serde::Serialize, Versionize)]
+#[versionize(XofSeedVersions)]
 pub struct XofSeed {
     // We store the domain separator concatenated with the seed bytes (str||seed)
     // as it makes it easier to create the iterator of u128 blocks
@@ -24,7 +25,7 @@ pub struct XofSeed {
 }
 
 impl XofSeed {
-    const DOMAIN_SEP_LEN: usize = 8;
+    pub const DOMAIN_SEP_LEN: usize = 8;
 
     // Creates a new seed of 128 bits
     pub fn new_u128(seed: u128, domain_separator: [u8; Self::DOMAIN_SEP_LEN]) -> Self {
@@ -68,13 +69,37 @@ impl XofSeed {
             u128::from_ne_bytes(buf)
         })
     }
+
+    /// Creates a new XofSeed from raw bytes.
+    ///
+    /// # Panics
+    ///
+    /// Panics if the provided data is smaller than the domain separator length
+    pub fn from_bytes(data: Vec<u8>) -> Self {
+        assert!(
+            data.len() >= Self::DOMAIN_SEP_LEN,
+            "XofSeed must be at least {} bytes long (got {})",
+            Self::DOMAIN_SEP_LEN,
+            data.len()
+        );
+        Self { data }
+    }
+
+    pub fn bytes(&self) -> &Vec<u8> {
+        &self.data
+    }
+
+    pub fn into_bytes(self) -> Vec<u8> {
+        self.data
+    }
 }
 
 pub enum SeedKind {
     /// Initializes the Aes-Ctr with a counter starting at 0
     /// and uses the seed as the Aes key.
     Ctr(Seed),
-    /// Seed that initialized the Aes-Ctr following the NIST document (see [XofSeed]).
+    /// Seed that initialized the Aes-Ctr following the Threshold (Fully) Homomorphic Encryption
+    /// document (see [XofSeed]).
     ///
     /// An Aes-Key and starting counter will be derived from the XofSeed, to
     /// then initialize the Aes-Ctr random generator
@@ -105,11 +130,15 @@ pub trait Seeder {
         Self: Sized;
 }
 
+pub mod backward_compatibility;
 mod implem;
 // This import statement can be empty if seeder features are disabled, rustc's behavior changed to
 // warn of empty modules, we know this can happen, so allow it.
 #[allow(unused_imports)]
 pub use implem::*;
+use tfhe_versionable::Versionize;
+
+use crate::seeders::backward_compatibility::XofSeedVersions;
 
 #[cfg(test)]
 mod generic_tests {

tfhe-zk-pok/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "tfhe-zk-pok"
-version = "0.7.2"
+version = "0.7.4"
 edition = "2021"
 keywords = ["zero", "knowledge", "proof", "vector-commitments"]
 homepage = "https://zama.ai/"
@@ -23,7 +23,7 @@ sha3 = "0.10.8"
 serde = { workspace = true, features = ["default", "derive"] }
 zeroize = "1.7.0"
 num-bigint = "0.4.5"
-tfhe-versionable = { version = "0.6.1", path = "../utils/tfhe-versionable" }
+tfhe-versionable = { version = "0.6.2", path = "../utils/tfhe-versionable" }
 
 [features]
 experimental = []

tfhe-zk-pok/src/curve_446/mod.rs

@@ -416,3 +416,27 @@ pub mod g2 {
     /// 107680854723992552431070996218129928499826544031468382031848626814251381379173928074140221537929995580031433096217223703806029068859074
     pub const G2_GENERATOR_Y_C1: Fq = MontFp!("107680854723992552431070996218129928499826544031468382031848626814251381379173928074140221537929995580031433096217223703806029068859074");
 }
+
+#[cfg(test)]
+mod tests {
+    use crate::serialization::{InvalidFpError, SerializableFp};
+
+    use super::Fq;
+    use ark_ff::Field;
+
+    #[test]
+    fn test_serialization() {
+        // This one is only used to have the correct number of serialized bytes
+        let a = Fq::ONE;
+        let s: SerializableFp = a.into();
+
+        let mut data = vec![];
+        bincode::serialize_into(&mut data, &s).unwrap();
+        // First u64 is the vec size
+        data[std::mem::size_of::<u64>()..].fill(u8::MAX);
+
+        let s2: SerializableFp = bincode::deserialize(&data).unwrap();
+        let a2 = Fq::try_from(s2);
+        assert!(matches!(a2, Err(InvalidFpError::GreaterThanModulus)));
+    }
+}

tfhe-zk-pok/src/curve_api/bls12_381.rs

@@ -529,7 +529,7 @@ mod g2 {
 }
 
 mod gt {
-    use crate::serialization::InvalidArraySizeError;
+    use crate::serialization::InvalidFpError;
 
     use super::*;
     use ark_ec::pairing::Pairing;
@@ -550,7 +550,7 @@ mod gt {
     }
 
     impl TryFrom<SerializableFp12> for Gt {
-        type Error = InvalidArraySizeError;
+        type Error = InvalidFpError;
 
         fn try_from(value: SerializableFp12) -> Result<Self, Self::Error> {
             Ok(Self {
@@ -682,7 +682,7 @@ mod gt {
 }
 
 mod zp {
-    use crate::serialization::InvalidArraySizeError;
+    use crate::serialization::InvalidFpError;
 
     use super::*;
     use ark_ff::Fp;
@@ -737,7 +737,7 @@ mod zp {
         }
     }
     impl TryFrom<SerializableFp> for Zp {
-        type Error = InvalidArraySizeError;
+        type Error = InvalidFpError;
 
         fn try_from(value: SerializableFp) -> Result<Self, Self::Error> {
             Ok(Self {

tfhe-zk-pok/src/curve_api/bls12_446.rs

@@ -958,7 +958,7 @@ mod gt {
 mod zp {
     use super::*;
     use crate::curve_446::FrConfig;
-    use crate::serialization::InvalidArraySizeError;
+    use crate::serialization::InvalidFpError;
     use ark_ff::{Fp, FpConfig, MontBackend, PrimeField};
     use tfhe_versionable::Versionize;
     use zeroize::{Zeroize, ZeroizeOnDrop};
@@ -1013,7 +1013,7 @@ mod zp {
         }
     }
     impl TryFrom<SerializableFp> for Zp {
-        type Error = InvalidArraySizeError;
+        type Error = InvalidFpError;
 
         fn try_from(value: SerializableFp) -> Result<Self, Self::Error> {
             Ok(Self {

tfhe-zk-pok/src/serialization.rs

@@ -12,7 +12,9 @@ use crate::backward_compatibility::{
 };
 use ark_ec::short_weierstrass::{Affine, SWCurveConfig};
 use ark_ec::AffineRepr;
-use ark_ff::{BigInt, Field, Fp, Fp2, Fp6, Fp6Config, FpConfig, QuadExtConfig, QuadExtField};
+use ark_ff::{
+    BigInt, Field, Fp, Fp2, Fp6, Fp6Config, FpConfig, PrimeField, QuadExtConfig, QuadExtField,
+};
 use serde::{Deserialize, Serialize};
 use tfhe_versionable::Versionize;
 
@@ -47,6 +49,34 @@ impl Display for InvalidArraySizeError {
 
 impl Error for InvalidArraySizeError {}
 
+#[derive(Debug)]
+pub enum InvalidFpError {
+    InvalidArraySizeError(InvalidArraySizeError),
+    GreaterThanModulus,
+}
+
+impl From<InvalidArraySizeError> for InvalidFpError {
+    fn from(value: InvalidArraySizeError) -> Self {
+        Self::InvalidArraySizeError(value)
+    }
+}
+
+impl Display for InvalidFpError {
+    fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
+        match self {
+            Self::InvalidArraySizeError(e) => e.fmt(f),
+            Self::GreaterThanModulus => {
+                write!(
+                    f,
+                    "The deserialized value was bigger than what its type modulus allowed"
+                )
+            }
+        }
+    }
+}
+
+impl Error for InvalidFpError {}
+
 /// Tries to convert a Vec into a constant size array, and returns an [`InvalidArraySizeError`] if
 /// the size does not match
 pub(crate) fn try_vec_to_array<T, const N: usize>(
@@ -77,16 +107,20 @@ impl<P: FpConfig<N>, const N: usize> From<Fp<P, N>> for SerializableFp {
 }
 
 impl<P: FpConfig<N>, const N: usize> TryFrom<SerializableFp> for Fp<P, N> {
-    type Error = InvalidArraySizeError;
+    type Error = InvalidFpError;
 
     fn try_from(value: SerializableFp) -> Result<Self, Self::Error> {
-        Ok(Fp(BigInt(try_vec_to_array(value.val)?), PhantomData))
+        let fp = BigInt(try_vec_to_array(value.val)?);
+        if fp >= Fp::<P, N>::MODULUS {
+            return Err(InvalidFpError::GreaterThanModulus);
+        }
+        Ok(Fp(fp, PhantomData))
     }
 }
 
 #[derive(Debug)]
 pub enum InvalidSerializedAffineError {
-    InvalidFp(InvalidArraySizeError),
+    InvalidFp(InvalidFpError),
     InvalidCompressedXCoordinate,
 }
 
@@ -115,8 +149,8 @@ impl Error for InvalidSerializedAffineError {
     }
 }
 
-impl From<InvalidArraySizeError> for InvalidSerializedAffineError {
-    fn from(value: InvalidArraySizeError) -> Self {
+impl From<InvalidFpError> for InvalidSerializedAffineError {
+    fn from(value: InvalidFpError) -> Self {
         Self::InvalidFp(value)
     }
 }
@@ -163,7 +197,7 @@ impl<F> SerializableAffine<F> {
 
 impl<F, C: SWCurveConfig> TryFrom<SerializableAffine<F>> for Affine<C>
 where
-    F: TryInto<C::BaseField, Error = InvalidArraySizeError>,
+    F: TryInto<C::BaseField, Error = InvalidFpError>,
 {
     type Error = InvalidSerializedAffineError;
 
@@ -207,9 +241,9 @@ where
 
 impl<F, P: QuadExtConfig> TryFrom<SerializableQuadExtField<F>> for QuadExtField<P>
 where
-    F: TryInto<P::BaseField, Error = InvalidArraySizeError>,
+    F: TryInto<P::BaseField, Error = InvalidFpError>,
 {
-    type Error = InvalidArraySizeError;
+    type Error = InvalidFpError;
 
     fn try_from(value: SerializableQuadExtField<F>) -> Result<Self, Self::Error> {
         Ok(QuadExtField {
@@ -244,9 +278,9 @@ where
 
 impl<F, P6: Fp6Config> TryFrom<SerializableCubicExtField<F>> for Fp6<P6>
 where
-    F: TryInto<Fp2<P6::Fp2Config>, Error = InvalidArraySizeError>,
+    F: TryInto<Fp2<P6::Fp2Config>, Error = InvalidFpError>,
 {
-    type Error = InvalidArraySizeError;
+    type Error = InvalidFpError;
 
     fn try_from(value: SerializableCubicExtField<F>) -> Result<Self, Self::Error> {
         Ok(Fp6 {

tfhe/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "tfhe"
-version = "1.4.0"
+version = "1.4.0-alpha.4"
 edition = "2021"
 readme = "../README.md"
 keywords = ["fully", "homomorphic", "encryption", "fhe", "cryptography"]
@@ -65,18 +65,20 @@ tfhe-fft = { version = "0.9.0", path = "../tfhe-fft", features = [
 ] }
 tfhe-ntt = { version = "0.6.1", path = "../tfhe-ntt" }
 pulp = { workspace = true, features = ["default"] }
-tfhe-cuda-backend = { version = "0.11.0", path = "../backends/tfhe-cuda-backend", optional = true }
+tfhe-cuda-backend = { version = "0.12.0-alpha.2", path = "../backends/tfhe-cuda-backend", optional = true }
 aligned-vec = { workspace = true, features = ["default", "serde"] }
 dyn-stack = { workspace = true, features = ["default"] }
 paste = "1.0.7"
 fs2 = { version = "0.4.3", optional = true }
-# Used for OPRF in shortint
+# Used for OPRF in shortint and rerand
 sha3 = { version = "0.10", optional = true }
+blake3 = { version = "1.8", optional = true }
+
 itertools = { workspace = true }
 rand_core = { version = "0.6.4", features = ["std"] }
 strum = { version = "0.27", features = ["derive"], optional = true }
-tfhe-zk-pok = { version = "0.7.2", path = "../tfhe-zk-pok", optional = true }
-tfhe-versionable = { version = "0.6.1", path = "../utils/tfhe-versionable" }
+tfhe-zk-pok = { version = "0.7.4", path = "../tfhe-zk-pok", optional = true }
+tfhe-versionable = { version = "0.6.2", path = "../utils/tfhe-versionable" }
 
 # wasm deps
 wasm-bindgen = { workspace = true, features = [
@@ -93,7 +95,7 @@ tfhe-hpu-backend = { version = "0.2", path = "../backends/tfhe-hpu-backend", opt
 
 [features]
 boolean = []
-shortint = ["dep:sha3"]
+shortint = ["dep:sha3", "dep:blake3"]
 integer = ["shortint", "dep:strum"]
 strings = ["integer"]
 internal-keycache = ["dep:fs2"]

tfhe/docs/getting-started/quick-start.md

@@ -127,7 +127,7 @@ $ cargo run --release
 
 You can learn more about homomorphic types and associated compilation features in the [configuration documentation.](../configuration/rust-configuration.md)
 
-## Perforance tips
+## Performance tips
 Performance can be further improved by setting `lto="fat"` in `Cargo.toml`
 ```toml
 [profile.release]

tfhe/examples/hpu/bench.rs

@@ -56,7 +56,7 @@ pub struct Args {
     #[arg(long, value_parser = maybe_hex::<u128>)]
     pub src: Vec<u128>,
 
-    /// Force immediat input values
+    /// Force immediate input values
     #[arg(long, value_parser = maybe_hex::<u128>)]
     pub imm: Vec<u128>,
 

tfhe/src/core_crypto/algorithms/lwe_bootstrap_key_generation.rs

@@ -11,6 +11,7 @@ use crate::core_crypto::commons::parameters::*;
 use crate::core_crypto::commons::traits::*;
 use crate::core_crypto::entities::*;
 use rayon::prelude::*;
+use tfhe_csprng::seeders::Seed;
 
 /// Fill an [`LWE bootstrap key`](`LweBootstrapKey`) with an actual bootstrapping key constructed
 /// from an input key [`LWE secret key`](`LweSecretKey`) and an output key
@@ -390,6 +391,41 @@ pub fn generate_seeded_lwe_bootstrap_key<
     OutputCont: ContainerMut<Element = OutputScalar>,
     // Maybe Sized allows to pass Box<dyn Seeder>.
     NoiseSeeder: Seeder + ?Sized,
+{
+    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
+        output.compression_seed().seed,
+        noise_seeder,
+    );
+
+    generate_seeded_lwe_bootstrap_key_with_pre_seeded_generator(
+        input_lwe_secret_key,
+        output_glwe_secret_key,
+        output,
+        noise_distribution,
+        &mut generator,
+    )
+}
+
+pub fn generate_seeded_lwe_bootstrap_key_with_pre_seeded_generator<
+    NoiseDistribution,
+    InputKeyCont,
+    OutputKeyCont,
+    OutputCont,
+    ByteGen,
+>(
+    input_lwe_secret_key: &LweSecretKey<InputKeyCont>,
+    output_glwe_secret_key: &GlweSecretKey<OutputKeyCont>,
+    output: &mut SeededLweBootstrapKey<OutputCont>,
+    noise_distribution: NoiseDistribution,
+    generator: &mut EncryptionRandomGenerator<ByteGen>,
+) where
+    NoiseDistribution: Distribution,
+    InputKeyCont: Container,
+    OutputKeyCont: Container,
+    OutputCont: ContainerMut<Element = OutputKeyCont::Element>,
+    InputKeyCont::Element: Copy + CastInto<OutputKeyCont::Element>,
+    OutputKeyCont::Element: Encryptable<Uniform, NoiseDistribution>,
+    ByteGen: ByteRandomGenerator,
 {
     assert!(
         output.input_lwe_dimension() == input_lwe_secret_key.lwe_dimension(),
@@ -415,11 +451,6 @@ pub fn generate_seeded_lwe_bootstrap_key<
         output.polynomial_size()
     );
 
-    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
-        output.compression_seed().seed,
-        noise_seeder,
-    );
-
     let gen_iter = generator
         .try_fork_from_config(output.encryption_fork_config(Uniform, noise_distribution))
         .unwrap();
@@ -492,6 +523,49 @@ where
     bsk
 }
 
+pub fn allocate_and_generate_lwe_bootstrapping_key_with_pre_seeded_generator<
+    LweCont,
+    GlweCont,
+    ByteGen,
+>(
+    input_lwe_secret_key: &LweSecretKey<LweCont>,
+    output_glwe_secret_key: &GlweSecretKey<GlweCont>,
+    decomp_base_log: DecompositionBaseLog,
+    decomp_level_count: DecompositionLevelCount,
+    noise_distribution: DynamicDistribution<GlweCont::Element>,
+    ciphertext_modulus: CiphertextModulus<GlweCont::Element>,
+    generator: &mut EncryptionRandomGenerator<ByteGen>,
+) -> SeededLweBootstrapKeyOwned<GlweCont::Element>
+where
+    LweCont: Container,
+    GlweCont: Container,
+    LweCont::Element: Copy + CastInto<GlweCont::Element>,
+    GlweCont::Element:
+        UnsignedInteger + Encryptable<Uniform, DynamicDistribution<GlweCont::Element>>,
+    ByteGen: ByteRandomGenerator,
+{
+    let mut lwe_bootstrapping_key = SeededLweBootstrapKeyOwned::new(
+        GlweCont::Element::ZERO,
+        output_glwe_secret_key.glwe_dimension().to_glwe_size(),
+        output_glwe_secret_key.polynomial_size(),
+        decomp_base_log,
+        decomp_level_count,
+        input_lwe_secret_key.lwe_dimension(),
+        CompressionSeed::from(Seed(0)),
+        ciphertext_modulus,
+    );
+
+    generate_seeded_lwe_bootstrap_key_with_pre_seeded_generator(
+        input_lwe_secret_key,
+        output_glwe_secret_key,
+        &mut lwe_bootstrapping_key,
+        noise_distribution,
+        generator,
+    );
+
+    lwe_bootstrapping_key
+}
+
 /// Parallel variant of [`generate_seeded_lwe_bootstrap_key`], it is recommended to use this
 /// function for better key generation times as LWE bootstrapping keys can be quite large.
 pub fn par_generate_seeded_lwe_bootstrap_key<

tfhe/src/core_crypto/algorithms/lwe_compact_public_key_generation.rs

@@ -110,6 +110,41 @@ pub fn generate_seeded_lwe_compact_public_key<
     OutputKeyCont: ContainerMut<Element = Scalar>,
     // Maybe Sized allows to pass Box<dyn Seeder>.
     NoiseSeeder: Seeder + ?Sized,
+{
+    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
+        output.compression_seed().seed,
+        noise_seeder,
+    );
+
+    generate_seeded_lwe_compact_public_key_with_pre_seeded_generator(
+        lwe_secret_key,
+        output,
+        noise_distribution,
+        &mut generator,
+    );
+}
+
+/// Fill a [`seeded LWE compact public key`](`LweCompactPublicKey`) with an actual public key
+/// constructed from a private [`LWE secret key`](`LweSecretKey`).
+///
+/// This uses an already seeded generator
+pub fn generate_seeded_lwe_compact_public_key_with_pre_seeded_generator<
+    Scalar,
+    NoiseDistribution,
+    InputKeyCont,
+    OutputKeyCont,
+    ByteGen,
+>(
+    lwe_secret_key: &LweSecretKey<InputKeyCont>,
+    output: &mut SeededLweCompactPublicKey<OutputKeyCont>,
+    noise_distribution: NoiseDistribution,
+    generator: &mut EncryptionRandomGenerator<ByteGen>,
+) where
+    Scalar: Encryptable<Uniform, NoiseDistribution>,
+    NoiseDistribution: Distribution,
+    InputKeyCont: Container<Element = Scalar>,
+    OutputKeyCont: ContainerMut<Element = Scalar>,
+    ByteGen: ByteRandomGenerator,
 {
     assert!(
         output.ciphertext_modulus().is_native_modulus(),
@@ -124,11 +159,6 @@ pub fn generate_seeded_lwe_compact_public_key<
         output.lwe_dimension()
     );
 
-    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
-        output.compression_seed().seed,
-        noise_seeder,
-    );
-
     let mut tmp_mask = vec![Scalar::ZERO; output.lwe_dimension().0];
     generator.fill_slice_with_random_uniform_mask(tmp_mask.as_mut());
 

tfhe/src/core_crypto/algorithms/lwe_encryption.rs

@@ -4,7 +4,9 @@
 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::core_crypto::commons::generators::{
+    EncryptionRandomGenerator, NoiseRandomGenerator, SecretRandomGenerator,
+};
 #[cfg(feature = "zk-pok")]
 use crate::core_crypto::commons::math::random::BoundedDistribution;
 use crate::core_crypto::commons::math::random::{
@@ -1852,13 +1854,6 @@ where
     KeyCont: Container<Element = Scalar>,
 {
     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 \
-            of the expected bounds of given to the CRS"
-                .into(),
-        );
-    }
 
     if mask_noise_distribution.contains(exclusive_max.cast_into()) {
         // The proof expect noise bound between [-b, b) (aka -b..b)
@@ -2301,7 +2296,6 @@ fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl<
     OutputCont,
     MaskDistribution,
     NoiseDistribution,
-    SecretGen,
     EncryptionGen,
 >(
     lwe_compact_public_key: &LweCompactPublicKey<KeyCont>,
@@ -2309,8 +2303,7 @@ fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl<
     encoded: &PlaintextList<InputCont>,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
 ) -> CompactPublicKeyRandomVectors<Scalar>
 where
     Scalar: Encryptable<MaskDistribution, NoiseDistribution> + RandomGenerable<UniformBinary>,
@@ -2319,7 +2312,6 @@ where
     OutputCont: ContainerMut<Element = Scalar>,
     MaskDistribution: Distribution,
     NoiseDistribution: Distribution,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ByteRandomGenerator,
 {
     assert!(
@@ -2355,14 +2347,14 @@ where
     let (mut output_mask_list, mut output_body_list) = output.get_mut_mask_and_body_list();
 
     let mut binary_random_vector = vec![Scalar::ZERO; output_mask_list.lwe_mask_list_size()];
-    secret_generator.fill_slice_with_random_uniform_binary(&mut binary_random_vector);
+    noise_generator.fill_slice_with_random_uniform_binary_bits(&mut binary_random_vector);
 
     let mut mask_noise = vec![Scalar::ZERO; output_mask_list.lwe_mask_list_size()];
-    encryption_generator
+    noise_generator
         .fill_slice_with_random_noise_from_distribution(&mut mask_noise, mask_noise_distribution);
 
     let mut body_noise = vec![Scalar::ZERO; encoded.plaintext_count().0];
-    encryption_generator
+    noise_generator
         .fill_slice_with_random_noise_from_distribution(&mut body_noise, body_noise_distribution);
 
     let max_ciphertext_per_bin = lwe_compact_public_key.lwe_dimension().0;
@@ -2452,9 +2444,9 @@ where
 /// // Create the PRNG
 /// let mut seeder = new_seeder();
 /// let seeder = seeder.as_mut();
+/// let mut secret_generator = SecretRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed());
 /// let mut encryption_generator =
 ///     EncryptionRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed(), seeder);
-/// let mut secret_generator = SecretRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed());
 ///
 /// // Create the LweSecretKey
 /// let lwe_secret_key =
@@ -2489,8 +2481,7 @@ where
 ///     &input_plaintext_list,
 ///     glwe_noise_distribution,
 ///     glwe_noise_distribution,
-///     &mut secret_generator,
-///     &mut encryption_generator,
+///     encryption_generator.noise_generator_mut(),
 /// );
 ///
 /// let mut output_plaintext_list = input_plaintext_list.clone();
@@ -2522,7 +2513,6 @@ pub fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     KeyCont,
     InputCont,
     OutputCont,
-    SecretGen,
     EncryptionGen,
 >(
     lwe_compact_public_key: &LweCompactPublicKey<KeyCont>,
@@ -2530,8 +2520,7 @@ pub fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     encoded: &PlaintextList<InputCont>,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
 ) where
     Scalar: Encryptable<MaskDistribution, NoiseDistribution> + RandomGenerable<UniformBinary>,
     MaskDistribution: Distribution,
@@ -2539,7 +2528,6 @@ pub fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     KeyCont: Container<Element = Scalar>,
     InputCont: Container<Element = Scalar>,
     OutputCont: ContainerMut<Element = Scalar>,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ByteRandomGenerator,
 {
     let _ = encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl(
@@ -2548,8 +2536,7 @@ pub fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
         encoded,
         mask_noise_distribution,
         body_noise_distribution,
-        secret_generator,
-        encryption_generator,
+        noise_generator,
     );
 }
 
@@ -2638,8 +2625,7 @@ pub fn encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
 ///     delta,
 ///     glwe_noise_distribution,
 ///     glwe_noise_distribution,
-///     &mut secret_generator,
-///     &mut encryption_generator,
+///     encryption_generator.noise_generator_mut(),
 ///     &mut random_generator,
 ///     &crs,
 ///     &metadata,
@@ -2690,7 +2676,6 @@ pub fn encrypt_and_prove_lwe_compact_ciphertext_list_with_compact_public_key<
     OutputCont,
     MaskDistribution,
     NoiseDistribution,
-    SecretGen,
     EncryptionGen,
     G,
 >(
@@ -2700,8 +2685,7 @@ pub fn encrypt_and_prove_lwe_compact_ciphertext_list_with_compact_public_key<
     delta: Scalar,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
     random_generator: &mut RandomGenerator<G>,
     crs: &CompactPkeCrs,
     metadata: &[u8],
@@ -2719,7 +2703,6 @@ where
     KeyCont: Container<Element = Scalar>,
     InputCont: Container<Element = Scalar>,
     OutputCont: ContainerMut<Element = Scalar>,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ByteRandomGenerator,
     G: ByteRandomGenerator,
 {
@@ -2752,8 +2735,7 @@ where
         &encoded,
         mask_noise_distribution,
         body_noise_distribution,
-        secret_generator,
-        encryption_generator,
+        noise_generator,
     );
 
     Ok(crs.prove(
@@ -2776,7 +2758,6 @@ fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl<
     OutputCont,
     MaskDistribution,
     NoiseDistribution,
-    SecretGen,
     EncryptionGen,
 >(
     lwe_compact_public_key: &LweCompactPublicKey<KeyCont>,
@@ -2784,8 +2765,7 @@ fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl<
     encoded: &PlaintextList<InputCont>,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
 ) -> CompactPublicKeyRandomVectors<Scalar>
 where
     Scalar: Encryptable<MaskDistribution, NoiseDistribution> + RandomGenerable<UniformBinary>,
@@ -2794,7 +2774,6 @@ where
     OutputCont: ContainerMut<Element = Scalar>,
     MaskDistribution: Distribution,
     NoiseDistribution: Distribution,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ByteRandomGenerator,
 {
     assert!(
@@ -2830,14 +2809,14 @@ where
     let (mut output_mask_list, mut output_body_list) = output.get_mut_mask_and_body_list();
 
     let mut binary_random_vector = vec![Scalar::ZERO; output_mask_list.lwe_mask_list_size()];
-    secret_generator.fill_slice_with_random_uniform_binary(&mut binary_random_vector);
+    noise_generator.fill_slice_with_random_uniform_binary_bits(&mut binary_random_vector);
 
     let mut mask_noise = vec![Scalar::ZERO; output_mask_list.lwe_mask_list_size()];
-    encryption_generator
+    noise_generator
         .fill_slice_with_random_noise_from_distribution(&mut mask_noise, mask_noise_distribution);
 
     let mut body_noise = vec![Scalar::ZERO; encoded.plaintext_count().0];
-    encryption_generator
+    noise_generator
         .fill_slice_with_random_noise_from_distribution(&mut body_noise, body_noise_distribution);
 
     let max_ciphertext_per_bin = lwe_compact_public_key.lwe_dimension().0;
@@ -2932,11 +2911,11 @@ where
 /// // Create the PRNG
 /// let mut seeder = new_seeder();
 /// let seeder = seeder.as_mut();
+/// let mut secret_generator = SecretRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed());
 /// let mut encryption_generator =
 ///     EncryptionRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed(), seeder);
-/// let mut secret_generator = SecretRandomGenerator::<DefaultRandomGenerator>::new(seeder.seed());
 ///
-/// // Create the LweSecretKey
+/// // create the LweSecretKey
 /// let lwe_secret_key =
 ///     allocate_and_generate_new_binary_lwe_secret_key(lwe_dimension, &mut secret_generator);
 ///
@@ -2969,8 +2948,7 @@ where
 ///     &input_plaintext_list,
 ///     glwe_noise_distribution,
 ///     glwe_noise_distribution,
-///     &mut secret_generator,
-///     &mut encryption_generator,
+///     encryption_generator.noise_generator_mut(),
 /// );
 ///
 /// let mut output_plaintext_list = input_plaintext_list.clone();
@@ -3002,7 +2980,6 @@ pub fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     KeyCont,
     InputCont,
     OutputCont,
-    SecretGen,
     EncryptionGen,
 >(
     lwe_compact_public_key: &LweCompactPublicKey<KeyCont>,
@@ -3010,8 +2987,7 @@ pub fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     encoded: &PlaintextList<InputCont>,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
 ) where
     Scalar: Encryptable<MaskDistribution, NoiseDistribution>
         + RandomGenerable<UniformBinary>
@@ -3022,7 +2998,6 @@ pub fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
     KeyCont: Container<Element = Scalar>,
     InputCont: Container<Element = Scalar>,
     OutputCont: ContainerMut<Element = Scalar>,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ParallelByteRandomGenerator,
 {
     let _ = par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key_impl(
@@ -3031,8 +3006,7 @@ pub fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
         encoded,
         mask_noise_distribution,
         body_noise_distribution,
-        secret_generator,
-        encryption_generator,
+        noise_generator,
     );
 }
 
@@ -3122,8 +3096,7 @@ pub fn par_encrypt_lwe_compact_ciphertext_list_with_compact_public_key<
 ///     delta,
 ///     glwe_noise_distribution,
 ///     glwe_noise_distribution,
-///     &mut secret_generator,
-///     &mut encryption_generator,
+///     encryption_generator.noise_generator_mut(),
 ///     &mut random_generator,
 ///     &crs,
 ///     &metadata,
@@ -3174,7 +3147,6 @@ pub fn par_encrypt_and_prove_lwe_compact_ciphertext_list_with_compact_public_key
     OutputCont,
     MaskDistribution,
     NoiseDistribution,
-    SecretGen,
     EncryptionGen,
     G,
 >(
@@ -3184,8 +3156,7 @@ pub fn par_encrypt_and_prove_lwe_compact_ciphertext_list_with_compact_public_key
     delta: Scalar,
     mask_noise_distribution: MaskDistribution,
     body_noise_distribution: NoiseDistribution,
-    secret_generator: &mut SecretRandomGenerator<SecretGen>,
-    encryption_generator: &mut EncryptionRandomGenerator<EncryptionGen>,
+    noise_generator: &mut NoiseRandomGenerator<EncryptionGen>,
     random_generator: &mut RandomGenerator<G>,
     crs: &CompactPkeCrs,
     metadata: &[u8],
@@ -3203,7 +3174,6 @@ where
     KeyCont: Container<Element = Scalar>,
     InputCont: Container<Element = Scalar>,
     OutputCont: ContainerMut<Element = Scalar>,
-    SecretGen: ByteRandomGenerator,
     EncryptionGen: ByteRandomGenerator,
     G: ByteRandomGenerator,
 {
@@ -3236,8 +3206,7 @@ where
         &encoded,
         mask_noise_distribution,
         body_noise_distribution,
-        secret_generator,
-        encryption_generator,
+        noise_generator,
     );
 
     Ok(crs.prove(
@@ -3389,8 +3358,7 @@ mod test {
                     &input_plaintext_list,
                     glwe_noise_distribution,
                     glwe_noise_distribution,
-                    &mut secret_random_generator,
-                    &mut encryption_random_generator,
+                    encryption_random_generator.noise_generator_mut(),
                 );
 
                 let mut output_plaintext_list = input_plaintext_list.clone();
@@ -3456,8 +3424,7 @@ mod test {
                     &input_plaintext_list,
                     glwe_noise_distribution,
                     glwe_noise_distribution,
-                    &mut secret_random_generator,
-                    &mut encryption_random_generator,
+                    encryption_random_generator.noise_generator_mut(),
                 );
 
                 let mut output_plaintext_list = input_plaintext_list.clone();

tfhe/src/core_crypto/algorithms/lwe_keyswitch_key_generation.rs

@@ -13,6 +13,7 @@ use crate::core_crypto::commons::math::random::{
 use crate::core_crypto::commons::parameters::*;
 use crate::core_crypto::commons::traits::*;
 use crate::core_crypto::entities::*;
+use tfhe_csprng::seeders::Seed;
 
 /// Fill an [`LWE keyswitch key`](`LweKeyswitchKey`) with an actual keyswitching key constructed
 /// from an input and an output key [`LWE secret key`](`LweSecretKey`).
@@ -407,6 +408,43 @@ pub fn generate_seeded_lwe_keyswitch_key<
     KSKeyCont: ContainerMut<Element = OutputScalar>,
     // Maybe Sized allows to pass Box<dyn Seeder>.
     NoiseSeeder: Seeder + ?Sized,
+{
+    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
+        lwe_keyswitch_key.compression_seed().seed,
+        noise_seeder,
+    );
+
+    generate_seeded_lwe_keyswitch_key_with_pre_seeded_generator(
+        input_lwe_sk,
+        output_lwe_sk,
+        lwe_keyswitch_key,
+        noise_distribution,
+        &mut generator,
+    )
+}
+
+pub fn generate_seeded_lwe_keyswitch_key_with_pre_seeded_generator<
+    InputScalar,
+    OutputScalar,
+    NoiseDistribution,
+    InputKeyCont,
+    OutputKeyCont,
+    KSKeyCont,
+    ByteGen,
+>(
+    input_lwe_sk: &LweSecretKey<InputKeyCont>,
+    output_lwe_sk: &LweSecretKey<OutputKeyCont>,
+    lwe_keyswitch_key: &mut SeededLweKeyswitchKey<KSKeyCont>,
+    noise_distribution: NoiseDistribution,
+    generator: &mut EncryptionRandomGenerator<ByteGen>,
+) where
+    InputScalar: UnsignedInteger + CastInto<OutputScalar>,
+    OutputScalar: Encryptable<Uniform, NoiseDistribution>,
+    NoiseDistribution: Distribution,
+    InputKeyCont: Container<Element = InputScalar>,
+    OutputKeyCont: Container<Element = OutputScalar>,
+    KSKeyCont: ContainerMut<Element = OutputScalar>,
+    ByteGen: ByteRandomGenerator,
 {
     assert!(
         lwe_keyswitch_key.input_key_lwe_dimension() == input_lwe_sk.lwe_dimension(),
@@ -439,11 +477,6 @@ pub fn generate_seeded_lwe_keyswitch_key<
     let mut decomposition_plaintexts_buffer =
         PlaintextListOwned::new(OutputScalar::ZERO, PlaintextCount(decomp_level_count.0));
 
-    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
-        lwe_keyswitch_key.compression_seed().seed,
-        noise_seeder,
-    );
-
     // Iterate over the input key elements and the destination lwe_keyswitch_key memory
     for (input_key_element, mut keyswitch_key_block) in input_lwe_sk
         .as_ref()
@@ -473,7 +506,7 @@ pub fn generate_seeded_lwe_keyswitch_key<
             &mut keyswitch_key_block,
             &decomposition_plaintexts_buffer,
             noise_distribution,
-            &mut generator,
+            generator,
         );
     }
 }
@@ -527,6 +560,46 @@ where
     new_lwe_keyswitch_key
 }
 
+pub fn allocate_and_generate_lwe_key_switching_key_with_pre_seeded_generator<
+    InputLweCont,
+    OutputLweCont,
+    Gen,
+>(
+    input_lwe_secret_key: &LweSecretKey<InputLweCont>,
+    output_lwe_secret_key: &LweSecretKey<OutputLweCont>,
+    decomp_base_log: DecompositionBaseLog,
+    decomp_level_count: DecompositionLevelCount,
+    noise_distribution: DynamicDistribution<InputLweCont::Element>,
+    ciphertext_modulus: CiphertextModulus<InputLweCont::Element>,
+    generator: &mut EncryptionRandomGenerator<Gen>,
+) -> SeededLweKeyswitchKeyOwned<InputLweCont::Element>
+where
+    InputLweCont: Container,
+    InputLweCont::Element:
+        UnsignedInteger + Encryptable<Uniform, DynamicDistribution<InputLweCont::Element>>,
+    OutputLweCont: Container<Element = InputLweCont::Element>,
+    Gen: ByteRandomGenerator,
+{
+    let mut key_switching_key = SeededLweKeyswitchKeyOwned::new(
+        InputLweCont::Element::ZERO,
+        decomp_base_log,
+        decomp_level_count,
+        input_lwe_secret_key.lwe_dimension(),
+        output_lwe_secret_key.lwe_dimension(),
+        CompressionSeed::from(Seed(0)),
+        ciphertext_modulus,
+    );
+    generate_seeded_lwe_keyswitch_key_with_pre_seeded_generator(
+        input_lwe_secret_key,
+        output_lwe_secret_key,
+        &mut key_switching_key,
+        noise_distribution,
+        generator,
+    );
+
+    key_switching_key
+}
+
 /// A generator for producing chunks of an LWE keyswitch key.
 ///
 /// This struct allows for the generation of LWE keyswitch key chunks, which can be used to

tfhe/src/core_crypto/algorithms/lwe_multi_bit_bootstrap_key_generation.rs

@@ -496,6 +496,47 @@ pub fn generate_seeded_lwe_multi_bit_bootstrap_key<
     OutputCont: ContainerMut<Element = Scalar>,
     // Maybe Sized allows to pass Box<dyn Seeder>.
     NoiseSeeder: Seeder + ?Sized,
+{
+    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
+        output.compression_seed().seed,
+        noise_seeder,
+    );
+
+    generate_seeded_lwe_multi_bit_bootstrap_key_with_pre_seeded_generator(
+        input_lwe_secret_key,
+        output_glwe_secret_key,
+        output,
+        noise_distribution,
+        &mut generator,
+    );
+}
+
+/// Fill a [`seeded LWE bootstrap key`](`SeededLweMultiBitBootstrapKey`) with an actual seeded
+/// bootstrapping key constructed from an input key [`LWE secret key`](`LweSecretKey`) and an output
+/// key [`GLWE secret key`](`GlweSecretKey`)
+///
+///
+/// This uses an already seeded generator
+pub fn generate_seeded_lwe_multi_bit_bootstrap_key_with_pre_seeded_generator<
+    Scalar,
+    NoiseDistribution,
+    InputKeyCont,
+    OutputKeyCont,
+    OutputCont,
+    ByteGen,
+>(
+    input_lwe_secret_key: &LweSecretKey<InputKeyCont>,
+    output_glwe_secret_key: &GlweSecretKey<OutputKeyCont>,
+    output: &mut SeededLweMultiBitBootstrapKey<OutputCont>,
+    noise_distribution: NoiseDistribution,
+    generator: &mut EncryptionRandomGenerator<ByteGen>,
+) where
+    Scalar: Encryptable<Uniform, NoiseDistribution> + CastFrom<usize>,
+    NoiseDistribution: Distribution,
+    InputKeyCont: Container<Element = Scalar>,
+    OutputKeyCont: Container<Element = Scalar>,
+    OutputCont: ContainerMut<Element = Scalar>,
+    ByteGen: ByteRandomGenerator,
 {
     assert!(
         output.input_lwe_dimension() == input_lwe_secret_key.lwe_dimension(),
@@ -521,11 +562,6 @@ pub fn generate_seeded_lwe_multi_bit_bootstrap_key<
         output.polynomial_size()
     );
 
-    let mut generator = EncryptionRandomGenerator::<DefaultRandomGenerator>::new(
-        output.compression_seed().seed,
-        noise_seeder,
-    );
-
     assert!(
         output.input_lwe_dimension() == input_lwe_secret_key.lwe_dimension(),
         "Mismatched LweDimension between input LWE secret key and LWE bootstrap key. \