chore(ci): add permission to github token to release crates

When using crates.io trusted publishing feature GitHub token `id-token: write` permission to be able to authenticate the workflow on the registry.

.gitattributes

@@ -1,3 +0,0 @@
-*.hpu filter=lfs diff=lfs merge=lfs -text
-*.bcode filter=lfs diff=lfs merge=lfs -text
-*.cbor filter=lfs diff=lfs merge=lfs -text

.github/actions/gpu_setup/action.yml

@@ -23,58 +23,38 @@ runs:
         echo "${CMAKE_SCRIPT_SHA} cmake-${CMAKE_VERSION}-linux-x86_64.sh" > checksum
         sha256sum -c checksum
         sudo bash cmake-"${CMAKE_VERSION}"-linux-x86_64.sh --skip-license --prefix=/usr/ --exclude-subdir
-        sudo apt remove -y unattended-upgrades
         sudo apt update
         sudo apt install -y cmake-format libclang-dev
       env:
         CMAKE_VERSION: 3.29.6
         CMAKE_SCRIPT_SHA: "6e4fada5cba3472ae503a11232b6580786802f0879cead2741672bf65d97488a"
 
-    - name: Install GCC
-      if: inputs.github-instance == 'true'
-      shell: bash
-      env:
-        GCC_VERSION: ${{ inputs.gcc-version }}
-      run: |
-        sudo apt-get install gcc-"{GCC_VERSION}" g++-"{GCC_VERSION}"
-        sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-"{GCC_VERSION}" 20
-        sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-"{GCC_VERSION}" 20
-
-    - name: Check GCC
-      shell: bash
-      env:
-        GCC_VERSION: ${{ inputs.gcc-version }}
-      run: |
-        which gcc-"${GCC_VERSION}"
-
     - name: Install CUDA
       if: inputs.github-instance == 'true'
       shell: bash
-      env:
-        CUDA_VERSION: ${{ inputs.cuda-version }}
-        CUDA_KEYRING_PACKAGE: cuda-keyring_1.1-1_all.deb
-        CUDA_KEYRING_SHA: "d93190d50b98ad4699ff40f4f7af50f16a76dac3bb8da1eaaf366d47898ff8df"
       run: |
         # Use Sed to extract a value from a string, this cannot be done with the ${variable//search/replace} pattern.
         # shellcheck disable=SC2001
         TOOLKIT_VERSION="$(echo "${CUDA_VERSION}" | sed 's/\(.*\)\.\(.*\)/\1-\2/')"
-        wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/${CUDA_KEYRING_PACKAGE}
+        wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/${env.CUDA_KEYRING_PACKAGE}
         echo "${CUDA_KEYRING_SHA} ${CUDA_KEYRING_PACKAGE}" > checksum
         sha256sum -c checksum
         sudo dpkg -i "${CUDA_KEYRING_PACKAGE}"
         sudo apt update
         sudo apt -y install cuda-toolkit-"${TOOLKIT_VERSION}"
+      env:
+        CUDA_VERSION: ${{ inputs.cuda-version }}
+        CUDA_KEYRING_PACKAGE: cuda-keyring_1.1-1_all.deb
+        CUDA_KEYRING_SHA: "d93190d50b98ad4699ff40f4f7af50f16a76dac3bb8da1eaaf366d47898ff8df"
 
     - name: Export CUDA variables
       shell: bash
       run: |
-        find /usr/local -executable -name "nvcc"
         CUDA_PATH=/usr/local/cuda-"${CUDA_VERSION}"
         {
           echo "CUDA_PATH=$CUDA_PATH";
           echo "LD_LIBRARY_PATH=$CUDA_PATH/lib64:$LD_LIBRARY_PATH";
           echo "CUDA_MODULE_LOADER=EAGER";
-          echo "PATH=$PATH:$CUDA_PATH/bin"; 
         } >> "${GITHUB_ENV}"
         {
           echo "PATH=$PATH:$CUDA_PATH/bin"; 
@@ -94,11 +74,6 @@ runs:
       env:
         GCC_VERSION: ${{ inputs.gcc-version }}
 
-    - name: Check setup
-      shell: bash
-      run: |
-        which nvcc
-
     - name: Check device is detected
       shell: bash
       run: nvidia-smi

.github/workflows/aws_tfhe_backward_compat_tests.yml

@@ -67,30 +67,49 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # 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
+      - name: Use specific data branch
+        if: ${{ contains(github.event.pull_request.labels.*.name, 'data_PR') }}
+        env:
+          PR_BRANCH: ${{ github.head_ref || github.ref_name }}
         run: |
-          SHA=$(git lfs ls-files -l -I utils/tfhe-backward-compat-data | sha256sum | cut -d' ' -f1)
+          echo "BACKWARD_COMPAT_DATA_BRANCH=${PR_BRANCH}" >> "${GITHUB_ENV}"
+
+      - name: Get backward compat branch
+        id: backward_compat_branch
+        run: |
+          BRANCH="$(make backward_compat_branch)"
+          echo "branch=${BRANCH}" >> "${GITHUB_OUTPUT}"
+
+      - name: Get backward compat branch head SHA
+        id: backward_compat_sha
+        run: |
+          SHA=$(git ls-remote "${REPO_URL}" refs/heads/"${BACKWARD_COMPAT_BRANCH}" | awk '{print $1}')
           echo "sha=${SHA}" >> "${GITHUB_OUTPUT}"
+        env:
+          REPO_URL: "https://github.com/zama-ai/tfhe-backward-compat-data"
+          BACKWARD_COMPAT_BRANCH: ${{ steps.backward_compat_branch.outputs.branch }}
 
       - name: Retrieve data from cache
         id: retrieve-data-cache
         uses: actions/cache/restore@5a3ec84eff668545956fd18022155c47e93e2684 #v4.2.3
         with:
-          path: |
-            utils/tfhe-backward-compat-data/**/*.cbor
-            utils/tfhe-backward-compat-data/**/*.bcode
-          key: ${{ steps.hash-lfs-data.outputs.sha }}
+          path: tests/tfhe-backward-compat-data
+          key: ${{ steps.backward_compat_branch.outputs.branch }}_${{ steps.backward_compat_sha.outputs.sha }}
 
-      - name: Pull test data
+      - name: Clone test data
         if: steps.retrieve-data-cache.outputs.cache-hit != 'true'
-        run: |
-          make pull_backward_compat_data
+        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
+        with:
+          persist-credentials: 'false'
+          token: ${{ env.CHECKOUT_TOKEN }}
+          repository: zama-ai/tfhe-backward-compat-data
+          path: tests/tfhe-backward-compat-data
+          lfs: 'true'
+          ref: ${{ steps.backward_compat_branch.outputs.branch }}
 
       - name: Run backward compatibility tests
         run: |
@@ -101,10 +120,8 @@ jobs:
         continue-on-error: true
         uses: actions/cache/save@5a3ec84eff668545956fd18022155c47e93e2684 #v4.2.3
         with:
-          path: |
-            utils/tfhe-backward-compat-data/**/*.cbor
-            utils/tfhe-backward-compat-data/**/*.bcode
-          key: ${{ steps.hash-lfs-data.outputs.sha }}
+          path: tests/tfhe-backward-compat-data
+          key: ${{ steps.backward_compat_branch.outputs.branch }}_${{ steps.backward_compat_sha.outputs.sha }}
 
       - name: Set pull-request URL
         if: ${{ failure() && github.event_name == 'pull_request' }}

.github/workflows/aws_tfhe_fast_tests.yml

@@ -174,7 +174,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -182,11 +182,9 @@ jobs:
         if: needs.should-run.outputs.csprng_test == 'true'
         run: |
           make test_tfhe_csprng
-          make test_tfhe_csprng_big_endian
 
       - name: Run tfhe-zk-pok tests
-        # Always run it to catch non deterministic bugs earlier
-        # if: needs.should-run.outputs.zk_pok_test == 'true'
+        if: needs.should-run.outputs.zk_pok_test == 'true'
         run: |
           make test_zk_pok
 

.github/workflows/aws_tfhe_integer_tests.yml

@@ -103,7 +103,7 @@ jobs:
     name: Unsigned integer tests
     needs: setup-instance
     concurrency:
-      group: ${{ github.workflow_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+      group: ${{ github.workflow_ref }}
       cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
     runs-on: ${{ needs.setup-instance.outputs.runner-name }}
     steps:
@@ -114,7 +114,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/aws_tfhe_noise_checks.yml

@@ -1,115 +0,0 @@
-name: Run noise checks on CPU
-
-env:
-  CARGO_TERM_COLOR: always
-  ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
-  RUSTFLAGS: "-C target-cpu=native"
-  RUST_BACKTRACE: "full"
-  RUST_MIN_STACK: "8388608"
-  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 }}
-  SLACKIFY_MARKDOWN: true
-  PULL_REQUEST_MD_LINK: ""
-  CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
-  # Secrets will be available only to zama-ai organization members
-  SECRETS_AVAILABLE: ${{ secrets.JOB_SECRET != '' }}
-
-on:
-  # Allows you to run this workflow manually from the Actions tab as an alternative.
-  workflow_dispatch:
-
-permissions:
-  contents: read
-
-jobs:
-  setup-instance:
-    name: Setup instance (noise-checks)
-    runs-on: ubuntu-latest
-    outputs:
-      runner-name: ${{ steps.start-remote-instance.outputs.label || steps.start-github-instance.outputs.runner_group }}
-    steps:
-      - name: Start remote instance
-        id: start-remote-instance
-        if: env.SECRETS_AVAILABLE == 'true'
-        uses: zama-ai/slab-github-runner@79939325c3c429837c10d6041e4fd8589d328bac
-        with:
-          mode: start
-          github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
-          slab-url: ${{ secrets.SLAB_BASE_URL }}
-          job-secret: ${{ secrets.JOB_SECRET }}
-          backend: aws
-          # We want an hpc7a more compute, will be faster
-          profile: bench
-
-      # This instance will be spawned especially for pull-request from forked repository
-      - name: Start GitHub instance
-        id: start-github-instance
-        if: env.SECRETS_AVAILABLE == 'false'
-        run: |
-          echo "Cannot run this without secrets"
-          exit 1
-
-  noise-checks:
-    name: CPU noise checks
-    needs: setup-instance
-    runs-on: ${{ needs.setup-instance.outputs.runner-name }}
-    timeout-minutes: 1440
-    steps:
-      - name: Checkout tfhe-rs
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
-        with:
-          persist-credentials: 'false'
-          token: ${{ env.CHECKOUT_TOKEN }}
-
-      - name: Install latest stable
-        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
-        with:
-          toolchain: stable
-
-      - name: Run noise checks
-        timeout-minutes: 1440
-        run: |
-          make test_noise_check
-
-      - name: Set pull-request URL
-        if: ${{ !success() }}
-        run: |
-          echo "PULL_REQUEST_MD_LINK=[pull-request](${PR_BASE_URL}${PR_NUMBER}), "  >> "${GITHUB_ENV}"
-        env:
-          PR_BASE_URL: ${{ vars.PR_BASE_URL }}
-          PR_NUMBER: ${{ github.event.pull_request.number }}
-
-      - name: Slack Notification
-        if: ${{ !success() }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "Noise checks tests finished with status: ${{ job.status }}. (${{ env.PULL_REQUEST_MD_LINK }}[action run](${{ env.ACTION_RUN_URL }}))"
-
-  teardown-instance:
-    name: Teardown instance (noise-checks)
-    if: ${{ always() && needs.setup-instance.result == 'success' }}
-    needs: [ setup-instance, noise-checks ]
-    runs-on: ubuntu-latest
-    steps:
-      - name: Stop remote instance
-        id: stop-instance
-        if: env.SECRETS_AVAILABLE == 'true'
-        uses: zama-ai/slab-github-runner@79939325c3c429837c10d6041e4fd8589d328bac
-        with:
-          mode: stop
-          github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
-          slab-url: ${{ secrets.SLAB_BASE_URL }}
-          job-secret: ${{ secrets.JOB_SECRET }}
-          label: ${{ needs.setup-instance.outputs.runner-name }}
-
-      - name: Slack Notification
-        if: ${{ !success() }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "Instance teardown (noise-checks) finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"

.github/workflows/aws_tfhe_signed_integer_tests.yml

@@ -104,7 +104,7 @@ jobs:
     name: Signed integer tests
     needs: setup-instance
     concurrency:
-      group: ${{ github.workflow_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+      group: ${{ github.workflow_ref }}
       cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
     runs-on: ${{ needs.setup-instance.outputs.runner-name }}
     steps:
@@ -115,7 +115,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/aws_tfhe_tests.yml

@@ -185,7 +185,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/aws_tfhe_wasm_tests.yml

@@ -68,7 +68,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/benchmark_boolean.yml

@@ -68,7 +68,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_core_crypto.yml

@@ -68,7 +68,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_dex.yml

@@ -68,7 +68,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -98,27 +98,15 @@ jobs:
         env:
           REF_NAME: ${{ github.ref_name }}
 
-      - name: Parse swap request update PBS counts
+      - name: Parse swap request PBS counts
         run: |
-          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_request_update_dex_balance_pbs_count.csv "${RESULTS_FILENAME}" \
+          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_request_pbs_count.csv "${RESULTS_FILENAME}" \
           --object-sizes \
           --append-results
 
-      - name: Parse swap request finalize PBS counts
+      - name: Parse swap claim PBS counts
         run: |
-          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_request_finalize_pbs_count.csv "${RESULTS_FILENAME}" \
-          --object-sizes \
-          --append-results
-
-      - name: Parse swap claim prepare PBS counts
-        run: |
-          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_claim_prepare_pbs_count.csv "${RESULTS_FILENAME}" \
-          --object-sizes \
-          --append-results
-
-      - name: Parse swap claim update PBS counts
-        run: |
-          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_claim_update_dex_balance_pbs_count.csv "${RESULTS_FILENAME}" \
+          python3 ./ci/benchmark_parser.py tfhe-benchmark/dex_swap_claim_pbs_count.csv "${RESULTS_FILENAME}" \
           --object-sizes \
           --append-results
 

.github/workflows/benchmark_erc20.yml

@@ -69,7 +69,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_gpu.yml

@@ -31,7 +31,6 @@ on:
           - ks
           - ks_pbs
           - integer_zk
-          - hlapi_noise_squash
       op_flavor:
         description: "Operations set to run"
         type: choice

.github/workflows/benchmark_gpu_4090.yml

@@ -57,7 +57,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -140,7 +140,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -168,8 +168,7 @@ jobs:
           --commit-date "${COMMIT_DATE}" \
           --bench-date "${BENCH_DATE}" \
           --walk-subdirs \
-        env:
-          REF_NAME: ${{ github.ref_name }}
+      
 
       - name: Upload parsed results artifact
         uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02

.github/workflows/benchmark_gpu_common.yml

@@ -84,7 +84,7 @@ jobs:
         run: |
           # Use Sed to extract a value from a string, this cannot be done with the ${variable//search/replace} pattern.
           # shellcheck disable=SC2001
-          PARSED_COMMAND=$(echo "${INPUTS_COMMAND}" | sed 's/[[:space:]]*,[[:space:]]*/\", \"/g')
+          PARSED_COMMAND=$(echo "${INPUTS_COMMAND}" | sed 's/[[:space:]]*,[[:space:]]*/\\", \\"/g')
           echo "COMMAND=[\"${PARSED_COMMAND}\"]" >> "${GITHUB_ENV}"
 
       - name: Set single operations flavor
@@ -122,23 +122,31 @@ jobs:
 
       - name: Set command output
         id: set_command
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "command=${{ toJSON(env.COMMAND) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "command=${COMMAND_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          COMMAND_OUTPUT: ${{ toJSON(env.COMMAND) }}
 
       - name: Set operation flavor output
         id: set_op_flavor
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "op_flavor=${{ toJSON(env.OP_FLAVOR) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "op_flavor=${OP_FLAVOR_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          OP_FLAVOR_OUTPUT: ${{ toJSON(env.OP_FLAVOR) }}
 
       - name: Set benchmark types output
         id: set_bench_type
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "bench_type=${{ toJSON(env.BENCH_TYPE) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "bench_type=${BENCH_TYPE_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          BENCH_TYPE_OUTPUT: ${{ toJSON(env.BENCH_TYPE) }}
 
       - name: Set parameters types output
         id: set_params_type
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "params_type=${{ toJSON(env.PARAMS_TYPE) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "params_type=${PARAMS_TYPE_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          PARAMS_TYPE_OUTPUT: ${{ toJSON(env.PARAMS_TYPE) }}
 
   setup-instance:
     name: Setup instance (cuda-${{ inputs.profile }}-benchmarks)
@@ -192,7 +200,7 @@ jobs:
       matrix:
         # explicit include-based build matrix, of known valid options
         include:
-          - cuda: "12.8"
+          - cuda: "12.2"
             gcc: 11
     steps:
       - name: Checkout tfhe-rs repo
@@ -224,7 +232,7 @@ jobs:
         params_type: ${{ fromJSON(needs.prepare-matrix.outputs.params_type) }}
         # explicit include-based build matrix, of known valid options
         include:
-          - cuda: "12.8"
+          - cuda: "12.2"
             gcc: 11
     env:
       CUDA_PATH: /usr/local/cuda-${{ matrix.cuda }}
@@ -269,7 +277,7 @@ jobs:
           GCC_VERSION: ${{ matrix.gcc }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_gpu_dex_common.yml

@@ -100,7 +100,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11
     steps:
       - name: Checkout tfhe-rs repo with tags
@@ -129,7 +129,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_gpu_erc20_common.yml

@@ -101,7 +101,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11
     steps:
       - name: Checkout tfhe-rs repo with tags
@@ -130,7 +130,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_gpu_weekly.yml

@@ -10,16 +10,37 @@ on:
 permissions: {}
 
 jobs:
-  run-benchmarks-8-h100-sxm5-integer:
-    name: Run integer benchmarks (8xH100-SXM5)
+  run-benchmarks-1-h100:
+    name: Run integer benchmarks (1xH100)
     if: github.repository == 'zama-ai/tfhe-rs'
     uses: ./.github/workflows/benchmark_gpu_common.yml
     with:
-      profile: multi-h100-sxm5
-      hardware_name: n3-H100x8-SXM5
+      profile: single-h100
+      hardware_name: n3-H100x1
+      command: integer,integer_multi_bit
+      op_flavor: default
+      bench_type: latency
+      all_precisions: true
+    secrets:
+      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
+      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
+      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
+      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
+      JOB_SECRET: ${{ secrets.JOB_SECRET }}
+      SLAB_ACTION_TOKEN: ${{ secrets.SLAB_ACTION_TOKEN }}
+      SLAB_URL: ${{ secrets.SLAB_URL }}
+      SLAB_BASE_URL: ${{ secrets.SLAB_BASE_URL }}
+
+  run-benchmarks-2-h100:
+    name: Run integer benchmarks (2xH100)
+    if: github.repository == 'zama-ai/tfhe-rs'
+    uses: ./.github/workflows/benchmark_gpu_common.yml
+    with:
+      profile: 2-h100
+      hardware_name: n3-H100x2
       command: integer_multi_bit
       op_flavor: default
-      bench_type: both
+      bench_type: latency
       all_precisions: true
     secrets:
       BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
@@ -31,16 +52,16 @@ jobs:
       SLAB_URL: ${{ secrets.SLAB_URL }}
       SLAB_BASE_URL: ${{ secrets.SLAB_BASE_URL }}
 
-  run-benchmarks-8-h100-sxm5-integer-compression:
-    name: Run integer compression benchmarks (8xH100-SXM5)
+  run-benchmarks-8-h100:
+    name: Run integer benchmarks (8xH100)
     if: github.repository == 'zama-ai/tfhe-rs'
     uses: ./.github/workflows/benchmark_gpu_common.yml
     with:
-      profile: multi-h100-sxm5
-      hardware_name: n3-H100x8-SXM5
-      command: integer_compression
+      profile: multi-h100
+      hardware_name: n3-H100x8
+      command: integer_multi_bit
       op_flavor: default
-      bench_type: both
+      bench_type: latency
       all_precisions: true
     secrets:
       BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
@@ -52,37 +73,16 @@ jobs:
       SLAB_URL: ${{ secrets.SLAB_URL }}
       SLAB_BASE_URL: ${{ secrets.SLAB_BASE_URL }}
 
-  run-benchmarks-8-h100-sxm5-integer-zk:
-    name: Run integer zk benchmarks (8xH100-SXM5)
+  run-benchmarks-l40:
+    name: Run integer benchmarks (L40)
     if: github.repository == 'zama-ai/tfhe-rs'
     uses: ./.github/workflows/benchmark_gpu_common.yml
     with:
-      profile: multi-h100-sxm5
-      hardware_name: n3-H100x8-SXM5
-      command: integer_zk
+      profile: l40
+      hardware_name: n3-L40x1
+      command: integer_multi_bit,integer_compression,pbs,ks
       op_flavor: default
-      bench_type: both
-      all_precisions: true
-    secrets:
-      BOT_USERNAME: ${{ secrets.BOT_USERNAME }}
-      SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }}
-      SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
-      REPO_CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN }}
-      JOB_SECRET: ${{ secrets.JOB_SECRET }}
-      SLAB_ACTION_TOKEN: ${{ secrets.SLAB_ACTION_TOKEN }}
-      SLAB_URL: ${{ secrets.SLAB_URL }}
-      SLAB_BASE_URL: ${{ secrets.SLAB_BASE_URL }}
-
-  run-benchmarks-8-h100-sxm5-noise-squash:
-    name: Run integer zk benchmarks (8xH100-SXM5)
-    if: github.repository == 'zama-ai/tfhe-rs'
-    uses: ./.github/workflows/benchmark_gpu_common.yml
-    with:
-      profile: multi-h100-sxm5
-      hardware_name: n3-H100x8-SXM5
-      command: hlapi_noise_squash
-      op_flavor: default
-      bench_type: both
+      bench_type: latency
       all_precisions: true
     secrets:
       BOT_USERNAME: ${{ secrets.BOT_USERNAME }}

.github/workflows/benchmark_hpu_integer.yml

@@ -33,7 +33,6 @@ jobs:
         with:
           fetch-depth: 0
           persist-credentials: 'false'
-          lfs: true
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
 
       - name: Get benchmark details
@@ -48,7 +47,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -62,7 +61,6 @@ jobs:
 
       - name: Run benchmarks
         run: |
-          make pull_hpu_files
           make bench_integer_hpu
           make bench_hlapi_erc20_hpu
 
@@ -81,7 +79,7 @@ jobs:
           REF_NAME: ${{ github.ref_name }}
 
       - name: Upload parsed results artifact
-        uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02
+        uses: actions/upload-artifact@65c4c4a1ddee5b72f698fdd19549f0f0fb45cf08
         with:
           name: ${{ github.sha }}_integer_benchmarks
           path: ${{ env.RESULTS_FILENAME }}

.github/workflows/benchmark_integer.yml

@@ -78,13 +78,17 @@ jobs:
 
       - name: Set operation flavor output
         id: set_op_flavor
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "op_flavor=${{ toJSON(env.OP_FLAVOR) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "op_flavor=${OP_FLAVOR_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          OP_FLAVOR_OUTPUT: ${{ toJSON(env.OP_FLAVOR) }}
 
       - name: Set benchmark types output
         id: set_bench_type
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "bench_type=${{ toJSON(env.BENCH_TYPE) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "bench_type=${BENCH_TYPE_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          BENCH_TYPE_OUTPUT: ${{ toJSON(env.BENCH_TYPE) }}
 
   setup-instance:
     name: Setup instance (integer-benchmarks)
@@ -138,7 +142,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_shortint.yml

@@ -47,8 +47,10 @@ jobs:
 
       - name: Set operation flavor output
         id: set_op_flavor
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "op_flavor=${{ toJSON(env.OP_FLAVOR) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "op_flavor=${OP_FLAVOR_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          OP_FLAVOR_OUTPUT: ${{ toJSON(env.OP_FLAVOR) }}
 
   setup-instance:
     name: Setup instance (shortint-benchmarks)
@@ -99,7 +101,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_signed_integer.yml

@@ -78,13 +78,17 @@ jobs:
 
       - name: Set operation flavor output
         id: set_op_flavor
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "op_flavor=${{ toJSON(env.OP_FLAVOR) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "op_flavor=${OP_FLAVOR_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          OP_FLAVOR_OUTPUT: ${{ toJSON(env.OP_FLAVOR) }}
 
       - name: Set benchmark types output
         id: set_bench_type
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "bench_type=${{ toJSON(env.BENCH_TYPE) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "bench_type=${BENCH_TYPE_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          BENCH_TYPE_OUTPUT: ${{ toJSON(env.BENCH_TYPE) }}
 
   setup-instance:
     name: Setup instance (signed-integer-benchmarks)
@@ -138,7 +142,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_tfhe_fft.yml

@@ -48,7 +48,7 @@ jobs:
     name: Execute FFT benchmarks in EC2
     needs: setup-ec2
     concurrency:
-      group: ${{ github.workflow_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+      group: ${{ github.workflow_ref }}
       cancel-in-progress: true
     runs-on: ${{ needs.setup-ec2.outputs.runner-name }}
     steps:

.github/workflows/benchmark_tfhe_ntt.yml

@@ -48,7 +48,7 @@ jobs:
     name: Execute NTT benchmarks in EC2
     needs: setup-ec2
     concurrency:
-      group: ${{ github.workflow_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+      group: ${{ github.workflow_ref }}
       cancel-in-progress: true
     runs-on: ${{ needs.setup-ec2.outputs.runner-name }}
     steps:

.github/workflows/benchmark_tfhe_zk_pok.yml

@@ -108,7 +108,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -144,7 +144,7 @@ jobs:
       - name: Upload parsed results artifact
         uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02
         with:
-          name: ${{ github.sha }}_tfhe_zk_pok_${{ env.BENCH_TYPE }}
+          name: ${{ github.sha }}_tfhe_zk_pok
           path: ${{ env.RESULTS_FILENAME }}
 
       - name: Checkout Slab repo

.github/workflows/benchmark_wasm_client.yml

@@ -106,7 +106,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 

.github/workflows/benchmark_zk_pke.yml

@@ -92,8 +92,10 @@ jobs:
 
       - name: Set benchmark types output
         id: set_bench_type
-        run: | # zizmor: ignore[template-injection] this env variable is safe
-          echo "bench_type=${{ toJSON(env.BENCH_TYPE) }}" >> "${GITHUB_OUTPUT}"
+        run: |
+          echo "bench_type=${BENCH_TYPE_OUTPUT}" >> "${GITHUB_OUTPUT}"
+        env:
+          BENCH_TYPE_OUTPUT: ${{ toJSON(env.BENCH_TYPE) }}
 
   setup-instance:
     name: Setup instance (pke-zk-benchmarks)
@@ -150,7 +152,7 @@ jobs:
           SHA: ${{ github.sha }}
 
       - name: Install rust
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: nightly
 
@@ -194,7 +196,7 @@ jobs:
       - name: Upload parsed results artifact
         uses: actions/upload-artifact@ea165f8d65b6e75b540449e92b4886f43607fa02
         with:
-          name: ${{ github.sha }}_integer_zk_${{ matrix.bench_type }}
+          name: ${{ github.sha }}_integer_zk
           path: ${{ env.RESULTS_FILENAME }}
 
       - name: Checkout Slab repo

.github/workflows/cargo_build.yml

@@ -35,7 +35,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -49,14 +49,6 @@ jobs:
           mv linelint-linux-amd64 /usr/local/bin/linelint
           make check_newline
 
-      # This is needed for the ws tests clippy checks
-      - name: Use specific data branch
-        if: ${{ contains(github.event.pull_request.labels.*.name, 'data_PR') }}
-        env:
-          PR_BRANCH: ${{ github.head_ref || github.ref_name }}
-        run: |
-          echo "BACKWARD_COMPAT_DATA_BRANCH=${PR_BRANCH}" >> "${GITHUB_ENV}"
-
       - name: Run pcc checks
         if: ${{ contains(matrix.os, 'ubuntu') }}
         run: |

.github/workflows/cargo_test_fft.yml

@@ -13,7 +13,7 @@ env:
   CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
 
 concurrency:
-  group: ${{ github.workflow }}-${{ github.head_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+  group: ${{ github.workflow }}-${{ github.head_ref }}
   cancel-in-progress: true
 
 permissions:

.github/workflows/cargo_test_ntt.yml

@@ -13,7 +13,7 @@ env:
   CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
 
 concurrency:
-  group: ${{ github.workflow }}-${{ github.head_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+  group: ${{ github.workflow }}-${{ github.head_ref }}
   cancel-in-progress: true
 
 permissions:

.github/workflows/check_commit.yml

@@ -3,15 +3,14 @@ name: Check commit and PR compliance
 on:
   pull_request:
 
-permissions: {}
+permissions:
+  contents: read
+  pull-requests: read # Permission needed to scan commits in a pull-request
 
 jobs:
   check-commit-pr:
     name: Check commit and PR
     runs-on: ubuntu-latest
-    permissions:
-      contents: read
-      pull-requests: write # Permission needed to scan commits in a pull-request and write issue comment
     steps:
       - name: Check first line
         uses: gsactions/commit-message-checker@16fa2d5de096ae0d35626443bcd24f1e756cafee

.github/workflows/ci_lint.yml

@@ -38,11 +38,9 @@ jobs:
       - name: Check workflows security
         run: |
           make check_workflow_security
-        env:
-          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@4830be28ce81da52ec70d65c552a7403821d98d4 # v3.0.23
         with:
           allowlist: |
             slsa-framework/slsa-github-generator

.github/workflows/code_coverage.yml

@@ -54,7 +54,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -90,7 +90,7 @@ jobs:
           make test_shortint_cov
 
       - name: Upload tfhe coverage to Codecov
-        uses: codecov/codecov-action@18283e04ce6e62d37312384ff67231eb8fd56d24
+        uses: codecov/codecov-action@ad3126e916f78f00edff4ed0317cf185271ccc2d
         if: steps.changed-files.outputs.tfhe_any_changed == 'true'
         with:
           token: ${{ secrets.CODECOV_TOKEN }}
@@ -104,7 +104,7 @@ jobs:
           make test_integer_cov
 
       - name: Upload tfhe coverage to Codecov
-        uses: codecov/codecov-action@18283e04ce6e62d37312384ff67231eb8fd56d24
+        uses: codecov/codecov-action@ad3126e916f78f00edff4ed0317cf185271ccc2d
         if: steps.changed-files.outputs.tfhe_any_changed == 'true'
         with:
           token: ${{ secrets.CODECOV_TOKEN }}

.github/workflows/csprng_randomness_tests.yml

@@ -66,7 +66,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/data_pr_close.yml

@@ -0,0 +1,142 @@
+name: Close or Merge corresponding PR on the data repo
+
+# When a PR with the data_PR tag is closed or merged, this will close the corresponding PR in the data repo.
+
+env:
+  TARGET_REPO_API_URL: ${{ github.api_url }}/repos/zama-ai/tfhe-backward-compat-data
+  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 }}
+  PR_BRANCH: ${{ github.head_ref || github.ref_name }}
+  CLOSE_TYPE: ${{ github.event.pull_request.merged && 'merge' || 'close' }}
+
+# only trigger on pull request closed events
+on:
+  pull_request:
+    types: [ closed ]
+
+# The same pattern is used for jobs that use the github api:
+# - save the result of the API call in the env var "GH_API_RES". Since the var is multiline
+# we use this trick: https://docs.github.com/en/actions/using-workflows/workflow-commands-for-github-actions#example-of-a-multiline-string
+# - "set +e" will make sure we reach the last "echo EOF" even in case of error
+# - "set -o" pipefail makes one line piped command return the error of the first failure
+# - 'RES="$?"' and 'exit $RES' are used to return the error code if a command failed. Without it, with "set +e"
+# the script will always return 0 because of the "echo EOF".
+
+
+
+permissions: {}
+
+jobs:
+  auto_close_job:
+    if: ${{ contains(github.event.pull_request.labels.*.name, 'data_PR') }}
+    runs-on: ubuntu-latest
+    steps:
+    - name: Find corresponding Pull Request in the data repo
+      run: |
+        {
+          set +e
+          set -o pipefail
+          echo 'TARGET_REPO_PR<<EOF'
+          curl --fail-with-body --no-progress-meter -L -X GET \
+          -H "Accept: application/vnd.github+json" \
+          -H "X-GitHub-Api-Version: 2022-11-28"  \
+          "${TARGET_REPO_API_URL}"/pulls\?head="${REPO_OWNER}":"${PR_BRANCH}" | jq -e '.[0]' | sed 's/null/{ "message": "corresponding PR not found" }/'
+          RES="$?"
+          echo EOF
+        } >> "${GITHUB_ENV}"
+        exit $RES
+      env:
+        REPO_OWNER: ${{ github.repository_owner }}
+
+    - name: Comment on the PR to indicate the reason of the close
+      run: |
+        BODY="'{ \"body\": \"PR ${CLOSE_TYPE}d because the corresponding PR in main repo was ${CLOSE_TYPE}d: ${REPO}#${EVENT_NUMBER}\" }'"
+        {
+          set +e
+          set -o pipefail
+          echo 'GH_API_RES<<EOF'
+          curl --fail-with-body --no-progress-meter -L -X POST \
+          -H "Accept: application/vnd.github+json" \
+          -H "Authorization: Bearer ${TOKEN}" \
+          -H "X-GitHub-Api-Version: 2022-11-28" \
+          "${COMMENTS_URL}" \
+          -d "${BODY}"
+          RES="$?"
+          echo EOF
+        } >> "${GITHUB_ENV}"
+        exit $RES
+      env:
+        REPO: ${{ github.repository }}
+        EVENT_NUMBER: ${{ github.event.number }}
+        COMMENTS_URL: ${{ fromJson(env.TARGET_REPO_PR).comments_url }}
+        TOKEN: ${{ secrets.FHE_ACTIONS_TOKEN }}
+
+    - name: Merge the Pull Request in the data repo
+      if: ${{ github.event.pull_request.merged }}
+      run: |
+        {
+          set +e
+          set -o pipefail
+          echo 'GH_API_RES<<EOF'
+          curl --fail-with-body --no-progress-meter -L -X PUT \
+          -H "Accept: application/vnd.github+json" \
+          -H "Authorization: Bearer ${TOKEN}" \
+          -H "X-GitHub-Api-Version: 2022-11-28" \
+          "${TARGET_REPO_PR_URL}"/merge \
+          -d '{ "merge_method": "rebase" }'
+          RES="$?"
+          echo EOF
+        } >> "${GITHUB_ENV}"
+        exit $RES
+      env:
+        TARGET_REPO_PR_URL: ${{ fromJson(env.TARGET_REPO_PR).url }}
+        TOKEN: ${{ secrets.FHE_ACTIONS_TOKEN }}
+
+    - name: Close the Pull Request in the data repo
+      if: ${{ !github.event.pull_request.merged }}
+      run: |
+        {
+          set +e
+          set -o pipefail
+          echo 'GH_API_RES<<EOF'
+          curl --fail-with-body --no-progress-meter -L -X PATCH \
+          -H "Accept: application/vnd.github+json" \
+          -H "Authorization: Bearer ${TOKEN}" \
+          -H "X-GitHub-Api-Version: 2022-11-28" \
+          "${TARGET_REPO_PR_URL}" \
+          -d '{ "state": "closed" }'
+          RES="$?"
+          echo EOF
+        } >> "${GITHUB_ENV}"
+        exit $RES
+      env:
+        TARGET_REPO_PR_URL: ${{ fromJson(env.TARGET_REPO_PR).url }}
+        TOKEN: ${{ secrets.FHE_ACTIONS_TOKEN }}
+
+    - name: Delete the associated branch in the data repo
+      run: |
+        {
+          set +e
+          set -o pipefail
+          echo 'GH_API_RES<<EOF'
+          curl --fail-with-body --no-progress-meter -L -X DELETE \
+          -H "Accept: application/vnd.github+json" \
+          -H "Authorization: Bearer ${TOKEN}" \
+          -H "X-GitHub-Api-Version: 2022-11-28" \
+          "${TARGET_REPO_API_URL}"/git/refs/heads/"${PR_BRANCH}"
+          RES="$?"
+          echo EOF
+        } >> "${GITHUB_ENV}"
+        exit $RES
+      env:
+        TOKEN: ${{ secrets.FHE_ACTIONS_TOKEN }}
+
+    - name: Slack Notification
+      if: ${{ always() && job.status == 'failure' }}
+      continue-on-error: true
+      uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661
+      env:
+        SLACK_COLOR: ${{ job.status }}
+        SLACK_MESSAGE: "Failed to auto-${{ env.CLOSE_TYPE }} PR on data repo: ${{ fromJson(env.GH_API_RES || env.TARGET_REPO_PR).message }}"

.github/workflows/gpu_4090_tests.yml

@@ -45,7 +45,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/gpu_code_validation_tests.yml

@@ -1,148 +0,0 @@
-# Compile and test tfhe-cuda-backend on an AWS instance
-name: Cuda - Code Validation
-
-env:
-  CARGO_TERM_COLOR: always
-  ACTION_RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}
-  RUSTFLAGS: "-C target-cpu=native"
-  RUST_BACKTRACE: "full"
-  RUST_MIN_STACK: "8388608"
-  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 }}
-  SLACKIFY_MARKDOWN: true
-  IS_PULL_REQUEST: ${{ github.event_name == 'pull_request' }}
-  PULL_REQUEST_MD_LINK: ""
-  CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
-  # Secrets will be available only to zama-ai organization members
-  SECRETS_AVAILABLE: ${{ secrets.JOB_SECRET != '' }}
-  EXTERNAL_CONTRIBUTION_RUNNER: "gpu_ubuntu-22.04"
-
-on:
-  # Allows you to run this workflow manually from the Actions tab as an alternative.
-  workflow_dispatch:
-  pull_request:
-    types: [ labeled ]
-
-permissions:
-  contents: read
-
-jobs:
-  setup-instance:
-    name: Setup instance (cuda-tests)
-    runs-on: ubuntu-latest
-    if: github.event_name != 'pull_request' ||
-      (github.event.action == 'labeled' && github.event.label.name == 'approved')
-    outputs:
-      runner-name: ${{ steps.start-remote-instance.outputs.label || steps.start-github-instance.outputs.runner_group }}
-    steps:
-      - name: Start remote instance
-        id: start-remote-instance
-        if: env.SECRETS_AVAILABLE == 'true'
-        uses: zama-ai/slab-github-runner@79939325c3c429837c10d6041e4fd8589d328bac
-        with:
-          mode: start
-          github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
-          slab-url: ${{ secrets.SLAB_BASE_URL }}
-          job-secret: ${{ secrets.JOB_SECRET }}
-          backend: hyperstack
-          profile: gpu-test
-
-      # This instance will be spawned especially for pull-request from forked repository
-      - name: Start GitHub instance
-        id: start-github-instance
-        if: env.SECRETS_AVAILABLE == 'false'
-        run: |
-          echo "runner_group=${EXTERNAL_CONTRIBUTION_RUNNER}" >> "$GITHUB_OUTPUT"
-
-  cuda-tests-linux:
-    name: CUDA Code Validation tests
-    needs: [ setup-instance ]
-    if: github.event_name != 'pull_request' ||
-      (github.event_name == 'pull_request' && needs.setup-instance.result != 'skipped')
-    concurrency:
-      group: ${{ github.workflow_ref }}
-      cancel-in-progress: ${{ github.ref != 'refs/heads/main' }}
-    runs-on: ${{ needs.setup-instance.outputs.runner-name }}
-    strategy:
-      fail-fast: false
-      # explicit include-based build matrix, of known valid options
-      matrix:
-        include:
-          - os: ubuntu-22.04
-            cuda: "12.8"
-            gcc: 11 
-    steps:
-      - name: Checkout tfhe-rs
-        uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683
-        with:
-          persist-credentials: 'false'
-          token: ${{ env.CHECKOUT_TOKEN }}
-
-      - name: Setup Hyperstack dependencies
-        uses: ./.github/actions/gpu_setup
-        with:
-          cuda-version: ${{ matrix.cuda }}
-          gcc-version: ${{ matrix.gcc }}
-          github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
-
-      - name: Find tools
-        run: |
-          find /usr -executable -name "compute-sanitizer"
-
-      - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
-        with:
-          toolchain: stable
-
-      - name: Run memory sanitizer
-        run: |
-          make test_high_level_api_gpu_debug
-
-  slack-notify:
-    name: Slack Notification
-    needs: [ setup-instance, cuda-tests-linux ]
-    runs-on: ubuntu-latest
-    if: ${{ always() && needs.cuda-tests-linux.result != 'skipped' && failure() }}
-    continue-on-error: true
-    steps:
-      - name: Set pull-request URL
-        if: env.SECRETS_AVAILABLE == 'true' && github.event_name == 'pull_request'
-        run: |
-          echo "PULL_REQUEST_MD_LINK=[pull-request](${PR_BASE_URL}${PR_NUMBER}), "  >> "${GITHUB_ENV}"
-        env:
-          PR_BASE_URL: ${{ vars.PR_BASE_URL }}
-          PR_NUMBER: ${{ github.event.pull_request.number }}
-
-      - name: Send message
-        if: env.SECRETS_AVAILABLE == 'true'
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661
-        env:
-          SLACK_COLOR: ${{ needs.cuda-tests-linux.result }}
-          SLACK_MESSAGE: "GPU code validation tests finished with status: ${{ needs.cuda-tests-linux.result }}. (${{ env.PULL_REQUEST_MD_LINK }}[action run](${{ env.ACTION_RUN_URL }}))"
-
-  teardown-instance:
-    name: Teardown instance (cuda-tests)
-    if: ${{ always() && needs.setup-instance.result == 'success' }}
-    needs: [ setup-instance, cuda-tests-linux ]
-    runs-on: ubuntu-latest
-    steps:
-      - name: Stop remote instance
-        id: stop-instance
-        if: env.SECRETS_AVAILABLE == 'true'
-        uses: zama-ai/slab-github-runner@79939325c3c429837c10d6041e4fd8589d328bac
-        with:
-          mode: stop
-          github-token: ${{ secrets.SLAB_ACTION_TOKEN }}
-          slab-url: ${{ secrets.SLAB_BASE_URL }}
-          job-secret: ${{ secrets.JOB_SECRET }}
-          label: ${{ needs.setup-instance.outputs.runner-name }}
-
-      - name: Slack Notification
-        if: ${{ failure() }}
-        continue-on-error: true
-        uses: rtCamp/action-slack-notify@e31e87e03dd19038e411e38ae27cbad084a90661
-        env:
-          SLACK_COLOR: ${{ job.status }}
-          SLACK_MESSAGE: "Instance teardown (cuda-tests) finished with status: ${{ job.status }}. (${{ env.ACTION_RUN_URL }})"

.github/workflows/gpu_fast_h100_tests.yml

@@ -122,7 +122,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -140,12 +140,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run core crypto and internal CUDA backend tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_core_crypto_gpu

.github/workflows/gpu_fast_tests.yml

@@ -107,7 +107,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -124,12 +124,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run core crypto and internal CUDA backend tests
         run: |
           make test_core_crypto_gpu

.github/workflows/gpu_full_h100_tests.yml

@@ -62,7 +62,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -79,12 +79,10 @@ jobs:
           gcc-version: ${{ matrix.gcc }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run core crypto, integer and internal CUDA backend tests
         run: |
           make test_gpu

.github/workflows/gpu_full_multi_gpu_tests.yml

@@ -109,7 +109,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -126,12 +126,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run multi-bit CUDA integer compression tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_integer_compression_gpu
@@ -139,7 +137,7 @@ jobs:
       # No need to test core_crypto and classic PBS in integer since it's already tested on single GPU.
       - name: Run multi-bit CUDA integer tests
         run: |
-          BIG_TESTS_INSTANCE=TRUE NO_BIG_PARAMS_GPU=TRUE make test_integer_multi_bit_gpu_ci
+          BIG_TESTS_INSTANCE=TRUE make test_integer_multi_bit_gpu_ci
 
       - name: Run user docs tests
         run: |
@@ -151,7 +149,7 @@ jobs:
 
       - name: Run High Level API Tests
         run: |
-          make test_high_level_api_gpu
+          BIG_TESTS_INSTANCE=FALSE make test_high_level_api_gpu
 
   slack-notify:
     name: Slack Notification

.github/workflows/gpu_integer_long_run_tests.yml

@@ -11,7 +11,6 @@ env:
   SLACK_USERNAME: ${{ secrets.BOT_USERNAME }}
   SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }}
   CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
-  IS_PR: ${{ github.event_name == 'pull_request' }}
 
 on:
   # Allows you to run this workflow manually from the Actions tab as an alternative.
@@ -19,8 +18,6 @@ on:
   schedule:
     # Nightly tests will be triggered each evening 8p.m.
     - cron: "0 20 * * *"
-  pull_request:
-
 
 permissions:
   contents: read
@@ -58,7 +55,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     timeout-minutes: 4320 # 72 hours
     steps:
@@ -75,19 +72,13 @@ jobs:
           gcc-version: ${{ matrix.gcc }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run tests
         run: |
-          if [[ "${IS_PR}" == "true" ]]; then
-            make test_integer_short_run_gpu
-          else
-            make test_integer_long_run_gpu
-          fi
+          make test_integer_long_run_gpu
 
   slack-notify:
     name: Slack Notification

.github/workflows/gpu_pcc.yml

@@ -1,4 +1,4 @@
-# Perform tfhe-cuda-backend post-commit checks on an AWS instance
+# Perfom tfhe-cuda-backend post-commit checks on an AWS instance
 name: Cuda - Post-commit Checks
 
 env:
@@ -94,7 +94,7 @@ jobs:
           CUDA_VERSION: ${{ matrix.cuda }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -128,10 +128,6 @@ jobs:
         run: |
           make pcc_gpu
 
-      - name: Check build with hpu enabled
-        run: |
-          make clippy_gpu_hpu
-
       - name: Set pull-request URL
         if: ${{ failure() && github.event_name == 'pull_request' }}
         run: |

.github/workflows/gpu_signed_integer_classic_tests.yml

@@ -109,7 +109,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -126,12 +126,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run signed integer tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_signed_integer_gpu_ci

.github/workflows/gpu_signed_integer_h100_tests.yml

@@ -122,7 +122,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -140,12 +140,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run signed integer multi-bit tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_signed_integer_multi_bit_gpu_ci

.github/workflows/gpu_signed_integer_tests.yml

@@ -25,6 +25,9 @@ on:
   # Allows you to run this workflow manually from the Actions tab as an alternative.
   workflow_dispatch:
   pull_request:
+  schedule:
+    # Nightly tests @ 1AM after each work day
+    - cron: "0 1 * * MON-FRI"
 
 permissions:
   contents: read
@@ -110,7 +113,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11
     steps:
       - name: Checkout tfhe-rs
@@ -127,12 +130,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Should run nightly tests
         if: github.event_name == 'schedule'
         run: |

.github/workflows/gpu_unsigned_integer_classic_tests.yml

@@ -109,7 +109,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -126,12 +126,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run unsigned integer tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_unsigned_integer_gpu_ci

.github/workflows/gpu_unsigned_integer_h100_tests.yml

@@ -122,7 +122,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11 
     steps:
       - name: Checkout tfhe-rs
@@ -140,12 +140,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Run unsigned integer multi-bit tests
         run: |
           BIG_TESTS_INSTANCE=TRUE make test_unsigned_integer_multi_bit_gpu_ci

.github/workflows/gpu_unsigned_integer_tests.yml

@@ -25,6 +25,9 @@ on:
   # Allows you to run this workflow manually from the Actions tab as an alternative.
   workflow_dispatch:
   pull_request:
+  schedule:
+    # Nightly tests @ 1AM after each work day
+    - cron: "0 1 * * MON-FRI"
 
 permissions:
   contents: read
@@ -110,7 +113,7 @@ jobs:
       matrix:
         include:
           - os: ubuntu-22.04
-            cuda: "12.8"
+            cuda: "12.2"
             gcc: 11
     steps:
       - name: Checkout tfhe-rs
@@ -127,12 +130,10 @@ jobs:
           github-instance: ${{ env.SECRETS_AVAILABLE == 'false' }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
-      - name: Enable nvidia multi-process service
-        run: |
-          nvidia-cuda-mps-control -d
+
       - name: Should run nightly tests
         if: github.event_name == 'schedule'
         run: |

.github/workflows/hpu_hlapi_tests.yml

@@ -13,7 +13,7 @@ env:
   CHECKOUT_TOKEN: ${{ secrets.REPO_CHECKOUT_TOKEN || secrets.GITHUB_TOKEN }}
 
 concurrency:
-  group: ${{ github.workflow }}-${{ github.head_ref }}${{ github.ref == 'refs/heads/main' && github.sha || '' }}
+  group: ${{ github.workflow }}-${{ github.head_ref }}
   cancel-in-progress: true
 
 
@@ -70,4 +70,4 @@ jobs:
           source setup_hpu.sh
           just -f mockups/tfhe-hpu-mockup/Justfile  BUILD_PROFILE=release mockup &
           make HPU_CONFIG=sim test_high_level_api_hpu
-          make HPU_CONFIG=sim test_user_doc_hpu
+     

.github/workflows/integer_long_run_tests.yml

@@ -57,7 +57,7 @@ jobs:
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/m1_tests.yml

@@ -46,7 +46,7 @@ jobs:
           token: ${{ env.CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 

.github/workflows/make_release.yml

@@ -87,7 +87,7 @@ jobs:
     # For provenance of npmjs publish
     permissions:
       contents: read
-      id-token: write
+      id-token: write # also needed for OIDC token exchange on crates.io
     steps:
       - name: Checkout
         uses: actions/checkout@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -104,16 +104,19 @@ jobs:
         with:
           name: crate
           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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          cargo publish -p tfhe ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash

.github/workflows/make_release_cuda.yml

@@ -67,7 +67,7 @@ jobs:
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
 
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -122,6 +122,9 @@ jobs:
     name: Publish CUDA Release
     needs: [setup-instance, package] # for comparing hashes
     runs-on: ${{ needs.setup-instance.outputs.runner-name }}
+    permissions:
+      # Needed for OIDC token exchange on crates.io
+      id-token: write
     strategy:
       fail-fast: false
       # explicit include-based build matrix, of known valid options
@@ -134,7 +137,7 @@ jobs:
       CUDA_PATH: /usr/local/cuda-${{ matrix.cuda }}
     steps:
       - name: Install latest stable
-        uses: dtolnay/rust-toolchain@b3b07ba8b418998c39fb20f53e8b695cdcc8de1b # zizmor: ignore[stale-action-refs] this action doesn't create releases
+        uses: dtolnay/rust-toolchain@888c2e1ea69ab0d4330cbf0af1ecc7b68f368cc1 # zizmor: ignore[stale-action-refs] this action doesn't create releases
         with:
           toolchain: stable
 
@@ -163,15 +166,19 @@ jobs:
         env:
           GCC_VERSION: ${{ matrix.gcc }}
 
+      - name: Authenticate on registry
+        uses: rust-lang/crates-io-auth-action@e919bc7605cde86df457cf5b93c5e103838bd879 # v1.0.1
+        id: auth
+
       - name: Publish crate.io package
         env:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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-cuda-backend --token "${CRATES_TOKEN}" ${DRY_RUN}
+          cargo publish -p tfhe-cuda-backend ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash

.github/workflows/make_release_hpu.yml

@@ -66,6 +66,9 @@ jobs:
     name: Publish tfhe-hpu-backend 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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -74,15 +77,19 @@ jobs:
           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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          cargo publish -p tfhe-hpu-backend ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash

.github/workflows/make_release_tfhe_csprng.yml

@@ -67,6 +67,9 @@ jobs:
     name: Publish tfhe-csprng 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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -79,15 +82,18 @@ jobs:
         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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          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}"

.github/workflows/make_release_tfhe_fft.yml

@@ -67,6 +67,9 @@ jobs:
     name: Publish tfhe-fft 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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -75,15 +78,19 @@ jobs:
           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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          cargo publish -p tfhe-fft ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash

.github/workflows/make_release_tfhe_ntt.yml

@@ -67,6 +67,9 @@ jobs:
     name: Publish tfhe-ntt 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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -75,15 +78,19 @@ jobs:
           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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          cargo publish -p tfhe-ntt ${DRY_RUN}
 
       - name: Generate hash
         id: published_hash

.github/workflows/make_release_tfhe_versionable.yml

@@ -60,6 +60,9 @@ jobs:
     name: Publish tfhe-versionable-derive Release
     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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -72,11 +75,14 @@ jobs:
         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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
         run: |
-          cargo publish -p tfhe-versionable-derive --token "${CRATES_TOKEN}"
+          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}"
@@ -149,11 +155,14 @@ jobs:
         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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          CARGO_REGISTRY_TOKEN: ${{ steps.auth.outputs.token }}
         run: |
-          cargo publish -p tfhe-versionable --token "${CRATES_TOKEN}"
+          cargo publish -p tfhe-versionable
       - name: Generate hash
         id: published_hash
         run: cd target/package && echo "pub_hash=$(sha256sum ./*.crate | base64 -w0)" >> "${GITHUB_OUTPUT}"

.github/workflows/make_release_zk_pok.yml

@@ -64,6 +64,9 @@ jobs:
     name: Publish tfhe-zk-pok 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@11bd71901bbe5b1630ceea73d27597364c9af683 # v4.2.2
@@ -76,15 +79,18 @@ jobs:
         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:
-          CRATES_TOKEN: ${{ secrets.CARGO_REGISTRY_TOKEN }}
+          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 --token "${CRATES_TOKEN}" ${DRY_RUN}
+          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}"

.github/workflows/sync_on_push.yml

@@ -21,7 +21,7 @@ jobs:
           persist-credentials: 'false'
           token: ${{ secrets.REPO_CHECKOUT_TOKEN }}
       - name: git-sync
-        uses: valtech-sd/git-sync@e734cfe9485a92e720eac5af8a4555dde5fecf88
+        uses: wei/git-sync@55c6b63b4f21607da0e9877ca9b4d11a29fc6d83
         with:
           source_repo: "zama-ai/tfhe-rs"
           source_branch: "main"

.github/workflows/unverified_prs.yml

@@ -1,26 +0,0 @@
-name: 'Close unverified PRs'
-on:
-  schedule:
-    - cron: '30 1 * * *'
-
-permissions: {}
-
-jobs:
-  stale:
-    runs-on: ubuntu-latest
-    permissions:
-      issues: read
-      pull-requests: write
-    steps:
-      - uses: actions/stale@5bef64f19d7facfb25b37b414482c7164d639639 # v9.1.0
-        with:
-          stale-pr-message: 'This PR is unverified and has been open for 2 days, it will now be closed. If you want to contribute please sign the CLA as indicated by the bot.'
-          days-before-stale: 2
-          days-before-close: 0
-          # We are not interested in suppressing issues so have a currently non existent label
-          # if we ever accept issues to become stale/closable this label will be the signal for that
-          only-issue-labels: can-be-auto-closed
-          # Only unverified PRs are an issue
-          exempt-pr-labels: cla-signed
-          # We don't want people commenting to keep an unverified PR
-          ignore-updates: true

.gitignore

@@ -36,7 +36,7 @@ package-lock.json
 .env
 __pycache__
 
+# Dir used for backward compatibility test data
+# First directive is to ignore symlinks
+tests/tfhe-backward-compat-data
 ci/
-
-# In case someone clones the lattice-estimator locally to verify security
-/lattice-estimator

.linelint.yml

@@ -10,7 +10,6 @@ ignore:
   - keys
   - coverage
   - utils/tfhe-lints/ui/main.stderr
-  - utils/tfhe-backward-compat-data/**/*.ron # ron files are autogenerated
 
 rules:
   # checks if file ends in a newline character

CODEOWNERS

@@ -1,28 +1,12 @@
 # Specifying a path without code owners means that path won't have owners and is akin to a negation
 # i.e. the `core_crypto` dir is owned and needs owner approval/review, but not the `gpu` sub dir
 # See https://docs.github.com/en/repositories/managing-your-repositorys-settings-and-features/customizing-your-repository/about-code-owners#example-of-a-codeowners-file
-
-/backends/tfhe-cuda-backend/            @agnesLeroy
-/backends/tfhe-hpu-backend/             @zama-ai/hardware
-
-/tfhe/examples/hpu                      @zama-ai/hardware
-
 /tfhe/src/core_crypto/                  @IceTDrinker
-/tfhe/src/core_crypto/gpu               @agnesLeroy
-/tfhe/src/core_crypto/hpu               @zama-ai/hardware
+/tfhe/src/core_crypto/gpu
 
 /tfhe/src/shortint/                     @mayeul-zama
 
 /tfhe/src/integer/                      @tmontaigu
-/tfhe/src/integer/gpu                   @agnesLeroy
-/tfhe/src/integer/hpu                   @zama-ai/hardware
+/tfhe/src/integer/gpu
 
 /tfhe/src/high_level_api/               @tmontaigu
-
-/Makefile                               @IceTDrinker @soonum
-
-/mockups/tfhe-hpu-mockup                @zama-ai/hardware
-
-/.github/                               @soonum
-
-/CODEOWNERS                             @IceTDrinker

CONTRIBUTING.md

@@ -170,8 +170,6 @@ On the contrary, these changes are *not* data breaking:
  * Renaming a type (unless it implements the `Named` trait).
  * Adding a variant to the end of an enum.
 
-Historical data from previous TFHE-rs versions are stored inside `utils/tfhe-backward-compat-data`. They are used to check on every PR that backward compatibility has been preserved.
-
 ## Example: adding a field
 
 Suppose you want to add an i32 field to a type named `MyType`. The original type is defined as:

Cargo.toml

@@ -18,7 +18,7 @@ members = [
 ]
 
 exclude = [
-    "utils/tfhe-backward-compat-data",
+    "tests/backward_compatibility_tests",
     "utils/tfhe-lints",
     "apps/trivium",
 ]
@@ -30,7 +30,7 @@ itertools = "0.14"
 num-complex = "0.4"
 pulp = { version = "0.21", default-features = false }
 rand = "0.8"
-rayon = "1"
+rayon = "1.11"
 serde = { version = "1.0", default-features = false }
 wasm-bindgen = "0.2.100"
 

Makefile

@@ -22,7 +22,10 @@ BENCH_TYPE?=latency
 BENCH_PARAM_TYPE?=classical
 BENCH_PARAMS_SET?=default
 NODE_VERSION=22.6
-BACKWARD_COMPAT_DATA_DIR=utils/tfhe-backward-compat-data
+BACKWARD_COMPAT_DATA_URL=https://github.com/zama-ai/tfhe-backward-compat-data.git
+BACKWARD_COMPAT_DATA_BRANCH?=$(shell ./scripts/backward_compat_data_version.py)
+BACKWARD_COMPAT_DATA_PROJECT=tfhe-backward-compat-data
+BACKWARD_COMPAT_DATA_DIR=$(BACKWARD_COMPAT_DATA_PROJECT)
 TFHE_SPEC:=tfhe
 WASM_PACK_VERSION="0.13.1"
 # We are kind of hacking the cut here, the version cannot contain a quote '"'
@@ -156,12 +159,12 @@ install_tarpaulin: install_rs_build_toolchain
 
 .PHONY: install_cargo_dylint # Install custom tfhe-rs lints
 install_cargo_dylint:
-	cargo install --locked cargo-dylint dylint-link
+	cargo install cargo-dylint dylint-link
 
 .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 || \
+	cargo $(CARGO_RS_BUILD_TOOLCHAIN) install typos-cli || \
 	( echo "Unable to install typos-cli, unknown error." && exit 1 )
 
 .PHONY: install_zizmor # Install zizmor workflow security checker
@@ -170,10 +173,6 @@ install_zizmor: install_rs_build_toolchain
 	cargo $(CARGO_RS_BUILD_TOOLCHAIN) install --locked zizmor --version ~1.9 || \
 	( echo "Unable to install zizmor, unknown error." && exit 1 )
 
-.PHONY: install_cargo_cross # Install cross for big endian tests
-install_cargo_cross: install_rs_build_toolchain
-	cargo $(CARGO_RS_BUILD_TOOLCHAIN) install --locked cross
-
 .PHONY: setup_venv # Setup Python virtualenv for wasm tests
 setup_venv:
 	python3 -m venv venv
@@ -249,9 +248,6 @@ install_mlc: install_rs_build_toolchain
 .PHONY: fmt # Format rust code
 fmt: install_rs_check_toolchain
 	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" fmt
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C $(BACKWARD_COMPAT_DATA_DIR) fmt
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C utils/tfhe-lints fmt
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C apps/trivium fmt
 
 .PHONY: fmt_js # Format javascript code
 fmt_js: check_nvm_installed
@@ -273,9 +269,6 @@ fmt_c_tests:
 .PHONY: check_fmt # Check rust code format
 check_fmt: install_rs_check_toolchain
 	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" fmt --check
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C $(BACKWARD_COMPAT_DATA_DIR) fmt --check
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C utils/tfhe-lints fmt --check
-	cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options -C apps/trivium fmt --check
 
 .PHONY: check_fmt_c_tests  # Check C tests format
 check_fmt_c_tests:
@@ -301,7 +294,7 @@ check_typos: install_typos_checker
 .PHONY: clippy_gpu # Run clippy lints on tfhe with "gpu" enabled
 clippy_gpu: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy \
-		--features=boolean,shortint,integer,internal-keycache,gpu,pbs-stats,extended-types,zk-pok \
+		--features=boolean,shortint,integer,internal-keycache,gpu,pbs-stats,extended-types \
 		--all-targets \
 		-p $(TFHE_SPEC) -- --no-deps -D warnings
 
@@ -319,13 +312,6 @@ clippy_hpu: install_rs_check_toolchain
 		--all-targets \
 		-p $(TFHE_SPEC) -- --no-deps -D warnings
 
-.PHONY: clippy_gpu_hpu # Run clippy lints on tfhe with "gpu" and "hpu" enabled
-clippy_gpu_hpu: install_rs_check_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy \
-		--features=boolean,shortint,integer,internal-keycache,gpu,hpu,pbs-stats,extended-types,zk-pok \
-		--all-targets \
-		-p $(TFHE_SPEC) -- --no-deps -D warnings
-
 .PHONY: fix_newline # Fix newline at end of file issues to be UNIX compliant
 fix_newline: check_linelint_installed
 	linelint -a .
@@ -465,8 +451,6 @@ clippy_tfhe_csprng: install_rs_check_toolchain
 clippy_zk_pok: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy --all-targets \
 		-p tfhe-zk-pok -- --no-deps -D warnings
-	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy --all-targets \
-		-p tfhe-zk-pok --features=experimental -- --no-deps -D warnings
 
 .PHONY: clippy_versionable # Run clippy lints on tfhe-versionable
 clippy_versionable: install_rs_check_toolchain
@@ -486,22 +470,10 @@ clippy_param_dedup: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy --all-targets \
 		-p param_dedup -- --no-deps -D warnings
 
-.PHONY: clippy_backward_compat_data # Run clippy lints on tfhe-backward-compat-data
-clippy_backward_compat_data: install_rs_check_toolchain # the toolchain is selected with toolchain.toml
-	@# Some old crates are x86 specific, only run in that case
-	@if uname -a | grep -q x86; then \
-		RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" -Z unstable-options \
-			-C $(BACKWARD_COMPAT_DATA_DIR) clippy --all-targets \
-			-- --no-deps -D warnings; \
-	else \
-		echo "Cannot run clippy for backward compat crate on non x86 platform for now."; \
-	fi
-
 .PHONY: clippy_all # Run all clippy targets
 clippy_all: clippy_rustdoc clippy clippy_boolean clippy_shortint clippy_integer clippy_all_targets \
 clippy_c_api clippy_js_wasm_api clippy_tasks clippy_core clippy_tfhe_csprng clippy_zk_pok clippy_trivium \
-clippy_versionable clippy_tfhe_lints clippy_ws_tests clippy_bench clippy_param_dedup \
-clippy_backward_compat_data
+clippy_versionable clippy_tfhe_lints clippy_ws_tests clippy_bench clippy_param_dedup
 
 .PHONY: clippy_fast # Run main clippy targets
 clippy_fast: clippy_rustdoc clippy clippy_all_targets clippy_c_api clippy_js_wasm_api clippy_tasks \
@@ -517,12 +489,6 @@ clippy_hpu_backend: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy --all-targets \
 		-p tfhe-hpu-backend -- --no-deps -D warnings
 
-.PHONY: clippy_hpu_mockup # Run clippy lints on tfhe-hpu-mockup
-clippy_hpu_mockup: install_rs_check_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" cargo "$(CARGO_RS_CHECK_TOOLCHAIN)" clippy \
-		--all-targets \
-		-p tfhe-hpu-mockup -- --no-deps -D warnings
-
 .PHONY: check_rust_bindings_did_not_change # Check rust bindings are up to date for tfhe-cuda-backend
 check_rust_bindings_did_not_change:
 	cargo build -p tfhe-cuda-backend && "$(MAKE)" fmt_gpu && \
@@ -535,9 +501,6 @@ check_rust_bindings_did_not_change:
 tfhe_lints: install_cargo_dylint
 	RUSTFLAGS="$(RUSTFLAGS)" cargo dylint --all -p tfhe --no-deps -- \
 		--features=boolean,shortint,integer,strings,zk-pok
-	RUSTFLAGS="$(RUSTFLAGS)" cargo dylint --all -p tfhe-zk-pok --no-deps -- \
-		--features=experimental
-
 
 .PHONY: build_core # Build core_crypto without experimental features
 build_core: install_rs_build_toolchain install_rs_check_toolchain
@@ -679,31 +642,6 @@ test_integer_gpu: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --doc --profile $(CARGO_PROFILE) \
 		--features=integer,gpu -p $(TFHE_SPEC) -- integer::gpu::server_key::
 
-.PHONY: test_integer_gpu_debug # Run the tests of the integer module with Debug flags for CUDA
-test_integer_gpu_debug: install_rs_build_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile release_lto_off \
-		--features=integer,gpu-debug -vv -p $(TFHE_SPEC) -- integer::gpu::server_key:: --test-threads=1 --nocapture
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --doc --profile release_lto_off \
-		--features=integer,gpu-debug -p $(TFHE_SPEC) -- integer::gpu::server_key::
-
-.PHONY: test_high_level_api_gpu_debug # Run the tests of the integer module with Debug flags for CUDA
-test_high_level_api_gpu_debug: install_rs_build_toolchain install_cargo_nextest
-	export RUSTFLAGS="$(RUSTFLAGS)" && \
-	export CARGO_RS_BUILD_TOOLCHAIN="$(CARGO_RS_BUILD_TOOLCHAIN)" && \
-	export TFHE_SPEC="$(TFHE_SPEC)" && \
-	export CARGO_PROFILE="$(CARGO_PROFILE)" &&	scripts/check_memory_errors.sh
-
-.PHONY: test_integer_hl_test_gpu_check_warnings
-test_integer_hl_test_gpu_check_warnings: install_rs_build_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) build \
-		--features=integer,internal-keycache,gpu-debug,zk-pok -vv -p $(TFHE_SPEC) &> /tmp/gpu_compile_output
-	WARNINGS=$$(cat /tmp/gpu_compile_output | grep ": warning #" | grep "\[tfhe-cuda-backend" | grep -v "inline qualifier" || true) && \
-	if [[ "$${WARNINGS}" != "" ]]; then \
-	    echo "FAILING BECAUSE CUDA COMPILATION WARNINGS WERE DETECTED: " && \
-		echo "$${WARNINGS}" && exit 1; \
-	fi
-
-
 .PHONY: test_integer_long_run_gpu # Run the long run integer tests on the gpu backend
 test_integer_long_run_gpu: install_rs_check_toolchain install_cargo_nextest
 	BIG_TESTS_INSTANCE="$(BIG_TESTS_INSTANCE)" \
@@ -712,12 +650,6 @@ test_integer_long_run_gpu: install_rs_check_toolchain install_cargo_nextest
 		--cargo-profile "$(CARGO_PROFILE)" --avx512-support "$(AVX512_SUPPORT)" \
 		--tfhe-package "$(TFHE_SPEC)" --backend "gpu"
 
-.PHONY: test_integer_short_run_gpu # Run the long run integer tests on the gpu backend
-test_integer_short_run_gpu: install_rs_check_toolchain install_cargo_nextest
-	TFHE_RS_TEST_LONG_TESTS_MINIMAL=TRUE \
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		--features=integer,gpu -p $(TFHE_SPEC) -- integer::gpu::server_key::radix::tests_long_run::test_random_op_sequence integer::gpu::server_key::radix::tests_long_run::test_signed_random_op_sequence --test-threads=1 --nocapture
-
 .PHONY: test_integer_compression
 test_integer_compression: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
@@ -928,15 +860,6 @@ test_integer_long_run: install_rs_check_toolchain install_cargo_nextest
 		--cargo-profile "$(CARGO_PROFILE)" --avx512-support "$(AVX512_SUPPORT)" \
 		--tfhe-package "$(TFHE_SPEC)"
 
-.PHONY: test_noise_check # Run dedicated noise and pfail check tests
-test_noise_check: install_rs_check_toolchain
-	@# First run the sanity checks to make sure the atomic patterns are correct
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_CHECK_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		--features=boolean,shortint,integer,nightly-avx512 -p $(TFHE_SPEC) -- sanity_check
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_CHECK_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		--features=boolean,shortint,integer,nightly-avx512 -p $(TFHE_SPEC) -- noise_check \
-		--test-threads=1 --nocapture
-
 .PHONY: test_safe_serialization # Run the tests for safe serialization
 test_safe_serialization: install_rs_build_toolchain install_cargo_nextest
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
@@ -969,7 +892,7 @@ test_high_level_api: install_rs_build_toolchain
 
 test_high_level_api_gpu: install_rs_build_toolchain install_cargo_nextest
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) nextest run --cargo-profile $(CARGO_PROFILE) \
-		--test-threads=4 --features=integer,internal-keycache,gpu,zk-pok -p $(TFHE_SPEC) \
+		--features=integer,internal-keycache,gpu -p $(TFHE_SPEC) \
 		-E "test(/high_level_api::.*gpu.*/)"
 
 test_high_level_api_hpu: install_rs_build_toolchain install_cargo_nextest
@@ -1005,21 +928,9 @@ test_user_doc: install_rs_build_toolchain
 .PHONY: test_user_doc_gpu # Run tests for GPU from the .md documentation
 test_user_doc_gpu: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) --doc \
-		--features=internal-keycache,integer,zk-pok,gpu -p $(TFHE_SPEC) \
+		--features=boolean,shortint,integer,internal-keycache,gpu,zk-pok -p $(TFHE_SPEC) \
 		-- test_user_docs::
 
-.PHONY: test_user_doc_hpu # Run tests for HPU from the .md documentation
-test_user_doc_hpu: install_rs_build_toolchain
-ifeq ($(HPU_CONFIG), v80)
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) --doc \
-		--features=internal-keycache,integer,hpu,hpu-v80 -p $(TFHE_SPEC) \
-		-- test_user_docs::
-else
-	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) --doc \
-		--features=internal-keycache,integer,hpu -p $(TFHE_SPEC) \
-		-- test_user_docs::
-endif
-
 
 
 .PHONY: test_regex_engine # Run tests for regex_engine example
@@ -1050,16 +961,10 @@ test_tfhe_csprng: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
 		-p tfhe-csprng
 
-.PHONY: test_tfhe_csprng_big_endian # Run tfhe-csprng tests on an emulated big endian system
-test_tfhe_csprng_big_endian: install_rs_build_toolchain install_cargo_cross
-	RUSTFLAGS="" cross $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		-p tfhe-csprng --target=powerpc64-unknown-linux-gnu
-
-
 .PHONY: test_zk_pok # Run tfhe-zk-pok tests
 test_zk_pok: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
-		-p tfhe-zk-pok --features experimental
+		-p tfhe-zk-pok
 
 .PHONY: test_zk_wasm_x86_compat_ci
 test_zk_wasm_x86_compat_ci: check_nvm_installed
@@ -1089,11 +994,16 @@ test_tfhe_lints: install_cargo_dylint
 # Here we use the "patch" functionality of Cargo to make sure the repo used for the data is the same as the one used for the code.
 .PHONY: test_backward_compatibility_ci
 test_backward_compatibility_ci: install_rs_build_toolchain
-	TFHE_BACKWARD_COMPAT_DATA_DIR="../$(BACKWARD_COMPAT_DATA_DIR)" RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
+	TFHE_BACKWARD_COMPAT_DATA_DIR="$(BACKWARD_COMPAT_DATA_DIR)" RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --profile $(CARGO_PROFILE) \
+		--config "patch.'$(BACKWARD_COMPAT_DATA_URL)'.$(BACKWARD_COMPAT_DATA_PROJECT).path=\"tests/$(BACKWARD_COMPAT_DATA_DIR)\"" \
 		--features=shortint,integer,zk-pok -p tests test_backward_compatibility -- --nocapture
 
 .PHONY: test_backward_compatibility # Same as test_backward_compatibility_ci but tries to clone the data repo first if needed
-test_backward_compatibility: pull_backward_compat_data test_backward_compatibility_ci
+test_backward_compatibility: tests/$(BACKWARD_COMPAT_DATA_DIR) test_backward_compatibility_ci
+
+.PHONY: backward_compat_branch # Prints the required backward compatibility branch
+backward_compat_branch:
+	@echo "$(BACKWARD_COMPAT_DATA_BRANCH)"
 
 .PHONY: doc # Build rust doc
 doc: install_rs_check_toolchain
@@ -1138,10 +1048,6 @@ check_intra_md_links: install_mlc
 check_md_links: install_mlc
 	mlc --match-file-extension tfhe/docs
 
-.PHONY: check_doc_paths_use_dash # Check paths use "-" instead of "_" in docs for gitbook compatibility
-check_doc_paths_use_dash:
-	python3 ./scripts/check_doc_paths_use_dash.py
-
 .PHONY: check_parameter_export_ok # Checks exported "current" shortint parameter module is correct
 check_parameter_export_ok:
 	python3 ./scripts/check_current_param_export.py
@@ -1160,7 +1066,7 @@ check_compile_tests: install_rs_build_toolchain
 .PHONY: check_compile_tests_benches_gpu # Build tests in debug without running them
 check_compile_tests_benches_gpu: install_rs_build_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_BUILD_TOOLCHAIN) test --no-run \
-		--features=experimental,boolean,shortint,integer,internal-keycache,gpu,zk-pok \
+		--features=experimental,boolean,shortint,integer,internal-keycache,gpu \
 		-p $(TFHE_SPEC)
 	mkdir -p "$(TFHECUDA_BUILD)" && \
 		cd "$(TFHECUDA_BUILD)" && \
@@ -1308,7 +1214,7 @@ bench_integer_compression_gpu: install_rs_check_toolchain
 	cargo $(CARGO_RS_CHECK_TOOLCHAIN) bench \
 	--bench	glwe_packing_compression-integer-bench \
 	--features=integer,internal-keycache,gpu,pbs-stats -p tfhe-benchmark --
-
+	
 .PHONY: bench_integer_zk_gpu
 bench_integer_zk_gpu: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" __TFHE_RS_BENCH_TYPE=$(BENCH_TYPE) \
@@ -1499,20 +1405,6 @@ bench_tfhe_zk_pok: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" \
 	cargo $(CARGO_RS_CHECK_TOOLCHAIN) bench -p tfhe-zk-pok --
 
-.PHONY: bench_hlapi_noise_squash # Run benchmarks for noise squash operation
-bench_hlapi_noise_squash: install_rs_check_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" __TFHE_RS_BENCH_TYPE=$(BENCH_TYPE) \
-	cargo $(CARGO_RS_CHECK_TOOLCHAIN) bench \
-	--bench hlapi-noise-squash \
-	--features=integer,internal-keycache,pbs-stats,nightly-avx512 -p tfhe-benchmark --
-
-.PHONY: bench_hlapi_noise_squash_gpu # Run benchmarks for noise squash operation on GPU
-bench_hlapi_noise_squash_gpu: install_rs_check_toolchain
-	RUSTFLAGS="$(RUSTFLAGS)" __TFHE_RS_BENCH_TYPE=$(BENCH_TYPE) \
-	cargo $(CARGO_RS_CHECK_TOOLCHAIN) bench \
-	--bench hlapi-noise-squash \
-	--features=integer,gpu,internal-keycache,pbs-stats,nightly-avx512 -p tfhe-benchmark --
-
 #
 # Utility tools
 #
@@ -1568,15 +1460,13 @@ parse_wasm_benchmarks: install_rs_check_toolchain
 .PHONY: write_params_to_file # Gather all crypto parameters into a file with a Sage readable format.
 write_params_to_file: install_rs_check_toolchain
 	RUSTFLAGS="$(RUSTFLAGS)" cargo $(CARGO_RS_CHECK_TOOLCHAIN) run \
-	--example write_params_to_file --features=boolean,shortint,hpu,internal-keycache
+	--example write_params_to_file --features=boolean,shortint,internal-keycache
 
-.PHONY: pull_backward_compat_data # Pull the data files needed for backward compatibility tests
-pull_backward_compat_data:
-	./scripts/pull_lfs_data.sh $(BACKWARD_COMPAT_DATA_DIR)
+.PHONY: clone_backward_compat_data # Clone the data repo needed for backward compatibility tests
+clone_backward_compat_data:
+	./scripts/clone_backward_compat_data.sh $(BACKWARD_COMPAT_DATA_URL) $(BACKWARD_COMPAT_DATA_BRANCH) tests/$(BACKWARD_COMPAT_DATA_DIR)
 
-.PHONY: pull_hpu_files # Pull the hpu files
-pull_hpu_files:
-	./scripts/pull_lfs_data.sh backends/tfhe-hpu-backend/
+tests/$(BACKWARD_COMPAT_DATA_DIR): clone_backward_compat_data
 
 #
 # Real use case examples
@@ -1602,20 +1492,19 @@ sha256_bool: install_rs_check_toolchain
 
 .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 \
+check_md_docs_are_tested check_intra_md_links clippy_all check_compile_tests test_tfhe_lints \
 tfhe_lints
 
 .PHONY: pcc_gpu # pcc stands for pre commit checks for GPU compilation
 pcc_gpu: check_rust_bindings_did_not_change clippy_rustdoc_gpu \
-clippy_gpu clippy_cuda_backend clippy_bench_gpu check_compile_tests_benches_gpu test_integer_hl_test_gpu_check_warnings
+clippy_gpu clippy_cuda_backend clippy_bench_gpu check_compile_tests_benches_gpu
 
 .PHONY: pcc_hpu # pcc stands for pre commit checks for HPU compilation
-pcc_hpu: clippy_hpu clippy_hpu_backend clippy_hpu_mockup test_integer_hpu_mockup_ci_fast
+pcc_hpu: clippy_hpu clippy_hpu_backend test_integer_hpu_mockup_ci_fast 
 
 .PHONY: fpcc # pcc stands for pre commit checks, the f stands for fast
 fpcc: 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_fast check_compile_tests
+check_md_docs_are_tested clippy_fast check_compile_tests
 
 .PHONY: conformance # Automatically fix problems that can be fixed
 conformance: fix_newline fmt fmt_js

README.md

@@ -18,7 +18,6 @@
   <a href="https://github.com/zama-ai/tfhe-rs/releases"><img src="https://img.shields.io/github/v/release/zama-ai/tfhe-rs?style=flat-square"></a>
   <a href="LICENSE"><img src="https://img.shields.io/badge/License-BSD--3--Clause--Clear-%23ffb243?style=flat-square"></a>
   <a href="https://github.com/zama-ai/bounty-program"><img src="https://img.shields.io/badge/Contribute-Zama%20Bounty%20Program-%23ffd208?style=flat-square"></a>
-  <a href="https://slsa.dev"><img alt="SLSA 3" src="https://slsa.dev/images/gh-badge-level3.svg" /></a>
 </p>
 
 ## About
@@ -149,7 +148,7 @@ To run this code, use the following command:
 > 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.
 
-*Find an example with more explanations in [this part of the documentation](https://docs.zama.ai/tfhe-rs/get-started/quick-start)*
+*Find an example with more explanations in [this part of the documentation](https://docs.zama.ai/tfhe-rs/get-started/quick_start)*
 
 <p align="right">
   <a href="#about" > ↑ Back to top </a>
@@ -204,7 +203,7 @@ When a new update is published in the Lattice Estimator, we update parameters ac
 By default, the parameter sets used in the High-Level API with the x86 CPU backend have a failure probability $\le 2^{128}$ to securely work in the IND-CPA^D model using the algorithmic techniques provided in our code base [1].
 If you want to work within the IND-CPA security model, which is less strict than the IND-CPA-D model, the parameter sets can easily be changed and would have slightly better performance. More details can be found in the [TFHE-rs documentation](https://docs.zama.ai/tfhe-rs).
 
-The default parameters used in the High-Level API with the GPU backend are chosen considering the IND-CPA security model, and are selected with a bootstrapping failure probability fixed at $p_{error} \le 2^{-128}$. In particular, it is assumed that the results of decrypted computations are not shared by the secret key owner with any third parties, as such an action can lead to leakage of the secret encryption key. If you are designing an application where decryptions must be shared, you will need to craft custom encryption parameters which are chosen in consideration of the IND-CPA^D security model [2].
+The default parameters used in the High-Level API with the GPU backend are chosen considering the IND-CPA security model, and are selected with a bootstrapping failure probability fixed at $p_{error} \le 2^{-64}$. In particular, it is assumed that the results of decrypted computations are not shared by the secret key owner with any third parties, as such an action can lead to leakage of the secret encryption key. If you are designing an application where decryptions must be shared, you will need to craft custom encryption parameters which are chosen in consideration of the IND-CPA^D security model [2].
 
 [1] Bernard, Olivier, et al. "Drifting Towards Better Error Probabilities in Fully Homomorphic Encryption Schemes". https://eprint.iacr.org/2024/1718.pdf
 

apps/trivium/README.md

@@ -129,7 +129,7 @@ Other sizes than 64 bit are expected to be available in the future.
 
 # FHE shortint Trivium implementation
 
-The same implementation is also available for generic Ciphertexts representing bits (meant to be used with parameters `V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128`).
+The same implementation is also available for generic Ciphertexts representing bits (meant to be used with parameters `V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128`).
 It uses a lower level API of tfhe-rs, so the syntax is a little bit different. It also implements the `TransCiphering` trait. For optimization purposes, it does not internally run
 on the same cryptographic parameters as the high level API of tfhe-rs. As such, it requires the usage of a casting key, to switch from one parameter space to another, which makes
 its setup a little more intricate.
@@ -137,10 +137,10 @@ its setup a little more intricate.
 Example code:
 ```rust
 use tfhe::shortint::prelude::*;
-use tfhe::shortint::parameters::current_params::{
-    V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+use tfhe::shortint::parameters::v1_2::{
+    V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
 };
 use tfhe::{ConfigBuilder, generate_keys, FheUint64};
 use tfhe::prelude::*;
@@ -148,17 +148,17 @@ use tfhe_trivium::TriviumStreamShortint;
 
 fn test_shortint() {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
-    let (client_key, server_key): (ClientKey, ServerKey) = gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+    let (client_key, server_key): (ClientKey, ServerKey) = gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB".to_string();

apps/trivium/benches/kreyvium_shortint.rs

@@ -1,9 +1,9 @@
 use criterion::Criterion;
 use tfhe::prelude::*;
-use tfhe::shortint::parameters::current_params::{
-    V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
+use tfhe::shortint::parameters::v1_2::{
+    V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
 };
 use tfhe::shortint::prelude::*;
 use tfhe::{generate_keys, ConfigBuilder, FheUint64};
@@ -11,19 +11,19 @@ use tfhe_trivium::{KreyviumStreamShortint, TransCiphering};
 
 pub fn kreyvium_shortint_warmup(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB000000000000".to_string();
@@ -64,19 +64,19 @@ pub fn kreyvium_shortint_warmup(c: &mut Criterion) {
 
 pub fn kreyvium_shortint_gen(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB000000000000".to_string();
@@ -112,19 +112,19 @@ pub fn kreyvium_shortint_gen(c: &mut Criterion) {
 
 pub fn kreyvium_shortint_trans(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB000000000000".to_string();

apps/trivium/benches/trivium_shortint.rs

@@ -1,9 +1,9 @@
 use criterion::Criterion;
 use tfhe::prelude::*;
-use tfhe::shortint::parameters::current_params::{
-    V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
+use tfhe::shortint::parameters::v1_2::{
+    V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
 };
 use tfhe::shortint::prelude::*;
 use tfhe::{generate_keys, ConfigBuilder, FheUint64};
@@ -11,19 +11,19 @@ use tfhe_trivium::{TransCiphering, TriviumStreamShortint};
 
 pub fn trivium_shortint_warmup(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB".to_string();
@@ -64,19 +64,19 @@ pub fn trivium_shortint_warmup(c: &mut Criterion) {
 
 pub fn trivium_shortint_gen(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB".to_string();
@@ -112,19 +112,19 @@ pub fn trivium_shortint_gen(c: &mut Criterion) {
 
 pub fn trivium_shortint_trans(c: &mut Criterion) {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB".to_string();

apps/trivium/src/kreyvium/test.rs

@@ -1,9 +1,9 @@
 use crate::{KreyviumStream, KreyviumStreamByte, KreyviumStreamShortint, TransCiphering};
 use tfhe::prelude::*;
-use tfhe::shortint::parameters::current_params::{
-    V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
+use tfhe::shortint::parameters::v1_2::{
+    V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
 };
 use tfhe::{generate_keys, ConfigBuilder, FheBool, FheUint64, FheUint8};
 // Values for these tests come from the github repo renaud1239/Kreyvium,
@@ -221,19 +221,19 @@ use tfhe::shortint::prelude::*;
 #[test]
 fn kreyvium_test_shortint_long() {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB000000000000".to_string();

apps/trivium/src/trivium/test.rs

@@ -1,9 +1,9 @@
 use crate::{TransCiphering, TriviumStream, TriviumStreamByte, TriviumStreamShortint};
 use tfhe::prelude::*;
-use tfhe::shortint::parameters::current_params::{
-    V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
-    V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
+use tfhe::shortint::parameters::v1_2::{
+    V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128,
+    V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128,
 };
 use tfhe::{generate_keys, ConfigBuilder, FheBool, FheUint64, FheUint8};
 // Values for these tests come from the github repo cantora/avr-crypto-lib, commit 2a5b018,
@@ -357,19 +357,19 @@ use tfhe::shortint::prelude::*;
 #[test]
 fn trivium_test_shortint_long() {
     let config = ConfigBuilder::default()
-        .use_custom_parameters(V1_3_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
+        .use_custom_parameters(V1_2_PARAM_MESSAGE_2_CARRY_2_KS_PBS_GAUSSIAN_2M128)
         .build();
     let (hl_client_key, hl_server_key) = generate_keys(config);
     let underlying_ck: tfhe::shortint::ClientKey = (*hl_client_key.as_ref()).clone().into();
     let underlying_sk: tfhe::shortint::ServerKey = (*hl_server_key.as_ref()).clone().into();
 
     let (client_key, server_key): (ClientKey, ServerKey) =
-        gen_keys(V1_3_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
+        gen_keys(V1_2_PARAM_MESSAGE_1_CARRY_1_KS_PBS_GAUSSIAN_2M128);
 
     let ksk = KeySwitchingKey::new(
         (&client_key, Some(&server_key)),
         (&underlying_ck, &underlying_sk),
-        V1_3_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
+        V1_2_PARAM_KEYSWITCH_1_1_KS_PBS_TO_2_2_KS_PBS_GAUSSIAN_2M128,
     );
 
     let key_string = "0053A6F94C9FF24598EB".to_string();

backends/tfhe-cuda-backend/Cargo.toml

@@ -1,6 +1,6 @@
 [package]
 name = "tfhe-cuda-backend"
-version = "0.11.0"
+version = "0.10.1"
 edition = "2021"
 authors = ["Zama team"]
 license = "BSD-3-Clause-Clear"
@@ -19,4 +19,3 @@ bindgen = "0.71"
 [features]
 experimental-multi-arch = []
 profile = []
-debug = []

backends/tfhe-cuda-backend/build.rs

@@ -53,10 +53,6 @@ fn main() {
             cmake_config.define("USE_NVTOOLS", "OFF");
         }
 
-        if cfg!(feature = "debug") {
-            cmake_config.define("CMAKE_BUILD_TYPE", "Debug");
-        }
-
         // Build the CMake project
         let dest = cmake_config.build();
         println!("cargo:rustc-link-search=native={}", dest.display());

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

@@ -52,8 +52,6 @@ endif()
 
 if(NOT CMAKE_BUILD_TYPE)
   set(CMAKE_BUILD_TYPE Release)
-else()
-  message("Building CUDA backend in ${CMAKE_BUILD_TYPE}")
 endif()
 
 # Add OpenMP support
@@ -78,10 +76,8 @@ endif()
 
 add_compile_definitions(CUDA_ARCH=${CUDA_ARCH})
 
-string(TOLOWER "${CMAKE_BUILD_TYPE}" CMAKE_BUILD_TYPE_LOWERCASE)
-
 # Check if the DEBUG flag is defined
-if(CMAKE_BUILD_TYPE_LOWERCASE STREQUAL "debug")
+if(CMAKE_BUILD_TYPE STREQUAL "Debug")
   # Debug mode
   message("Compiling in Debug mode")
   add_definitions(-DDEBUG)

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

@@ -28,10 +28,9 @@ void cuda_modulus_switch_inplace_64(void *stream, uint32_t gpu_index,
 
 void cuda_improve_noise_modulus_switch_64(
     void *stream, uint32_t gpu_index, void *lwe_array_out,
-    void const *lwe_array_in, void const *lwe_array_indexes,
-    void const *encrypted_zeros, uint32_t lwe_size, uint32_t num_lwes,
-    uint32_t num_zeros, double input_variance, double r_sigma, double bound,
-    uint32_t log_modulus);
+    void const *lwe_array_in, void const *encrypted_zeros, uint32_t lwe_size,
+    uint32_t num_lwes, uint32_t num_zeros, double input_variance,
+    double r_sigma, double bound, uint32_t log_modulus);
 
 void cuda_glwe_sample_extract_128(
     void *stream, uint32_t gpu_index, void *lwe_array_out,

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

@@ -26,7 +26,6 @@ inline void cuda_error(cudaError_t code, const char *file, int line) {
     std::abort();                                                              \
   }
 
-uint32_t cuda_get_device();
 void cuda_set_device(uint32_t gpu_index);
 
 cudaEvent_t cuda_create_event(uint32_t gpu_index);
@@ -50,13 +49,12 @@ void *cuda_malloc(uint64_t size, uint32_t gpu_index);
 
 void *cuda_malloc_with_size_tracking_async(uint64_t size, cudaStream_t stream,
                                            uint32_t gpu_index,
-                                           uint64_t &size_tracker,
+                                           uint64_t *size_tracker,
                                            bool allocate_gpu_memory);
 
 void *cuda_malloc_async(uint64_t size, cudaStream_t stream, uint32_t gpu_index);
 
 bool cuda_check_valid_malloc(uint64_t size, uint32_t gpu_index);
-uint64_t cuda_device_total_memory(uint32_t gpu_index);
 
 void cuda_memcpy_with_size_tracking_async_to_gpu(void *dest, const void *src,
                                                  uint64_t size,

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

@@ -6,10 +6,9 @@
 
 extern std::mutex m;
 extern bool p2p_enabled;
-extern const int THRESHOLD_MULTI_GPU;
 
 extern "C" {
-int32_t cuda_setup_multi_gpu(int device_0_id);
+int32_t cuda_setup_multi_gpu();
 }
 
 // Define a variant type that can be either a vector or a single pointer
@@ -25,15 +24,7 @@ using LweArrayVariant = std::variant<std::vector<Torus *>, Torus *>;
       return std::get<Torus *>(variant);                                       \
     }                                                                          \
   }()
-// Macro to define the visitor logic using std::holds_alternative for vectors
-#define GET_VARIANT_ELEMENT_64BIT(variant, index)                              \
-  [&] {                                                                        \
-    if (std::holds_alternative<std::vector<uint64_t *>>(variant)) {            \
-      return std::get<std::vector<uint64_t *>>(variant)[index];                \
-    } else {                                                                   \
-      return std::get<uint64_t *>(variant);                                    \
-    }                                                                          \
-  }()
+
 int get_active_gpu_count(int num_inputs, int gpu_count);
 
 int get_num_inputs_on_gpu(int total_num_inputs, int gpu_index, int gpu_count);

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

@@ -20,15 +20,15 @@ template <typename Torus> struct int_compression {
                   uint32_t gpu_count, int_radix_params compression_params,
                   uint32_t num_radix_blocks, uint32_t lwe_per_glwe,
                   uint32_t storage_log_modulus, bool allocate_gpu_memory,
-                  uint64_t &size_tracker) {
+                  uint64_t *size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     this->compression_params = compression_params;
     this->lwe_per_glwe = lwe_per_glwe;
     this->storage_log_modulus = storage_log_modulus;
     this->body_count = num_radix_blocks;
 
-    uint64_t glwe_accumulator_size = (compression_params.glwe_dimension + 1) *
-                                     compression_params.polynomial_size;
+    Torus glwe_accumulator_size = (compression_params.glwe_dimension + 1) *
+                                  compression_params.polynomial_size;
 
     tmp_lwe = (Torus *)cuda_malloc_with_size_tracking_async(
         num_radix_blocks * (compression_params.small_lwe_dimension + 1) *
@@ -38,7 +38,7 @@ template <typename Torus> struct int_compression {
         lwe_per_glwe * glwe_accumulator_size * sizeof(Torus), streams[0],
         gpu_indexes[0], size_tracker, allocate_gpu_memory);
 
-    size_tracker += scratch_packing_keyswitch_lwe_list_to_glwe_64(
+    *size_tracker += scratch_packing_keyswitch_lwe_list_to_glwe_64(
         streams[0], gpu_indexes[0], &fp_ks_buffer,
         compression_params.small_lwe_dimension,
         compression_params.glwe_dimension, compression_params.polynomial_size,
@@ -76,7 +76,7 @@ template <typename Torus> struct int_decompression {
                     int_radix_params compression_params,
                     uint32_t num_radix_blocks, uint32_t body_count,
                     uint32_t storage_log_modulus, bool allocate_gpu_memory,
-                    uint64_t &size_tracker) {
+                    uint64_t *size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     this->encryption_params = encryption_params;
     this->compression_params = compression_params;
@@ -84,11 +84,11 @@ template <typename Torus> struct int_decompression {
     this->num_radix_blocks = num_radix_blocks;
     this->body_count = body_count;
 
-    uint64_t glwe_accumulator_size = (compression_params.glwe_dimension + 1) *
-                                     compression_params.polynomial_size;
-    uint64_t lwe_accumulator_size = (compression_params.glwe_dimension *
-                                         compression_params.polynomial_size +
-                                     1);
+    Torus glwe_accumulator_size = (compression_params.glwe_dimension + 1) *
+                                  compression_params.polynomial_size;
+    Torus lwe_accumulator_size = (compression_params.glwe_dimension *
+                                      compression_params.polynomial_size +
+                                  1);
     decompression_rescale_lut = new int_radix_lut<Torus>(
         streams, gpu_indexes, gpu_count, encryption_params, 1, num_radix_blocks,
         allocate_gpu_memory, size_tracker);
@@ -124,7 +124,7 @@ template <typename Torus> struct int_decompression {
         encryption_params.carry_modulus, decompression_rescale_f,
         gpu_memory_allocated);
 
-    decompression_rescale_lut->broadcast_lut(streams, gpu_indexes);
+    decompression_rescale_lut->broadcast_lut(streams, gpu_indexes, 0);
   }
   void release(cudaStream_t const *streams, uint32_t const *gpu_indexes,
                uint32_t gpu_count) {

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

@@ -48,34 +48,6 @@ typedef struct {
   uint32_t lwe_dimension;
 } CudaRadixCiphertextFFI;
 
-typedef struct {
-  uint64_t const *chosen_multiplier_has_at_least_one_set;
-  uint64_t const *decomposed_chosen_multiplier;
-
-  uint32_t const num_scalars;
-  uint32_t const active_bits;
-  uint64_t const shift_pre;
-  uint32_t const shift_post;
-  uint32_t const ilog2_chosen_multiplier;
-  uint32_t const chosen_multiplier_num_bits;
-
-  bool const is_chosen_multiplier_zero;
-  bool const is_abs_chosen_multiplier_one;
-  bool const is_chosen_multiplier_negative;
-  bool const is_chosen_multiplier_pow2;
-  bool const chosen_multiplier_has_more_bits_than_numerator;
-  // if signed: test if chosen_multiplier >= 2^{num_bits - 1}
-  bool const is_chosen_multiplier_geq_two_pow_numerator;
-
-  uint32_t const ilog2_divisor;
-
-  bool const is_divisor_zero;
-  bool const is_abs_divisor_one;
-  bool const is_divisor_negative;
-  bool const is_divisor_pow2;
-  bool const divisor_has_more_bits_than_numerator;
-} CudaScalarDivisorFFI;
-
 uint64_t scratch_cuda_apply_univariate_lut_kb_64(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
     int8_t **mem_ptr, void const *input_lut, uint32_t lwe_dimension,
@@ -423,8 +395,7 @@ uint64_t scratch_cuda_integer_radix_partial_sum_ciphertexts_vec_kb_64(
     uint32_t pbs_level, uint32_t pbs_base_log, uint32_t grouping_factor,
     uint32_t num_blocks_in_radix, uint32_t max_num_radix_in_vec,
     uint32_t message_modulus, uint32_t carry_modulus, PBS_TYPE pbs_type,
-    bool reduce_degrees_for_single_carry_propagation, bool allocate_gpu_memory,
-    bool allocate_ms_array);
+    bool allocate_gpu_memory, bool allocate_ms_array);
 
 void cuda_integer_radix_partial_sum_ciphertexts_vec_kb_64(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
@@ -443,8 +414,7 @@ uint64_t scratch_cuda_integer_scalar_mul_kb_64(
     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, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, uint32_t num_scalar_bits, bool allocate_gpu_memory,
-    bool allocate_ms_array);
+    PBS_TYPE pbs_type, bool allocate_gpu_memory, bool allocate_ms_array);
 
 void cuda_scalar_multiplication_integer_radix_ciphertext_64_inplace(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
@@ -568,189 +538,5 @@ void cleanup_cuda_integer_is_at_least_one_comparisons_block_true(
     void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
     int8_t **mem_ptr_void);
 
-void extend_radix_with_trivial_zero_blocks_msb_64(
-    CudaRadixCiphertextFFI *output, CudaRadixCiphertextFFI const *input,
-    void *const *streams, uint32_t const *gpu_indexes);
-
-void trim_radix_blocks_lsb_64(CudaRadixCiphertextFFI *output,
-                              CudaRadixCiphertextFFI const *input,
-                              void *const *streams,
-                              uint32_t const *gpu_indexes);
-
-uint64_t scratch_cuda_apply_noise_squashing_kb(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t lwe_dimension, uint32_t glwe_dimension,
-    uint32_t polynomial_size, uint32_t input_glwe_dimension,
-    uint32_t input_polynomial_size, 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_radix_blocks, uint32_t num_original_blocks,
-    uint32_t message_modulus, uint32_t carry_modulus, PBS_TYPE pbs_type,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-void cuda_apply_noise_squashing_kb(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *output_radix_lwe,
-    CudaRadixCiphertextFFI const *input_radix_lwe, int8_t *mem_ptr,
-    void *const *ksks,
-    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key,
-    void *const *bsks);
-
-void cleanup_cuda_apply_noise_squashing_kb(void *const *streams,
-                                           uint32_t const *gpu_indexes,
-                                           uint32_t gpu_count,
-                                           int8_t **mem_ptr_void);
-
-uint64_t scratch_cuda_sub_and_propagate_single_carry_kb_64_inplace(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, uint32_t polynomial_size,
-    uint32_t big_lwe_dimension, 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,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-void cuda_sub_and_propagate_single_carry_kb_64_inplace(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *lhs_array, const CudaRadixCiphertextFFI *rhs_array,
-    CudaRadixCiphertextFFI *carry_out, const CudaRadixCiphertextFFI *carry_in,
-    int8_t *mem_ptr, void *const *bsks, void *const *ksks,
-    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key,
-    uint32_t requested_flag, uint32_t uses_carry);
-
-void cleanup_cuda_sub_and_propagate_single_carry(void *const *streams,
-                                                 uint32_t const *gpu_indexes,
-                                                 uint32_t gpu_count,
-                                                 int8_t **mem_ptr_void);
-
-uint64_t scratch_cuda_integer_unsigned_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-void cuda_integer_unsigned_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *numerator_ct, int8_t *mem_ptr, void *const *bsks,
-    void *const *ksks,
-    const CudaModulusSwitchNoiseReductionKeyFFI *ms_noise_reduction_key,
-    const CudaScalarDivisorFFI *scalar_divisor_ffi);
-
-void cleanup_cuda_integer_unsigned_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr_void);
-
-uint64_t scratch_cuda_extend_radix_with_sign_msb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, uint32_t num_additional_blocks,
-    uint32_t message_modulus, uint32_t carry_modulus, PBS_TYPE pbs_type,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-void cuda_extend_radix_with_sign_msb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *output, CudaRadixCiphertextFFI const *input,
-    int8_t *mem_ptr, uint32_t num_additional_blocks, void *const *bsks,
-    void *const *ksks,
-    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key);
-
-void cleanup_cuda_extend_radix_with_sign_msb_64(void *const *streams,
-                                                uint32_t const *gpu_indexes,
-                                                uint32_t gpu_count,
-                                                int8_t **mem_ptr_void);
-
-uint64_t scratch_cuda_integer_signed_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-void cuda_integer_signed_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *numerator_ct, int8_t *mem_ptr, void *const *bsks,
-    void *const *ksks,
-    const CudaModulusSwitchNoiseReductionKeyFFI *ms_noise_reduction_key,
-    const CudaScalarDivisorFFI *scalar_divisor_ffi, uint32_t numerator_bits);
-
-void cleanup_cuda_integer_signed_scalar_div_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr_void);
-
-uint64_t scratch_integer_unsigned_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    uint32_t const active_bits_divisor, bool allocate_gpu_memory,
-    bool allocate_ms_array);
-
-void cuda_integer_unsigned_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *quotient_ct, CudaRadixCiphertextFFI *remainder_ct,
-    int8_t *mem_ptr, void *const *bsks, void *const *ksks,
-    const CudaModulusSwitchNoiseReductionKeyFFI *ms_noise_reduction_key,
-    const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    uint64_t const *divisor_has_at_least_one_set,
-    uint64_t const *decomposed_divisor, uint32_t const num_scalars_divisor,
-    void const *clear_blocks, void const *h_clear_blocks,
-    uint32_t num_clear_blocks);
-
-void cleanup_cuda_integer_unsigned_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr_void);
-
-uint64_t scratch_integer_signed_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, uint32_t message_modulus, uint32_t carry_modulus,
-    PBS_TYPE pbs_type, const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    uint32_t const active_bits_divisor, bool allocate_gpu_memory,
-    bool allocate_ms_array);
-
-void cuda_integer_signed_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *quotient_ct, CudaRadixCiphertextFFI *remainder_ct,
-    int8_t *mem_ptr, void *const *bsks, void *const *ksks,
-    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key,
-    const CudaScalarDivisorFFI *scalar_divisor_ffi,
-    uint64_t const *divisor_has_at_least_one_set,
-    uint64_t const *decomposed_divisor, uint32_t const num_scalars_divisor,
-    uint32_t numerator_bits);
-
-void cleanup_cuda_integer_signed_scalar_div_rem_radix_kb_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr_void);
-
-uint64_t scratch_cuda_integer_grouped_oprf_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    int8_t **mem_ptr, uint32_t glwe_dimension, 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, bool allocate_ms_array);
-
-void cuda_integer_grouped_oprf_async_64(
-    void *const *streams, uint32_t const *gpu_indexes, uint32_t gpu_count,
-    CudaRadixCiphertextFFI *radix_lwe_out, const void *seeded_lwe_input,
-    uint32_t num_blocks_to_process, int8_t *mem, void *const *bsks,
-    CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key);
-
-void cleanup_cuda_integer_grouped_oprf_64(void *const *streams,
-                                          uint32_t const *gpu_indexes,
-                                          uint32_t gpu_count,
-                                          int8_t **mem_ptr_void);
 } // extern C
 #endif // CUDA_INTEGER_H

backends/tfhe-cuda-backend/cuda/include/pbs/pbs_128_utilities.h

@@ -1,13 +0,0 @@
-#ifndef CUDA_BOOTSTRAP_128_H
-#define CUDA_BOOTSTRAP_128_H
-
-#include "pbs_enums.h"
-#include <stdint.h>
-
-uint64_t scratch_cuda_programmable_bootstrap_128_vector_64(
-    void *stream, uint32_t gpu_index, int8_t **pbs_buffer,
-    uint32_t lwe_dimension, uint32_t glwe_dimension, uint32_t polynomial_size,
-    uint32_t level_count, uint32_t input_lwe_ciphertext_count,
-    bool allocate_gpu_memory, bool allocate_ms_array);
-
-#endif // CUDA_BOOTSTRAP_128_H

backends/tfhe-cuda-backend/cuda/include/pbs/pbs_multibit_utilities.h

@@ -66,9 +66,6 @@ void cuda_multi_bit_programmable_bootstrap_lwe_ciphertext_vector(
     uint32_t num_many_lut, uint32_t lut_stride);
 
 template <typename Torus>
-uint64_t get_buffer_size_full_sm_multibit_programmable_bootstrap_128_keybundle(
-    uint32_t polynomial_size);
-template <typename Torus>
 uint64_t get_buffer_size_full_sm_multibit_programmable_bootstrap_keybundle(
     uint32_t polynomial_size);
 template <typename Torus>
@@ -98,12 +95,8 @@ uint64_t get_buffer_size_full_sm_tbc_multibit_programmable_bootstrap(
 
 template <typename Torus, class params>
 uint32_t get_lwe_chunk_size(uint32_t gpu_index, uint32_t max_num_pbs,
-                            uint32_t polynomial_size,
-                            uint64_t full_sm_keybundle);
-template <typename Torus, class params>
-uint32_t get_lwe_chunk_size_128(uint32_t gpu_index, uint32_t max_num_pbs,
-                                uint32_t polynomial_size,
-                                uint64_t full_sm_keybundle);
+                            uint32_t polynomial_size);
+
 template <typename Torus> struct pbs_buffer<Torus, PBS_TYPE::MULTI_BIT> {
   int8_t *d_mem_keybundle = NULL;
   int8_t *d_mem_acc_step_one = NULL;
@@ -122,7 +115,7 @@ template <typename Torus> struct pbs_buffer<Torus, PBS_TYPE::MULTI_BIT> {
              uint32_t polynomial_size, uint32_t level_count,
              uint32_t input_lwe_ciphertext_count, uint32_t lwe_chunk_size,
              PBS_VARIANT pbs_variant, bool allocate_gpu_memory,
-             uint64_t &size_tracker) {
+             uint64_t *size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     cuda_set_device(gpu_index);
 
@@ -288,146 +281,4 @@ template <typename Torus> struct pbs_buffer<Torus, PBS_TYPE::MULTI_BIT> {
   }
 };
 
-template <typename InputTorus>
-struct pbs_buffer_128<InputTorus, PBS_TYPE::MULTI_BIT> {
-  int8_t *d_mem_keybundle = NULL;
-  int8_t *d_mem_acc_step_one = NULL;
-  int8_t *d_mem_acc_step_two = NULL;
-  int8_t *d_mem_acc_cg = NULL;
-  int8_t *d_mem_acc_tbc = NULL;
-  uint32_t lwe_chunk_size;
-  double *keybundle_fft;
-  __uint128_t *global_accumulator;
-  double *global_join_buffer;
-
-  PBS_VARIANT pbs_variant;
-  bool gpu_memory_allocated;
-
-  pbs_buffer_128(cudaStream_t stream, uint32_t gpu_index,
-                 uint32_t glwe_dimension, uint32_t polynomial_size,
-                 uint32_t level_count, uint32_t input_lwe_ciphertext_count,
-                 uint32_t lwe_chunk_size, PBS_VARIANT pbs_variant,
-                 bool allocate_gpu_memory, uint64_t &size_tracker) {
-    gpu_memory_allocated = allocate_gpu_memory;
-    cuda_set_device(gpu_index);
-
-    this->pbs_variant = pbs_variant;
-    this->lwe_chunk_size = lwe_chunk_size;
-    auto max_shared_memory = cuda_get_max_shared_memory(gpu_index);
-
-    // default
-    uint64_t full_sm_keybundle =
-        get_buffer_size_full_sm_multibit_programmable_bootstrap_128_keybundle<
-            __uint128_t>(polynomial_size);
-    uint64_t full_sm_accumulate_step_one =
-        get_buffer_size_full_sm_multibit_programmable_bootstrap_step_one<
-            __uint128_t>(polynomial_size);
-    uint64_t full_sm_accumulate_step_two =
-        get_buffer_size_full_sm_multibit_programmable_bootstrap_step_two<
-            __uint128_t>(polynomial_size);
-    uint64_t partial_sm_accumulate_step_one =
-        get_buffer_size_partial_sm_multibit_programmable_bootstrap_step_one<
-            __uint128_t>(polynomial_size);
-    // cg
-    uint64_t full_sm_cg_accumulate =
-        get_buffer_size_full_sm_cg_multibit_programmable_bootstrap<__uint128_t>(
-            polynomial_size);
-    uint64_t partial_sm_cg_accumulate =
-        get_buffer_size_partial_sm_cg_multibit_programmable_bootstrap<
-            __uint128_t>(polynomial_size);
-
-    auto num_blocks_keybundle = input_lwe_ciphertext_count * lwe_chunk_size *
-                                (glwe_dimension + 1) * (glwe_dimension + 1) *
-                                level_count;
-    auto num_blocks_acc_step_one =
-        level_count * (glwe_dimension + 1) * input_lwe_ciphertext_count;
-    auto num_blocks_acc_step_two =
-        input_lwe_ciphertext_count * (glwe_dimension + 1);
-    auto num_blocks_acc_cg =
-        level_count * (glwe_dimension + 1) * input_lwe_ciphertext_count;
-
-    // Keybundle
-    if (max_shared_memory < full_sm_keybundle)
-      d_mem_keybundle = (int8_t *)cuda_malloc_with_size_tracking_async(
-          num_blocks_keybundle * full_sm_keybundle, stream, gpu_index,
-          size_tracker, allocate_gpu_memory);
-
-    switch (pbs_variant) {
-    case PBS_VARIANT::CG:
-      // Accumulator CG
-      if (max_shared_memory < partial_sm_cg_accumulate)
-        d_mem_acc_cg = (int8_t *)cuda_malloc_with_size_tracking_async(
-            num_blocks_acc_cg * full_sm_cg_accumulate, stream, gpu_index,
-            size_tracker, allocate_gpu_memory);
-      else if (max_shared_memory < full_sm_cg_accumulate)
-        d_mem_acc_cg = (int8_t *)cuda_malloc_with_size_tracking_async(
-            num_blocks_acc_cg * partial_sm_cg_accumulate, stream, gpu_index,
-            size_tracker, allocate_gpu_memory);
-      break;
-    case PBS_VARIANT::DEFAULT:
-      // Accumulator step one
-      if (max_shared_memory < partial_sm_accumulate_step_one)
-        d_mem_acc_step_one = (int8_t *)cuda_malloc_with_size_tracking_async(
-            num_blocks_acc_step_one * full_sm_accumulate_step_one, stream,
-            gpu_index, size_tracker, allocate_gpu_memory);
-      else if (max_shared_memory < full_sm_accumulate_step_one)
-        d_mem_acc_step_one = (int8_t *)cuda_malloc_with_size_tracking_async(
-            num_blocks_acc_step_one * partial_sm_accumulate_step_one, stream,
-            gpu_index, size_tracker, allocate_gpu_memory);
-
-      // Accumulator step two
-      if (max_shared_memory < full_sm_accumulate_step_two)
-        d_mem_acc_step_two = (int8_t *)cuda_malloc_with_size_tracking_async(
-            num_blocks_acc_step_two * full_sm_accumulate_step_two, stream,
-            gpu_index, size_tracker, allocate_gpu_memory);
-      break;
-    default:
-      PANIC("Cuda error (PBS): unsupported implementation variant.")
-    }
-
-    keybundle_fft = (double *)cuda_malloc_with_size_tracking_async(
-        num_blocks_keybundle * (polynomial_size / 2) * 4 * sizeof(double),
-        stream, gpu_index, size_tracker, allocate_gpu_memory);
-    global_accumulator = (__uint128_t *)cuda_malloc_with_size_tracking_async(
-        input_lwe_ciphertext_count * (glwe_dimension + 1) * polynomial_size *
-            sizeof(__uint128_t),
-        stream, gpu_index, size_tracker, allocate_gpu_memory);
-    global_join_buffer = (double *)cuda_malloc_with_size_tracking_async(
-        level_count * (glwe_dimension + 1) * input_lwe_ciphertext_count *
-            (polynomial_size / 2) * 4 * sizeof(double),
-        stream, gpu_index, size_tracker, allocate_gpu_memory);
-  }
-
-  void release(cudaStream_t stream, uint32_t gpu_index) {
-
-    if (d_mem_keybundle)
-      cuda_drop_with_size_tracking_async(d_mem_keybundle, stream, gpu_index,
-                                         gpu_memory_allocated);
-    switch (pbs_variant) {
-    case DEFAULT:
-      if (d_mem_acc_step_one)
-        cuda_drop_with_size_tracking_async(d_mem_acc_step_one, stream,
-                                           gpu_index, gpu_memory_allocated);
-      if (d_mem_acc_step_two)
-        cuda_drop_with_size_tracking_async(d_mem_acc_step_two, stream,
-                                           gpu_index, gpu_memory_allocated);
-      break;
-    case CG:
-      if (d_mem_acc_cg)
-        cuda_drop_with_size_tracking_async(d_mem_acc_cg, stream, gpu_index,
-                                           gpu_memory_allocated);
-      break;
-    default:
-      PANIC("Cuda error (PBS): unsupported implementation variant.")
-    }
-
-    cuda_drop_with_size_tracking_async(keybundle_fft, stream, gpu_index,
-                                       gpu_memory_allocated);
-    cuda_drop_with_size_tracking_async(global_accumulator, stream, gpu_index,
-                                       gpu_memory_allocated);
-    cuda_drop_with_size_tracking_async(global_join_buffer, stream, gpu_index,
-                                       gpu_memory_allocated);
-  }
-};
-
 #endif // CUDA_MULTI_BIT_UTILITIES_H

backends/tfhe-cuda-backend/cuda/include/pbs/pbs_utilities.h

@@ -90,7 +90,7 @@ template <typename Torus> struct pbs_buffer<Torus, PBS_TYPE::CLASSICAL> {
              uint32_t glwe_dimension, uint32_t polynomial_size,
              uint32_t level_count, uint32_t input_lwe_ciphertext_count,
              PBS_VARIANT pbs_variant, bool allocate_gpu_memory,
-             bool allocate_ms_array, uint64_t &size_tracker) {
+             bool allocate_ms_array, uint64_t *size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     cuda_set_device(gpu_index);
     this->uses_noise_reduction = allocate_ms_array;
@@ -240,16 +240,14 @@ template <typename Torus> struct pbs_buffer<Torus, PBS_TYPE::CLASSICAL> {
   }
 };
 
-template <typename Torus, PBS_TYPE pbs_type> struct pbs_buffer_128;
+template <PBS_TYPE pbs_type> struct pbs_buffer_128;
 
-template <typename InputTorus>
-struct pbs_buffer_128<InputTorus, PBS_TYPE::CLASSICAL> {
+template <> struct pbs_buffer_128<PBS_TYPE::CLASSICAL> {
   int8_t *d_mem;
 
   __uint128_t *global_accumulator;
   double *global_join_buffer;
-  InputTorus *temp_lwe_array_in;
-  uint64_t *trivial_indexes;
+  __uint128_t *temp_lwe_array_in;
 
   PBS_VARIANT pbs_variant;
   bool uses_noise_reduction;
@@ -260,30 +258,16 @@ struct pbs_buffer_128<InputTorus, PBS_TYPE::CLASSICAL> {
                  uint32_t polynomial_size, uint32_t level_count,
                  uint32_t input_lwe_ciphertext_count, PBS_VARIANT pbs_variant,
                  bool allocate_gpu_memory, bool allocate_ms_array,
-                 uint64_t &size_tracker) {
+                 uint64_t *size_tracker) {
     gpu_memory_allocated = allocate_gpu_memory;
     cuda_set_device(gpu_index);
     this->pbs_variant = pbs_variant;
     this->uses_noise_reduction = allocate_ms_array;
-    if (allocate_ms_array) {
-      this->temp_lwe_array_in = (InputTorus *)cuda_malloc_async(
-          (lwe_dimension + 1) * input_lwe_ciphertext_count * sizeof(InputTorus),
-          stream, gpu_index);
-      this->trivial_indexes = (uint64_t *)cuda_malloc_with_size_tracking_async(
-          input_lwe_ciphertext_count * sizeof(uint64_t), stream, gpu_index,
-          size_tracker, allocate_ms_array);
-      uint64_t *h_trivial_indexes = new uint64_t[input_lwe_ciphertext_count];
-      for (uint32_t i = 0; i < input_lwe_ciphertext_count; i++)
-        h_trivial_indexes[i] = i;
-
-      cuda_memcpy_with_size_tracking_async_to_gpu(
-          trivial_indexes, h_trivial_indexes,
-          input_lwe_ciphertext_count * sizeof(uint64_t), stream, gpu_index,
-          allocate_gpu_memory);
-
-      cuda_synchronize_stream(stream, gpu_index);
-      delete[] h_trivial_indexes;
-    }
+    this->temp_lwe_array_in =
+        (__uint128_t *)cuda_malloc_with_size_tracking_async(
+            (lwe_dimension + 1) * input_lwe_ciphertext_count *
+                sizeof(__uint128_t),
+            stream, gpu_index, size_tracker, allocate_ms_array);
     auto max_shared_memory = cuda_get_max_shared_memory(gpu_index);
     size_t global_join_buffer_size = (glwe_dimension + 1) * level_count *
                                      input_lwe_ciphertext_count *
@@ -420,12 +404,9 @@ struct pbs_buffer_128<InputTorus, PBS_TYPE::CLASSICAL> {
       cuda_drop_with_size_tracking_async(global_accumulator, stream, gpu_index,
                                          gpu_memory_allocated);
 
-    if (uses_noise_reduction) {
+    if (uses_noise_reduction)
       cuda_drop_with_size_tracking_async(temp_lwe_array_in, stream, gpu_index,
                                          gpu_memory_allocated);
-      cuda_drop_with_size_tracking_async(trivial_indexes, stream, gpu_index,
-                                         gpu_memory_allocated);
-    }
   }
 };
 
@@ -521,12 +502,7 @@ template <typename Torus>
 bool has_support_to_cuda_programmable_bootstrap_tbc(uint32_t num_samples,
                                                     uint32_t glwe_dimension,
                                                     uint32_t polynomial_size,
-                                                    uint32_t level_count,
-                                                    uint32_t max_shared_memory);
-
-bool has_support_to_cuda_programmable_bootstrap_128_cg(
-    uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t level_count,
-    uint32_t num_samples, uint32_t max_shared_memory);
+                                                    uint32_t level_count);
 
 #ifdef __CUDACC__
 __device__ inline int get_start_ith_ggsw(int i, uint32_t polynomial_size,
@@ -544,11 +520,6 @@ __device__ T *get_ith_mask_kth_block(T *ptr, int i, int k, int level,
                                      uint32_t polynomial_size,
                                      int glwe_dimension, uint32_t level_count);
 
-template <typename T, uint32_t polynomial_size, uint32_t glwe_dimension,
-          uint32_t level_count, uint32_t level_id>
-__device__ const T *get_ith_mask_kth_block_2_2_params(const T *ptr,
-                                                      int iteration, int k);
-
 template <typename T>
 __device__ T *get_ith_body_kth_block(T *ptr, int i, int k, int level,
                                      uint32_t polynomial_size,

backends/tfhe-cuda-backend/cuda/include/pbs/programmable_bootstrap.h

@@ -100,7 +100,7 @@ void cuda_programmable_bootstrap_lwe_ciphertext_vector_128(
     void const *lut_vector, void const *lwe_array_in,
     void const *bootstrapping_key,
     CudaModulusSwitchNoiseReductionKeyFFI const *ms_noise_reduction_key,
-    void const *ms_noise_reduction_ptr, int8_t *buffer, uint32_t lwe_dimension,
+    void *ms_noise_reduction_ptr, int8_t *buffer, uint32_t lwe_dimension,
     uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t base_log,
     uint32_t level_count, uint32_t num_samples);
 

backends/tfhe-cuda-backend/cuda/include/pbs/programmable_bootstrap_multibit.h

@@ -15,11 +15,6 @@ void cuda_convert_lwe_multi_bit_programmable_bootstrap_key_64(
     uint32_t input_lwe_dim, uint32_t glwe_dim, uint32_t level_count,
     uint32_t polynomial_size, uint32_t grouping_factor);
 
-void cuda_convert_lwe_multi_bit_programmable_bootstrap_key_128(
-    void *stream, uint32_t gpu_index, void *dest, void const *src,
-    uint32_t input_lwe_dim, uint32_t glwe_dim, uint32_t level_count,
-    uint32_t polynomial_size, uint32_t grouping_factor);
-
 uint64_t scratch_cuda_multi_bit_programmable_bootstrap_64(
     void *stream, uint32_t gpu_index, int8_t **pbs_buffer,
     uint32_t glwe_dimension, uint32_t polynomial_size, uint32_t level_count,
@@ -38,25 +33,6 @@ void cuda_multi_bit_programmable_bootstrap_lwe_ciphertext_vector_64(
 void cleanup_cuda_multi_bit_programmable_bootstrap(void *stream,
                                                    uint32_t gpu_index,
                                                    int8_t **pbs_buffer);
-
-uint64_t scratch_cuda_multi_bit_programmable_bootstrap_128_vector_64(
-    void *stream, uint32_t gpu_index, int8_t **buffer, uint32_t glwe_dimension,
-    uint32_t polynomial_size, uint32_t level_count,
-    uint32_t input_lwe_ciphertext_count, bool allocate_gpu_memory);
-
-void cuda_multi_bit_programmable_bootstrap_lwe_ciphertext_vector_128(
-    void *stream, uint32_t gpu_index, void *lwe_array_out,
-    void const *lwe_output_indexes, void const *lut_vector,
-    void const *lut_vector_indexes, void const *lwe_array_in,
-    void const *lwe_input_indexes, void const *bootstrapping_key,
-    int8_t *mem_ptr, uint32_t lwe_dimension, uint32_t glwe_dimension,
-    uint32_t polynomial_size, uint32_t grouping_factor, uint32_t base_log,
-    uint32_t level_count, uint32_t num_samples, uint32_t num_many_lut,
-    uint32_t lut_stride);
-
-void cleanup_cuda_multi_bit_programmable_bootstrap_128(void *stream,
-                                                       const uint32_t gpu_index,
-                                                       int8_t **buffer);
 }
 
 #endif // CUDA_MULTI_BIT_H

backends/tfhe-cuda-backend/cuda/include/zk/zk_utilities.h

@@ -27,7 +27,7 @@ template <typename Torus> struct zk_expand_mem {
                 int_radix_params casting_params, KS_TYPE casting_key_type,
                 const uint32_t *num_lwes_per_compact_list,
                 const bool *is_boolean_array, uint32_t num_compact_lists,
-                bool allocate_gpu_memory, uint64_t &size_tracker)
+                bool allocate_gpu_memory, uint64_t *size_tracker)
       : computing_params(computing_params), casting_params(casting_params),
         num_compact_lists(num_compact_lists),
         casting_key_type(casting_key_type) {
@@ -112,15 +112,15 @@ template <typename Torus> struct zk_expand_mem {
 
     // Hint for future readers: if message_modulus == 4 then
     // packed_messages_per_lwe becomes 2
-    auto num_packed_msgs = log2_int(params.message_modulus);
+    auto packed_messages_per_lwe = log2_int(params.message_modulus);
 
     // Adjust indexes to permute the output and access the correct LUT
     auto h_indexes_in = static_cast<Torus *>(
-        malloc(num_packed_msgs * num_lwes * sizeof(Torus)));
+        malloc(packed_messages_per_lwe * num_lwes * sizeof(Torus)));
     auto h_indexes_out = static_cast<Torus *>(
-        malloc(num_packed_msgs * num_lwes * sizeof(Torus)));
+        malloc(packed_messages_per_lwe * num_lwes * sizeof(Torus)));
     auto h_lut_indexes = static_cast<Torus *>(
-        malloc(num_packed_msgs * num_lwes * sizeof(Torus)));
+        malloc(packed_messages_per_lwe * num_lwes * sizeof(Torus)));
     auto h_body_id_per_compact_list =
         static_cast<uint32_t *>(malloc(num_lwes * sizeof(uint32_t)));
     auto h_lwe_compact_input_indexes =
@@ -138,10 +138,6 @@ template <typename Torus> struct zk_expand_mem {
     auto compact_list_id = 0;
     auto idx = 0;
     auto count = 0;
-    // During flattening, all num_lwes LWEs from all compact lists are stored
-    // sequentially on a Torus array. h_lwe_compact_input_indexes stores the
-    // index of the first LWE related to the compact list that contains the i-th
-    // LWE
     for (int i = 0; i < num_lwes; i++) {
       h_lwe_compact_input_indexes[i] = idx;
       count++;
@@ -152,8 +148,6 @@ template <typename Torus> struct zk_expand_mem {
       }
     }
 
-    // Stores the index of the i-th LWE (within each compact list) related to
-    // the k-th compact list.
     auto offset = 0;
     for (int k = 0; k < num_compact_lists; k++) {
       auto num_lwes_in_kth_compact_list = num_lwes_per_compact_list[k];
@@ -165,75 +159,46 @@ template <typename Torus> struct zk_expand_mem {
       offset += num_lwes_in_kth_compact_list;
     }
 
-    /*
-     * Each LWE contains encrypted data in both carry and message spaces
-     * that needs to be extracted.
-     *
-     * The loop processes each compact list (k) and for each LWE within that
-     * list:
-     * 1. Sets input indexes to read each LWE twice (for carry and message
-     * extraction)
-     * 2. Creates output indexes to properly reorder the results
-     * 3. Selects appropriate LUT index based on whether boolean sanitization is
-     * needed
-     *
-     * We want the output to have always first the content of the message part
-     * and then the content of the carry part of each LWE.
-     *
-     * i.e. msg_extract(LWE_0), carry_extract(LWE_0), msg_extract(LWE_1),
-     * carry_extract(LWE_1), ...
-     *
-     * Aiming that behavior, with 4 LWEs we would have:
-     *
-     * // Each LWE is processed twice
-     * h_indexes_in   = {0, 1, 2, 3, 0, 1, 2, 3}
-     *
-     * // First 4 use message LUT, last 4 use carry LUT
-     * h_lut_indexes  = {0, 0, 0, 0, 1, 1, 1, 1}
-     *
-     * // Reorders output so message and carry for each LWE appear together
-     * h_indexes_out  = {0, 2, 4, 6, 1, 3, 5, 7}
-     *
-     * If an LWE contains a boolean value, its LUT index is shifted by
-     * num_packed_msgs to use the sanitization LUT (which ensures output is
-     * exactly 0 or 1).
-     */
     offset = 0;
     for (int k = 0; k < num_compact_lists; k++) {
-      auto num_lwes_in_kth = num_lwes_per_compact_list[k];
-      for (int i = 0; i < num_packed_msgs * num_lwes_in_kth; i++) {
-        auto lwe_index = i + num_packed_msgs * offset;
-        auto lwe_index_in_list = i % num_lwes_in_kth;
-        h_indexes_in[lwe_index] = lwe_index_in_list + offset;
-        h_indexes_out[lwe_index] =
-            num_packed_msgs * h_indexes_in[lwe_index] + i / num_lwes_in_kth;
+      auto num_lwes_in_kth_compact_list = num_lwes_per_compact_list[k];
+      for (int i = 0;
+           i < packed_messages_per_lwe * num_lwes_in_kth_compact_list; i++) {
+        Torus j = i % num_lwes_in_kth_compact_list;
+        h_indexes_in[i + packed_messages_per_lwe * offset] = j + offset;
+        h_indexes_out[i + packed_messages_per_lwe * offset] =
+            packed_messages_per_lwe * (j + offset) +
+            (i / num_lwes_in_kth_compact_list);
         // If the input relates to a boolean, shift the LUT so the correct one
         // with sanitization is used
-        auto boolean_offset =
-            is_boolean_array[h_indexes_out[lwe_index]] ? num_packed_msgs : 0;
-        h_lut_indexes[lwe_index] = i / num_lwes_in_kth + boolean_offset;
+        h_lut_indexes[i + packed_messages_per_lwe * offset] =
+            (is_boolean_array[h_indexes_out[i +
+                                            packed_messages_per_lwe * offset]]
+                 ? packed_messages_per_lwe
+                 : 0) +
+            i / num_lwes_in_kth_compact_list;
       }
-      offset += num_lwes_in_kth;
+      offset += num_lwes_in_kth_compact_list;
     }
 
     message_and_carry_extract_luts->set_lwe_indexes(
         streams[0], gpu_indexes[0], h_indexes_in, h_indexes_out);
     auto lut_indexes = message_and_carry_extract_luts->get_lut_indexes(0, 0);
+    message_and_carry_extract_luts->broadcast_lut(streams, gpu_indexes, 0);
 
     cuda_memcpy_with_size_tracking_async_to_gpu(
         d_lwe_compact_input_indexes, h_lwe_compact_input_indexes,
         num_lwes * sizeof(uint32_t), streams[0], gpu_indexes[0],
         allocate_gpu_memory);
     cuda_memcpy_with_size_tracking_async_to_gpu(
-        lut_indexes, h_lut_indexes, num_packed_msgs * num_lwes * sizeof(Torus),
-        streams[0], gpu_indexes[0], allocate_gpu_memory);
+        lut_indexes, h_lut_indexes,
+        packed_messages_per_lwe * num_lwes * sizeof(Torus), streams[0],
+        gpu_indexes[0], allocate_gpu_memory);
     cuda_memcpy_with_size_tracking_async_to_gpu(
         d_body_id_per_compact_list, h_body_id_per_compact_list,
         num_lwes * sizeof(uint32_t), streams[0], gpu_indexes[0],
         allocate_gpu_memory);
 
-    message_and_carry_extract_luts->broadcast_lut(streams, gpu_indexes);
-
     // The expanded LWEs will always be on the casting key format
     tmp_expanded_lwes = (Torus *)cuda_malloc_with_size_tracking_async(
         num_lwes * (casting_params.big_lwe_dimension + 1) * sizeof(Torus),

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

@@ -1,6 +1,5 @@
 file(GLOB_RECURSE SOURCES "*.cu")
-add_library(tfhe_cuda_backend STATIC ${SOURCES} pbs/programmable_bootstrap_multibit_128.cuh
-                                     pbs/programmable_bootstrap_multibit_128.cu)
+add_library(tfhe_cuda_backend STATIC ${SOURCES})
 set_target_properties(tfhe_cuda_backend PROPERTIES CUDA_SEPARABLE_COMPILATION ON CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 target_link_libraries(tfhe_cuda_backend PUBLIC cudart OpenMP::OpenMP_CXX)
 target_include_directories(tfhe_cuda_backend PRIVATE .)

backends/tfhe-cuda-backend/cuda/src/crypto/ciphertext.cu

@@ -84,19 +84,15 @@ void cuda_modulus_switch_inplace_64(void *stream, uint32_t gpu_index,
       static_cast<uint64_t *>(lwe_array_out), size, log_modulus);
 }
 
-// This end point is used only for testing purposes
-// its output always follows trivial ordering
 void cuda_improve_noise_modulus_switch_64(
     void *stream, uint32_t gpu_index, void *lwe_array_out,
-    void const *lwe_array_in, void const *lwe_array_indexes,
-    void const *encrypted_zeros, uint32_t lwe_size, uint32_t num_lwes,
-    uint32_t num_zeros, double input_variance, double r_sigma, double bound,
-    uint32_t log_modulus) {
+    void const *lwe_array_in, void const *encrypted_zeros, uint32_t lwe_size,
+    uint32_t num_lwes, uint32_t num_zeros, double input_variance,
+    double r_sigma, double bound, uint32_t log_modulus) {
   host_improve_noise_modulus_switch<uint64_t>(
       static_cast<cudaStream_t>(stream), gpu_index,
       static_cast<uint64_t *>(lwe_array_out),
       static_cast<uint64_t const *>(lwe_array_in),
-      static_cast<uint64_t const *>(lwe_array_indexes),
       static_cast<const uint64_t *>(encrypted_zeros), lwe_size, num_lwes,
       num_zeros, input_variance, r_sigma, bound, log_modulus);
 }

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

@@ -1,5 +1,5 @@
 #ifndef CNCRT_CRYPTO_CUH
-#define CNCRT_CRYPTO_CUH
+#define CNCRT_CRPYTO_CUH
 
 #include "crypto/torus.cuh"
 #include "device.h"
@@ -137,34 +137,6 @@ public:
   }
 };
 
-// Performs the decomposition for 2_2 params, assumes level_count = 1
-// this specialized version it is needed if we plan to keep everything in regs
-template <typename T, class params, uint32_t base_log>
-__device__ void decompose_and_compress_level_2_2_params(double2 *result,
-                                                        T *state) {
-  constexpr T mask_mod_b = (1ll << base_log) - 1ll;
-  for (int i = 0; i < params::opt / 2; i++) {
-    auto input1 = state[i];
-    auto input2 = state[i + params::opt / 2];
-    T res_re = input1 & mask_mod_b;
-    T res_im = input2 & mask_mod_b;
-
-    input1 >>= base_log; // Update state
-    input2 >>= base_log; // Update state
-
-    T carry_re = ((res_re - 1ll) | input1) & res_re;
-    T carry_im = ((res_im - 1ll) | input2) & res_im;
-    carry_re >>= (base_log - 1);
-    carry_im >>= (base_log - 1);
-
-    res_re -= carry_re << base_log;
-    res_im -= carry_im << base_log;
-
-    typecast_torus_to_double(res_re, result[i].x);
-    typecast_torus_to_double(res_im, result[i].y);
-  }
-}
-
 template <typename Torus>
 __device__ Torus decompose_one(Torus &state, Torus mask_mod_b, int base_log) {
   Torus res = state & mask_mod_b;
@@ -176,4 +148,4 @@ __device__ Torus decompose_one(Torus &state, Torus mask_mod_b, int base_log) {
   return res;
 }
 
-#endif // CNCRT_CRYPTO_CUH
+#endif // CNCRT_CRPYTO_H

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

@@ -38,16 +38,6 @@ __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.
 // 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
-// of accumulating subtractions using -=, we accumulate additions using += in
-// the local_lwe_out. This seems to work better cause profits madd ops and save
-// some regs. For this to work, we need to negate the input
-// lwe_array_in[lwe_dimension_in], and negate back the output at the end to get
-// the correct results. Additionally, we split the calculation of the ksk offset
-// in two parts, a constant part is calculated before the loop, and a variable
-// part is calculated inside the loop. This seems to help with the register
-// pressure as well.
 template <typename Torus>
 __global__ void
 keyswitch(Torus *lwe_array_out, const Torus *__restrict__ lwe_output_indexes,
@@ -70,7 +60,7 @@ keyswitch(Torus *lwe_array_out, const Torus *__restrict__ lwe_output_indexes,
         lwe_array_in, lwe_input_indexes[blockIdx.x], lwe_dimension_in + 1);
 
     if (tid == lwe_dimension_out && threadIdx.y == 0) {
-      local_lwe_out = -block_lwe_array_in[lwe_dimension_in];
+      local_lwe_out = block_lwe_array_in[lwe_dimension_in];
     }
     const Torus mask_mod_b = (1ll << base_log) - 1ll;
 
@@ -83,12 +73,12 @@ keyswitch(Torus *lwe_array_out, const Torus *__restrict__ lwe_output_indexes,
     for (int i = start_i; i < end_i; i++) {
       Torus state =
           init_decomposer_state(block_lwe_array_in[i], base_log, level_count);
-      uint32_t offset = i * level_count * (lwe_dimension_out + 1);
-      for (int j = 0; j < level_count; j++) {
 
+      for (int j = 0; j < level_count; j++) {
+        auto ksk_block =
+            get_ith_block(ksk, i, j, lwe_dimension_out, level_count);
         Torus decomposed = decompose_one<Torus>(state, mask_mod_b, base_log);
-        local_lwe_out +=
-            (Torus)ksk[tid + j * (lwe_dimension_out + 1) + offset] * decomposed;
+        local_lwe_out -= (Torus)ksk_block[tid] * decomposed;
       }
     }
 
@@ -103,7 +93,7 @@ keyswitch(Torus *lwe_array_out, const Torus *__restrict__ lwe_output_indexes,
           lwe_acc_out[shmem_index + offset * blockDim.x];
     }
     if (threadIdx.y == 0)
-      block_lwe_array_out[tid] = -lwe_acc_out[shmem_index];
+      block_lwe_array_out[tid] = lwe_acc_out[shmem_index];
   }
 }
 
@@ -125,7 +115,7 @@ __host__ void host_keyswitch_lwe_ciphertext_vector(
 
   int shared_mem = sizeof(Torus) * num_threads_y * num_threads_x;
   if (num_blocks_per_sample > 65536)
-    PANIC("Cuda error (Keyswitch): number of blocks per sample is too large");
+    PANIC("Cuda error (Keyswith): number of blocks per sample is too large");
 
   // In multiplication of large integers (512, 1024, 2048), the number of
   // samples can be larger than 65536, so we need to set it in the first
@@ -182,14 +172,14 @@ __host__ uint64_t scratch_packing_keyswitch_lwe_list_to_glwe(
 
   // allocate at least LWE-mask times two: to keep both decomposition state and
   // decomposed intermediate value
-  uint64_t memory_unit = glwe_accumulator_size > lwe_dimension * 2
-                             ? glwe_accumulator_size
-                             : lwe_dimension * 2;
+  int memory_unit = glwe_accumulator_size > lwe_dimension * 2
+                        ? glwe_accumulator_size
+                        : lwe_dimension * 2;
 
-  uint64_t size_tracker = 0;
-  uint64_t buffer_size = 2 * num_lwes * memory_unit * sizeof(Torus);
+  uint64_t size_tracker;
   *fp_ks_buffer = (int8_t *)cuda_malloc_with_size_tracking_async(
-      buffer_size, stream, gpu_index, size_tracker, allocate_gpu_memory);
+      2 * num_lwes * memory_unit * sizeof(Torus), stream, gpu_index,
+      &size_tracker, allocate_gpu_memory);
   return size_tracker;
 }
 

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

@@ -55,7 +55,7 @@ __global__ void decompose_vectorize_init(Torus const *lwe_in, Torus *lwe_out,
   lwe_out[write_state_idx] = state;
 }
 
-// Continue decomposition of an array of Torus elements in place. Supposes
+// Continue decomposiion of an array of Torus elements in place. Supposes
 // that the array contains already decomposed elements and
 // computes the new decomposed level in place.
 template <typename Torus>

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

@@ -66,13 +66,6 @@ __device__ inline void typecast_torus_to_double<uint64_t>(uint64_t x,
   r = __ll2double_rn(x);
 }
 
-template <>
-__device__ inline void typecast_torus_to_double<__uint128_t>(__uint128_t x,
-                                                             double &r) {
-  // We truncate x
-  r = __ll2double_rn(static_cast<uint64_t>(x));
-}
-
 template <typename T>
 __device__ inline T init_decomposer_state(T input, uint32_t base_log,
                                           uint32_t level_count) {
@@ -91,23 +84,6 @@ __device__ inline T init_decomposer_state(T input, uint32_t base_log,
   return res - (need_balance << rep_bit_count);
 }
 
-template <typename T, uint32_t base_log, uint32_t level_count>
-__device__ inline T init_decomposer_state_2_2_params(T input) {
-  constexpr T rep_bit_count = level_count * base_log;
-  constexpr T non_rep_bit_count = sizeof(T) * 8 - rep_bit_count;
-  T res = input >> (non_rep_bit_count - 1);
-  T rounding_bit = res & (T)(1);
-  res++;
-  res >>= 1;
-  constexpr T torus_max = scalar_max<T>();
-  constexpr T mod_mask = torus_max >> non_rep_bit_count;
-  res &= mod_mask;
-  T shifted_random = rounding_bit << (rep_bit_count - 1);
-  T need_balance =
-      (((res - (T)(1)) | shifted_random) & res) >> (rep_bit_count - 1);
-  return res - (need_balance << rep_bit_count);
-}
-
 template <typename T>
 __device__ __forceinline__ void modulus_switch(T input, T &output,
                                                uint32_t log_modulus) {
@@ -202,12 +178,11 @@ __device__ __forceinline__ double measure_modulus_switch_noise(
 
 // Each thread processes two elements of the lwe array
 template <typename Torus>
-__global__ void __launch_bounds__(512)
-    improve_noise_modulus_switch(Torus *array_out, const Torus *array_in,
-                                 const uint64_t *indexes, const Torus *zeros,
-                                 int lwe_size, int num_zeros,
-                                 double input_variance, double r_sigma,
-                                 double bound, uint32_t log_modulus) {
+__global__ void
+improve_noise_modulus_switch(Torus *array_out, const Torus *array_in,
+                             const Torus *zeros, int lwe_size, int num_zeros,
+                             double input_variance, double r_sigma,
+                             double bound, uint32_t log_modulus) {
 
   // First we will assume size is less than the number of threads per block
   // I should switch this to dynamic shared memory
@@ -223,14 +198,13 @@ __global__ void __launch_bounds__(512)
   // This probably are not needed cause we are setting the values
   sum_mask_errors[threadIdx.x] = 0.f;
   sum_squared_mask_errors[threadIdx.x] = 0.f;
-  auto this_block_lwe_in = array_in + indexes[blockIdx.x] * lwe_size;
-  // We use modulus switch to gather the output in trivial order
-  auto this_block_lwe_out = array_out + blockIdx.x * lwe_size;
-  Torus input_element1 = this_block_lwe_in[threadIdx.x];
 
-  Torus input_element2 = threadIdx.x + blockDim.x < lwe_size
-                             ? this_block_lwe_in[threadIdx.x + blockDim.x]
-                             : 0;
+  Torus input_element1 = array_in[threadIdx.x + blockIdx.x * lwe_size];
+
+  Torus input_element2 =
+      threadIdx.x + blockDim.x < lwe_size
+          ? array_in[threadIdx.x + blockDim.x + blockIdx.x * lwe_size]
+          : 0;
 
   // Base noise is only handled by thread 0
   double base_noise = measure_modulus_switch_noise<Torus>(
@@ -244,10 +218,11 @@ __global__ void __launch_bounds__(512)
   __syncthreads();
 
   if (found)
-    this_block_lwe_out[threadIdx.x] = input_element1;
+    array_out[threadIdx.x + blockIdx.x * lwe_size] = input_element1;
 
   if (found && (threadIdx.x + blockDim.x) < lwe_size)
-    this_block_lwe_out[threadIdx.x + blockDim.x] = input_element2;
+    array_out[threadIdx.x + blockDim.x + blockIdx.x * lwe_size] =
+        input_element2;
 
   __syncthreads();
   // If we found a zero element we stop iterating (in avg 20 times are
@@ -278,10 +253,11 @@ __global__ void __launch_bounds__(512)
     // Assumption we always have at least 512 elements
     // If we find a useful zero encryption we replace the lwe by lwe + zero
     if (found)
-      this_block_lwe_out[threadIdx.x] = zero_element1;
+      array_out[threadIdx.x + blockIdx.x * lwe_size] = zero_element1;
 
     if (found && (threadIdx.x + blockDim.x) < lwe_size)
-      this_block_lwe_out[threadIdx.x + blockDim.x] = zero_element2;
+      array_out[threadIdx.x + blockDim.x + blockIdx.x * lwe_size] =
+          zero_element2;
 
     __syncthreads();
     // If we found a zero element we stop iterating (in avg 20 times are
@@ -294,10 +270,9 @@ __global__ void __launch_bounds__(512)
 template <typename Torus>
 __host__ void host_improve_noise_modulus_switch(
     cudaStream_t stream, uint32_t gpu_index, Torus *array_out,
-    Torus const *array_in, uint64_t const *indexes, const Torus *zeros,
-    uint32_t lwe_size, uint32_t num_lwes, const uint32_t num_zeros,
-    const double input_variance, const double r_sigma, const double bound,
-    uint32_t log_modulus) {
+    Torus const *array_in, const Torus *zeros, uint32_t lwe_size,
+    uint32_t num_lwes, const uint32_t num_zeros, const double input_variance,
+    const double r_sigma, const double bound, uint32_t log_modulus) {
 
   if (lwe_size < 512) {
     PANIC("The lwe_size is less than 512, this is not supported\n");
@@ -314,8 +289,8 @@ __host__ void host_improve_noise_modulus_switch(
   int num_threads = 512, num_blocks = num_lwes;
 
   improve_noise_modulus_switch<Torus><<<num_blocks, num_threads, 0, stream>>>(
-      array_out, array_in, indexes, zeros, lwe_size, num_zeros, input_variance,
-      r_sigma, bound, log_modulus);
+      array_out, array_in, zeros, lwe_size, num_zeros, input_variance, r_sigma,
+      bound, log_modulus);
   check_cuda_error(cudaGetLastError());
 }
 

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

@@ -2,52 +2,8 @@
 #include <cstdint>
 #include <cuda_runtime.h>
 
-uint32_t cuda_get_device() {
-  int device;
-  check_cuda_error(cudaGetDevice(&device));
-  return static_cast<uint32_t>(device);
-}
-
 void cuda_set_device(uint32_t gpu_index) {
   check_cuda_error(cudaSetDevice(gpu_index));
-  static bool SETUP_MEM_AND_WARMUP = 1;
-  if (SETUP_MEM_AND_WARMUP){
-    const size_t warmup_size = 10L * 1024 * 1024 * 1024;  // 10 GB just for testing
-    // Get default memory pool
-    cudaMemPool_t default_pool;
-    check_cuda_error(cudaDeviceGetDefaultMemPool(&default_pool, gpu_index));
-
-    // Enable opportunistic reuse (may be on by default, but explicitly setting it is good practice)
-    int reuse = 1;
-    check_cuda_error(cudaMemPoolSetAttribute(
-        default_pool,
-        cudaMemPoolReuseAllowOpportunistic,
-        &reuse));
-
-    //Prevent memory from being released back to the OS too soon
-    size_t threshold = warmup_size;
-    check_cuda_error(cudaMemPoolSetAttribute(
-        default_pool,
-        cudaMemPoolAttrReleaseThreshold,
-        &threshold));
-
-    // Warm up the pool by allocating and freeing a large block
-    cudaStream_t stream;
-    check_cuda_error(cudaStreamCreate(&stream));
-
-    void* warmup_ptr = nullptr;
-    check_cuda_error(cudaMallocAsync(&warmup_ptr, warmup_size, stream));
-    check_cuda_error(cudaFreeAsync(warmup_ptr, stream));
-
-    // Sync to ensure pool is grown
-    check_cuda_error(cudaStreamSynchronize(stream));
-
-    printf("Default CUDA memory pool warmed up with 10 GB and opportunistic reuse enabled.\n");
-
-    // Clean up
-    check_cuda_error(cudaStreamDestroy(stream));
-    SETUP_MEM_AND_WARMUP = 0;
-  }
 }
 
 cudaEvent_t cuda_create_event(uint32_t gpu_index) {
@@ -118,9 +74,10 @@ void *cuda_malloc(uint64_t size, uint32_t gpu_index) {
 /// asynchronously.
 void *cuda_malloc_with_size_tracking_async(uint64_t size, cudaStream_t stream,
                                            uint32_t gpu_index,
-                                           uint64_t &size_tracker,
+                                           uint64_t *size_tracker,
                                            bool allocate_gpu_memory) {
-  size_tracker += size;
+  if (size_tracker != nullptr)
+    *size_tracker += size;
   void *ptr = nullptr;
   if (!allocate_gpu_memory)
     return ptr;
@@ -149,9 +106,8 @@ void *cuda_malloc_with_size_tracking_async(uint64_t size, cudaStream_t stream,
 /// asynchronously.
 void *cuda_malloc_async(uint64_t size, cudaStream_t stream,
                         uint32_t gpu_index) {
-  uint64_t size_tracker = 0;
-  return cuda_malloc_with_size_tracking_async(size, stream, gpu_index,
-                                              size_tracker, true);
+  return cuda_malloc_with_size_tracking_async(size, stream, gpu_index, nullptr,
+                                              true);
 }
 
 /// Check that allocation is valid
@@ -166,13 +122,6 @@ bool cuda_check_valid_malloc(uint64_t size, uint32_t gpu_index) {
   }
 }
 
-uint64_t cuda_device_total_memory(uint32_t gpu_index) {
-  cuda_set_device(gpu_index);
-  size_t total_mem = 0, free_mem = 0;
-  check_cuda_error(cudaMemGetInfo(&free_mem, &total_mem));
-  return total_mem;
-}
-
 /// Returns
 ///  false if Cooperative Groups is not supported.
 ///  true otherwise

backends/tfhe-cuda-backend/cuda/src/fft/bnsmfft.cuh

@@ -63,7 +63,7 @@ template <class params> __device__ void NSMFFT_direct(double2 *A) {
   }
 
   Index twiddle_shift = 1;
-  for (Index l = LOG2_DEGREE - 1; l >= 5; --l) {
+  for (Index l = LOG2_DEGREE - 1; l >= 1; --l) {
     Index lane_mask = 1 << (l - 1);
     Index thread_mask = (1 << l) - 1;
     twiddle_shift <<= 1;
@@ -96,43 +96,8 @@ template <class params> __device__ void NSMFFT_direct(double2 *A) {
       tid = tid + STRIDE;
     }
   }
-
-  for (Index l = 4; l >= 1; --l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    twiddle_shift <<= 1;
-
-    tid = threadIdx.x;
-    __syncwarp();
-    double2 reg_A[BUTTERFLY_DEPTH];
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      reg_A[i] = (u_stays_in_register) ? v[i] : u[i];
-      tid = tid + STRIDE;
-    }
-    __syncwarp();
-
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      w = shfl_xor_double2(reg_A[i], 1 << (l - 1), 0xFFFFFFFF);
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-      w = negtwiddles[tid / lane_mask + twiddle_shift];
-
-      w *= v[i];
-
-      v[i] = u[i] - w;
-      u[i] = u[i] + w;
-      tid = tid + STRIDE;
-    }
-  }
-
   __syncthreads();
+
   // store registers in SM
   tid = threadIdx.x;
 #pragma unroll
@@ -144,119 +109,6 @@ template <class params> __device__ void NSMFFT_direct(double2 *A) {
   __syncthreads();
 }
 
-/*
- * negacyclic fft optimized for 2_2 params
-   it uses the twiddles from shared memory for extra performance
-   this is possible cause we know for 2_2 params will have memory available
-   the fft is returned in registers to avoid extra synchronizations
- */
-template <class params>
-__device__ void NSMFFT_direct_2_2_params(double2 *A, double2 *fft_out,
-                                         double2 *shared_twiddles) {
-
-  /* We don't make bit reverse here, since twiddles are already reversed
-   *  Each thread is always in charge of "opt/2" pairs of coefficients,
-   *  which is why we always loop through N/2 by N/opt strides
-   *  The pragma unroll instruction tells the compiler to unroll the
-   *  full loop, which should increase performance
-   */
-
-  constexpr Index BUTTERFLY_DEPTH = params::opt >> 1;
-  constexpr Index LOG2_DEGREE = params::log2_degree;
-  constexpr Index HALF_DEGREE = params::degree >> 1;
-  constexpr Index STRIDE = params::degree / params::opt;
-
-  Index tid = threadIdx.x;
-  double2 u[BUTTERFLY_DEPTH], v[BUTTERFLY_DEPTH], w;
-
-  // switch register order
-  for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-    u[i] = fft_out[i];
-    v[i] = fft_out[i + params::opt / 2];
-  }
-
-  // level 1
-  // we don't make actual complex multiplication on level1 since we have only
-  // one twiddle, it's real and image parts are equal, so we can multiply
-  // it with simpler operations
-#pragma unroll
-  for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-    w = v[i] * (double2){0.707106781186547461715008466854,
-                         0.707106781186547461715008466854};
-    v[i] = u[i] - w;
-    u[i] = u[i] + w;
-  }
-
-  Index twiddle_shift = 1;
-  for (Index l = LOG2_DEGREE - 1; l >= 5; --l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    twiddle_shift <<= 1;
-
-    tid = threadIdx.x;
-    __syncthreads();
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      A[tid] = (u_stays_in_register) ? v[i] : u[i];
-      tid = tid + STRIDE;
-    }
-    __syncthreads();
-
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      w = A[tid ^ lane_mask];
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-      w = shared_twiddles[tid / lane_mask + twiddle_shift];
-
-      w *= v[i];
-
-      v[i] = u[i] - w;
-      u[i] = u[i] + w;
-      tid = tid + STRIDE;
-    }
-  }
-
-  for (Index l = 4; l >= 1; --l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    twiddle_shift <<= 1;
-
-    tid = threadIdx.x;
-    double2 reg_A[BUTTERFLY_DEPTH];
-
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      reg_A[i] = (u_stays_in_register) ? v[i] : u[i];
-      w = shfl_xor_double2(reg_A[i], 1 << (l - 1), 0xFFFFFFFF);
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-      w = shared_twiddles[tid / lane_mask + twiddle_shift];
-
-      w *= v[i];
-
-      v[i] = u[i] - w;
-      u[i] = u[i] + w;
-      tid = tid + STRIDE;
-    }
-  }
-
-  // Return result in registers, no need to synchronize here
-  // only with we need to use the same shared memory afterwards
-  for (Index i = 0; i < BUTTERFLY_DEPTH; i++) {
-    fft_out[i] = u[i];
-    fft_out[i + params::opt / 2] = v[i];
-  }
-}
-
 /*
  * negacyclic inverse fft
  */
@@ -292,46 +144,7 @@ template <class params> __device__ void NSMFFT_inverse(double2 *A) {
   }
 
   Index twiddle_shift = DEGREE;
-  for (Index l = 1; l <= 4; ++l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    tid = threadIdx.x;
-    twiddle_shift >>= 1;
-
-    // at this point registers are ready for the  butterfly
-    tid = threadIdx.x;
-    __syncwarp();
-    double2 reg_A[BUTTERFLY_DEPTH];
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      w = (u[i] - v[i]);
-      u[i] += v[i];
-      v[i] = w * conjugate(negtwiddles[tid / lane_mask + twiddle_shift]);
-
-      // keep one of the register for next iteration and store another one in sm
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      reg_A[i] = (u_stays_in_register) ? v[i] : u[i];
-
-      tid = tid + STRIDE;
-    }
-    __syncwarp();
-
-    // prepare registers for next butterfly iteration
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      w = shfl_xor_double2(reg_A[i], 1 << (l - 1), 0xFFFFFFFF);
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-
-      tid = tid + STRIDE;
-    }
-  }
-
-  for (Index l = 5; l <= LOG2_DEGREE - 1; ++l) {
+  for (Index l = 1; l <= LOG2_DEGREE - 1; ++l) {
     Index lane_mask = 1 << (l - 1);
     Index thread_mask = (1 << l) - 1;
     tid = threadIdx.x;
@@ -388,126 +201,6 @@ template <class params> __device__ void NSMFFT_inverse(double2 *A) {
   __syncthreads();
 }
 
-/*
- * negacyclic inverse fft optimized for 2_2 params
- * it uses the twiddles from shared memory for extra performance
- * this is possible cause we know for 2_2 params will have memory available
- * the input comes from registers to avoid some synchronizations and shared mem
- * usage
- */
-template <class params>
-__device__ void NSMFFT_inverse_2_2_params(double2 *A, double2 *buffer_regs,
-                                          double2 *shared_twiddles) {
-
-  /* We don't make bit reverse here, since twiddles are already reversed
-   *  Each thread is always in charge of "opt/2" pairs of coefficients,
-   *  which is why we always loop through N/2 by N/opt strides
-   *  The pragma unroll instruction tells the compiler to unroll the
-   *  full loop, which should increase performance
-   */
-
-  constexpr Index BUTTERFLY_DEPTH = params::opt >> 1;
-  constexpr Index LOG2_DEGREE = params::log2_degree;
-  constexpr Index DEGREE = params::degree;
-  constexpr Index HALF_DEGREE = params::degree >> 1;
-  constexpr Index STRIDE = params::degree / params::opt;
-
-  size_t tid = threadIdx.x;
-  double2 u[BUTTERFLY_DEPTH], v[BUTTERFLY_DEPTH], w;
-
-  // load into registers and divide by compressed polynomial size
-  for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-    u[i] = buffer_regs[i];
-    v[i] = buffer_regs[i + params::opt / 2];
-
-    u[i] /= DEGREE;
-    v[i] /= DEGREE;
-  }
-
-  Index twiddle_shift = DEGREE;
-  for (Index l = 1; l <= 4; ++l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    tid = threadIdx.x;
-    twiddle_shift >>= 1;
-
-    // at this point registers are ready for the  butterfly
-    tid = threadIdx.x;
-    double2 reg_A[BUTTERFLY_DEPTH];
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      w = (u[i] - v[i]);
-      u[i] += v[i];
-      v[i] = w * conjugate(shared_twiddles[tid / lane_mask + twiddle_shift]);
-
-      tid = tid + STRIDE;
-    }
-    __syncwarp();
-
-    // prepare registers for next butterfly iteration
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      reg_A[i] = (u_stays_in_register) ? v[i] : u[i];
-      w = shfl_xor_double2(reg_A[i], 1 << (l - 1), 0xFFFFFFFF);
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-
-      tid = tid + STRIDE;
-    }
-  }
-
-  for (Index l = 5; l <= LOG2_DEGREE - 1; ++l) {
-    Index lane_mask = 1 << (l - 1);
-    Index thread_mask = (1 << l) - 1;
-    tid = threadIdx.x;
-    twiddle_shift >>= 1;
-
-    // at this point registers are ready for the  butterfly
-    tid = threadIdx.x;
-    __syncthreads();
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      w = (u[i] - v[i]);
-      u[i] += v[i];
-      v[i] = w * conjugate(shared_twiddles[tid / lane_mask + twiddle_shift]);
-
-      // keep one of the register for next iteration and store another one in sm
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      A[tid] = (u_stays_in_register) ? v[i] : u[i];
-
-      tid = tid + STRIDE;
-    }
-    __syncthreads();
-
-    // prepare registers for next butterfly iteration
-    tid = threadIdx.x;
-#pragma unroll
-    for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-      Index rank = tid & thread_mask;
-      bool u_stays_in_register = rank < lane_mask;
-      w = A[tid ^ lane_mask];
-      u[i] = (u_stays_in_register) ? u[i] : w;
-      v[i] = (u_stays_in_register) ? w : v[i];
-
-      tid = tid + STRIDE;
-    }
-  }
-
-// last iteration
-#pragma unroll
-  for (Index i = 0; i < BUTTERFLY_DEPTH; ++i) {
-    w = (u[i] - v[i]);
-    buffer_regs[i] = u[i] + v[i];
-    buffer_regs[i + params::opt / 2] =
-        w * (double2){0.707106781186547461715008466854,
-                      -0.707106781186547461715008466854};
-  }
-}
-
 /*
  * global batch fft
  * does fft in half size

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

@@ -234,29 +234,6 @@ __device__ void convert_u128_to_f128_as_torus(
   }
 }
 
-// params is expected to be full degree not half degree
-// same as convert_u128_to_f128_as_torus() but expects input to be on registers
-template <class params>
-__device__ void convert_u128_on_regs_to_f128_as_torus(
-    double *out_re_hi, double *out_re_lo, double *out_im_hi, double *out_im_lo,
-    const __uint128_t *in_re_on_regs, const __uint128_t *in_im_on_regs) {
-
-  const double normalization = pow(2., -128.);
-  Index tid = threadIdx.x;
-  // #pragma unroll
-  for (Index i = 0; i < params::opt / 2; i++) {
-    auto out_re = u128_to_signed_to_f128(in_re_on_regs[i]);
-    auto out_im = u128_to_signed_to_f128(in_im_on_regs[i]);
-
-    out_re_hi[tid] = out_re.hi * normalization;
-    out_re_lo[tid] = out_re.lo * normalization;
-    out_im_hi[tid] = out_im.hi * normalization;
-    out_im_lo[tid] = out_im.lo * normalization;
-
-    tid += params::degree / params::opt;
-  }
-}
-
 template <class params>
 __device__ void
 convert_f128_to_u128_as_torus(__uint128_t *out_re, __uint128_t *out_im,
@@ -295,7 +272,7 @@ batch_convert_u128_to_f128_as_integer(double *out_re_hi, double *out_re_lo,
 }
 
 // params is expected to be full degree not half degree
-// converts standard input into complex<128> represented by 4 double
+// converts standqard input into complex<128> represented by 4 double
 // with following pattern: [re_hi_0, re_hi_1, ... re_hi_n, re_lo_0, re_lo_1,
 // ... re_lo_n, im_hi_0, im_hi_1, ..., im_hi_n,  im_lo_0, im_lo_1, ..., im_lo_n]
 template <class params>
@@ -314,7 +291,7 @@ batch_convert_u128_to_f128_as_torus(double *out_re_hi, double *out_re_lo,
 }
 
 // params is expected to be full degree not half degree
-// converts standard input into complex<128> represented by 4 double
+// converts standqard input into complex<128> represented by 4 double
 // with following pattern: [re_hi_0, re_lo_0, im_hi_0, im_lo_0, re_hi_1,
 // re_lo_1, im_hi_1, im_lo_1,
 // ...,re_hi_n, re_lo_n, im_hi_n, im_lo_n, ]
@@ -515,7 +492,6 @@ __host__ void host_fourier_transform_forward_as_integer_f128(
   batch_convert_u128_to_f128_as_integer<params>
       <<<grid_size, block_size, 0, stream>>>(d_re0, d_re1, d_im0, d_im1,
                                              d_standard);
-  check_cuda_error(cudaGetLastError());
 
   // call negacyclic 128 bit forward fft.
   if (full_sm) {
@@ -527,7 +503,6 @@ __host__ void host_fourier_transform_forward_as_integer_f128(
         <<<grid_size, block_size, shared_memory_size, stream>>>(
             d_re0, d_re1, d_im0, d_im1, d_re0, d_re1, d_im0, d_im1, buffer);
   }
-  check_cuda_error(cudaGetLastError());
 
   cuda_memcpy_async_to_cpu(re0, d_re0, N / 2 * sizeof(double), stream,
                            gpu_index);

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

@@ -26,7 +26,7 @@ __host__ uint64_t scratch_cuda_integer_abs_kb(
   if (is_signed) {
     *mem_ptr = new int_abs_buffer<Torus>(streams, gpu_indexes, gpu_count,
                                          params, num_blocks,
-                                         allocate_gpu_memory, size_tracker);
+                                         allocate_gpu_memory, &size_tracker);
   }
   return size_tracker;
 }
@@ -53,8 +53,7 @@ __host__ void host_integer_abs_kb(
       streams, gpu_indexes, gpu_count, mask, num_bits_in_ciphertext - 1,
       mem_ptr->arithmetic_scalar_shift_mem, bsks, ksks, ms_noise_reduction_key);
   host_addition<Torus>(streams[0], gpu_indexes[0], ct, mask, ct,
-                       ct->num_radix_blocks, mem_ptr->params.message_modulus,
-                       mem_ptr->params.carry_modulus);
+                       ct->num_radix_blocks);
 
   uint32_t requested_flag = outputFlag::FLAG_NONE;
   uint32_t uses_carry = 0;