From 629e22aaacf4f6faccbbf63b5322cc00030749e6 Mon Sep 17 00:00:00 2001 From: Agnes Leroy Date: Wed, 15 Mar 2023 10:55:08 +0100 Subject: [PATCH] chore(concrete_cuda): fix tests, reduce test time - Update concrete-cuda with latest commits in concrete-core - Add C++ tests based on google test - Run the tests in the CI with Cuda 11.8 and Cuda 11.1 - Test for many PBS/KS parameters - Repetitions & samples are set for each parameter set in the PBS test --- .github/workflows/concrete_cuda_test.yml | 16 +- .../implementation/CMakeLists.txt | 3 - .../implementation/include/device.h | 6 +- .../implementation/parameters/CMakeLists.txt | 4 - .../implementation/parameters/parameters.cpp | 380 ------------------ .../implementation/src/boolean_gates.cu | 4 +- .../implementation/src/device.cu | 13 +- .../test/test_bit_extraction.cpp | 38 +- .../implementation/test/test_bootstrap.cpp | 115 +++--- .../test/test_circuit_bootstrap.cpp | 4 +- .../implementation/test/test_cmux_tree.cpp | 3 +- .../implementation/test/test_keyswitch.cpp | 20 +- .../test/test_wop_bootstrap.cpp | 50 ++- .../implementation/test/utils.cpp | 35 +- 14 files changed, 142 insertions(+), 549 deletions(-) delete mode 100644 backends/concrete-cuda/implementation/parameters/CMakeLists.txt delete mode 100644 backends/concrete-cuda/implementation/parameters/parameters.cpp diff --git a/.github/workflows/concrete_cuda_test.yml b/.github/workflows/concrete_cuda_test.yml index 15839c6e9..c93a76fdc 100644 --- a/.github/workflows/concrete_cuda_test.yml +++ b/.github/workflows/concrete_cuda_test.yml @@ -108,7 +108,7 @@ jobs: echo "CUDA_PATH=$OLD_CUDA_PATH" >> "${GITHUB_ENV}" echo "$CUDA_PATH/bin" >> "${GITHUB_PATH}" echo "LD_LIBRARY_PATH=$OLD_CUDA_PATH/lib:$LD_LIBRARY_PATH" >> "${GITHUB_ENV}" - echo "CUDACXX=$OLD_CUDA_PATH/bin/nvcc" >> "${GITHUB_ENV}" + echo "CUDACXX=/usr/local/cuda-${{ matrix.old_cuda }}/bin/nvcc" >> "${GITHUB_ENV}" - name: Build concrete-cuda with Cuda 11.1 if: ${{ !cancelled() }} run: | @@ -121,19 +121,7 @@ jobs: if: ${{ !cancelled() }} run: | cd backends/concrete-cuda/implementation/build-old-cuda - ./test/test_concrete_cuda - - - name: Slack Notification - if: ${{ always() }} - continue-on-error: true - uses: rtCamp/action-slack-notify@12e36fc18b0689399306c2e0b3e0f2978b7f1ee7 - env: - SLACK_COLOR: ${{ job.status }} - SLACK_CHANNEL: ${{ secrets.SLACK_CHANNEL }} - SLACK_ICON: https://pbs.twimg.com/profile_images/1274014582265298945/OjBKP9kn_400x400.png - SLACK_MESSAGE: "Cuda AWS tests finished with status ${{ job.status }}. (${{ env.ACTION_RUN_URL }})" - SLACK_USERNAME: ${{ secrets.BOT_USERNAME }} - SLACK_WEBHOOK: ${{ secrets.SLACK_WEBHOOK }} + ./test/test_concrete_cuda --gtest_filter="Wop*" stop-runner: name: Stop EC2 runner diff --git a/backends/concrete-cuda/implementation/CMakeLists.txt b/backends/concrete-cuda/implementation/CMakeLists.txt index 38bd4d0f1..15d4b32a1 100644 --- a/backends/concrete-cuda/implementation/CMakeLists.txt +++ b/backends/concrete-cuda/implementation/CMakeLists.txt @@ -68,14 +68,11 @@ set(INCLUDE_DIR include) add_subdirectory(src) add_subdirectory(test) -add_subdirectory(parameters) target_include_directories(concrete_cuda PRIVATE ${INCLUDE_DIR}) # This is required for rust cargo build install(TARGETS concrete_cuda DESTINATION .) install(TARGETS concrete_cuda DESTINATION lib) -install(TARGETS cuda_parameters DESTINATION .) -install(TARGETS cuda_parameters DESTINATION lib) # Define a function to add a lint target. find_file(CPPLINT NAMES cpplint cpplint.exe) diff --git a/backends/concrete-cuda/implementation/include/device.h b/backends/concrete-cuda/implementation/include/device.h index 9bcdcc886..706d44043 100644 --- a/backends/concrete-cuda/implementation/include/device.h +++ b/backends/concrete-cuda/implementation/include/device.h @@ -4,10 +4,10 @@ #pragma once #include +#include +#include +#include #include -#include -#include -#include extern "C" { cudaStream_t *cuda_create_stream(uint32_t gpu_index); diff --git a/backends/concrete-cuda/implementation/parameters/CMakeLists.txt b/backends/concrete-cuda/implementation/parameters/CMakeLists.txt deleted file mode 100644 index 28e1cf9da..000000000 --- a/backends/concrete-cuda/implementation/parameters/CMakeLists.txt +++ /dev/null @@ -1,4 +0,0 @@ -file(GLOB SOURCES - "parameters.cpp") -add_library(cuda_parameters STATIC ${SOURCES}) -set_target_properties(cuda_parameters PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) diff --git a/backends/concrete-cuda/implementation/parameters/parameters.cpp b/backends/concrete-cuda/implementation/parameters/parameters.cpp deleted file mode 100644 index 4795db8bc..000000000 --- a/backends/concrete-cuda/implementation/parameters/parameters.cpp +++ /dev/null @@ -1,380 +0,0 @@ - - -#include -using namespace std; - -const int NORM2_MAX = 31; -const int P_MAX = 7; - -typedef struct V0Parameter { - int k; - int polynomialSize; - int nSmall; - int brLevel; - int brLogBase; - int ksLevel; - int ksLogBase; - - V0Parameter(int k_, int polynomialSize_, int nSmall_, int brLevel_, - int brLogBase_, int ksLevel_, int ksLogBase_) { - k = k_; - polynomialSize = polynomialSize_; - nSmall = nSmall_; - brLevel = brLevel_; - brLogBase = brLogBase_; - ksLevel = ksLevel_; - ksLogBase = ksLogBase_; - } - -} V0Parameter; - -typedef struct V0Variances { - float logstdEncrypt; - float logstdDecrypt; - - V0Variances(float logstdEncrypt_, float logstdDecrypt_) { - logstdEncrypt = logstdEncrypt_; - logstdDecrypt = logstdDecrypt_; - } - -} V0Variances; - -V0Parameter parameters[NORM2_MAX][P_MAX] = { - {V0Parameter(1, 10, 472, 2, 8, 4, 2), V0Parameter(1, 10, 514, 2, 8, 5, 2), - V0Parameter(1, 10, 564, 2, 8, 5, 2), V0Parameter(1, 10, 599, 3, 6, 6, 2), - V0Parameter(1, 10, 686, 3, 6, 7, 2), V0Parameter(1, 11, 736, 1, 23, 5, 3), - V0Parameter(1, 12, 830, 1, 23, 4, 4)}, - {V0Parameter(1, 10, 474, 2, 8, 4, 2), V0Parameter(1, 10, 519, 2, 8, 5, 2), - V0Parameter(1, 10, 558, 3, 6, 5, 2), V0Parameter(1, 10, 610, 3, 6, 6, 2), - V0Parameter(1, 11, 689, 1, 23, 4, 3), V0Parameter(1, 11, 736, 1, 23, 5, 3), - V0Parameter(1, 12, 831, 1, 23, 4, 4)}, - {V0Parameter(1, 10, 479, 2, 8, 4, 2), V0Parameter(1, 10, 515, 3, 6, 5, 2), - V0Parameter(1, 10, 569, 3, 6, 5, 2), V0Parameter(1, 11, 638, 1, 23, 4, 3), - V0Parameter(1, 11, 689, 1, 23, 4, 3), V0Parameter(1, 11, 737, 1, 23, 5, 3), - V0Parameter(1, 12, 840, 1, 23, 4, 4)}, - {V0Parameter(1, 10, 531, 2, 8, 5, 2), V0Parameter(1, 10, 523, 3, 6, 5, 2), - V0Parameter(1, 11, 598, 1, 23, 4, 3), V0Parameter(1, 11, 639, 1, 23, 4, 3), - V0Parameter(1, 11, 690, 1, 23, 4, 3), V0Parameter(1, 11, 739, 1, 23, 5, 3), - V0Parameter(1, 12, 806, 2, 16, 5, 3)}, - {V0Parameter(1, 10, 483, 3, 6, 4, 2), V0Parameter(1, 11, 563, 1, 23, 3, 3), - V0Parameter(1, 11, 598, 1, 23, 4, 3), V0Parameter(1, 11, 639, 1, 23, 4, 3), - V0Parameter(1, 11, 691, 1, 23, 4, 3), V0Parameter(1, 11, 748, 1, 23, 5, 3), - V0Parameter(1, 12, 806, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 497, 1, 23, 4, 2), V0Parameter(1, 11, 563, 1, 23, 3, 3), - V0Parameter(1, 11, 598, 1, 23, 4, 3), V0Parameter(1, 11, 640, 1, 23, 4, 3), - V0Parameter(1, 11, 699, 1, 23, 4, 3), V0Parameter(1, 11, 736, 2, 15, 5, 3), - V0Parameter(1, 12, 806, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 497, 1, 23, 4, 2), V0Parameter(1, 11, 563, 1, 23, 3, 3), - V0Parameter(1, 11, 599, 1, 23, 4, 3), V0Parameter(1, 11, 643, 1, 23, 4, 3), - V0Parameter(1, 11, 721, 1, 23, 5, 3), V0Parameter(1, 11, 736, 2, 15, 5, 3), - V0Parameter(1, 12, 806, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 497, 1, 23, 4, 2), V0Parameter(1, 11, 564, 1, 23, 3, 3), - V0Parameter(1, 11, 602, 1, 23, 4, 3), V0Parameter(1, 11, 671, 1, 23, 4, 3), - V0Parameter(1, 11, 689, 2, 15, 4, 3), V0Parameter(1, 11, 736, 2, 15, 5, 3), - V0Parameter(1, 12, 807, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 498, 1, 23, 4, 2), V0Parameter(1, 11, 569, 1, 23, 3, 3), - V0Parameter(1, 11, 622, 1, 23, 4, 3), V0Parameter(1, 11, 638, 2, 15, 4, 3), - V0Parameter(1, 11, 689, 2, 16, 4, 3), V0Parameter(1, 11, 736, 2, 16, 5, 3), - V0Parameter(1, 12, 809, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 502, 1, 23, 4, 2), V0Parameter(1, 11, 555, 1, 23, 5, 2), - V0Parameter(1, 11, 579, 2, 15, 5, 2), V0Parameter(1, 11, 638, 2, 15, 4, 3), - V0Parameter(1, 11, 689, 2, 15, 4, 3), V0Parameter(1, 11, 737, 2, 15, 5, 3), - V0Parameter(1, 12, 818, 2, 15, 5, 3)}, - {V0Parameter(1, 11, 537, 1, 23, 3, 3), V0Parameter(1, 11, 532, 2, 15, 5, 2), - V0Parameter(1, 11, 579, 2, 15, 5, 2), V0Parameter(1, 11, 638, 2, 15, 4, 3), - V0Parameter(1, 11, 690, 2, 15, 4, 3), V0Parameter(1, 11, 738, 2, 15, 5, 3), - V0Parameter(1, 12, 832, 2, 15, 9, 2)}, - {V0Parameter(1, 11, 497, 2, 15, 4, 2), V0Parameter(1, 11, 532, 2, 15, 5, 2), - V0Parameter(1, 11, 579, 2, 15, 5, 2), V0Parameter(1, 11, 639, 2, 15, 4, 3), - V0Parameter(1, 11, 691, 2, 15, 4, 3), V0Parameter(1, 11, 743, 2, 16, 5, 3), - V0Parameter(1, 12, 807, 3, 12, 5, 3)}, - {V0Parameter(1, 11, 497, 2, 15, 4, 2), V0Parameter(1, 11, 532, 2, 15, 5, 2), - V0Parameter(1, 11, 579, 2, 16, 5, 2), V0Parameter(1, 11, 639, 2, 15, 4, 3), - V0Parameter(1, 11, 695, 2, 16, 4, 3), - V0Parameter(1, 11, 757, 2, 16, 16, 1), - V0Parameter(1, 12, 811, 3, 12, 5, 3)}, - {V0Parameter(1, 11, 497, 2, 16, 4, 2), V0Parameter(1, 11, 533, 2, 15, 5, 2), - V0Parameter(1, 11, 580, 2, 15, 5, 2), V0Parameter(1, 11, 641, 2, 16, 4, 3), - V0Parameter(1, 11, 699, 2, 16, 5, 3), V0Parameter(1, 11, 737, 3, 12, 5, 3), - V0Parameter(1, 12, 788, 3, 12, 8, 2)}, - {V0Parameter(1, 11, 497, 2, 16, 4, 2), V0Parameter(1, 11, 533, 2, 15, 5, 2), - V0Parameter(1, 11, 583, 2, 16, 5, 2), V0Parameter(1, 11, 653, 2, 16, 4, 3), - V0Parameter(1, 11, 665, 3, 12, 6, 2), V0Parameter(1, 11, 738, 3, 12, 5, 3), - V0Parameter(1, 12, 775, 4, 9, 8, 2)}, - {V0Parameter(1, 11, 498, 2, 15, 4, 2), V0Parameter(1, 11, 535, 2, 16, 5, 2), - V0Parameter(1, 11, 610, 2, 16, 4, 3), V0Parameter(1, 11, 614, 3, 12, 6, 2), - V0Parameter(1, 11, 666, 3, 12, 6, 2), V0Parameter(1, 11, 747, 3, 12, 5, 3), - V0Parameter(1, 12, 782, 4, 9, 8, 2)}, - {V0Parameter(1, 11, 500, 2, 16, 4, 2), V0Parameter(1, 11, 544, 2, 16, 5, 2), - V0Parameter(1, 11, 580, 3, 12, 5, 2), V0Parameter(1, 11, 615, 3, 12, 6, 2), - V0Parameter(1, 11, 661, 3, 12, 7, 2), V0Parameter(1, 11, 715, 4, 9, 7, 2), - V0Parameter(1, 12, 778, 5, 8, 8, 2)}, - {V0Parameter(1, 11, 513, 2, 16, 4, 2), V0Parameter(1, 11, 533, 3, 12, 5, 2), - V0Parameter(1, 11, 581, 3, 12, 5, 2), V0Parameter(1, 11, 618, 3, 12, 6, 2), - V0Parameter(1, 11, 687, 3, 12, 7, 2), V0Parameter(1, 11, 726, 4, 9, 7, 2), - V0Parameter(1, 12, 809, 5, 8, 8, 2)}, - {V0Parameter(1, 11, 497, 3, 12, 4, 2), V0Parameter(1, 11, 533, 3, 12, 5, 2), - V0Parameter(1, 11, 585, 3, 12, 5, 2), V0Parameter(1, 11, 639, 3, 12, 6, 2), - V0Parameter(1, 11, 662, 4, 9, 7, 2), V0Parameter(1, 11, 717, 5, 8, 7, 2), - V0Parameter(1, 12, 820, 6, 7, 9, 2)}, - {V0Parameter(1, 11, 498, 3, 12, 4, 2), V0Parameter(1, 11, 536, 3, 12, 5, 2), - V0Parameter(1, 11, 593, 3, 12, 6, 2), V0Parameter(1, 11, 619, 4, 9, 6, 2), - V0Parameter(1, 11, 693, 4, 9, 7, 2), V0Parameter(1, 11, 737, 5, 8, 7, 2), - V0Parameter(1, 12, 788, 8, 5, 8, 2)}, - {V0Parameter(1, 11, 502, 3, 12, 4, 2), V0Parameter(1, 11, 552, 3, 12, 5, 2), - V0Parameter(1, 11, 585, 4, 9, 5, 2), V0Parameter(1, 11, 644, 4, 9, 6, 2), - V0Parameter(1, 11, 665, 5, 8, 7, 2), V0Parameter(1, 11, 736, 6, 7, 8, 2), - V0Parameter(1, 12, 786, 11, 4, 8, 2)}, - {V0Parameter(1, 11, 508, 3, 12, 5, 2), V0Parameter(1, 11, 536, 4, 9, 5, 2), - V0Parameter(1, 11, 596, 4, 9, 6, 2), V0Parameter(1, 11, 621, 5, 8, 6, 2), - V0Parameter(1, 11, 667, 6, 7, 7, 2), V0Parameter(1, 11, 746, 7, 6, 8, 2), - V0Parameter(1, 12, 798, 14, 3, 9, 2)}, - {V0Parameter(1, 11, 502, 4, 9, 4, 2), V0Parameter(1, 11, 555, 4, 9, 5, 2), - V0Parameter(1, 11, 580, 5, 8, 6, 2), V0Parameter(1, 11, 623, 6, 7, 6, 2), - V0Parameter(1, 11, 669, 7, 6, 7, 2), V0Parameter(1, 11, 723, 11, 4, 7, 2), - V0Parameter(1, 12, 814, 22, 2, 9, 2)}, - {V0Parameter(1, 11, 510, 4, 9, 5, 2), V0Parameter(1, 11, 539, 5, 8, 5, 2), - V0Parameter(1, 11, 636, 5, 8, 6, 2), V0Parameter(1, 11, 625, 7, 6, 6, 2), - V0Parameter(1, 11, 674, 9, 5, 7, 2), V0Parameter(1, 11, 735, 14, 3, 8, 2), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 498, 5, 8, 5, 2), V0Parameter(1, 11, 579, 5, 8, 6, 2), - V0Parameter(1, 11, 583, 7, 6, 6, 2), V0Parameter(1, 11, 661, 8, 5, 7, 2), - V0Parameter(1, 11, 681, 11, 4, 7, 2), V0Parameter(1, 11, 736, 22, 2, 8, 2), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 530, 5, 8, 5, 2), V0Parameter(1, 11, 541, 7, 6, 5, 2), - V0Parameter(1, 11, 611, 8, 5, 6, 2), V0Parameter(1, 11, 635, 11, 4, 6, 2), - V0Parameter(1, 11, 704, 15, 3, 7, 2), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 565, 6, 7, 5, 2), V0Parameter(1, 11, 569, 8, 5, 5, 2), - V0Parameter(1, 11, 590, 11, 4, 6, 2), V0Parameter(1, 11, 647, 15, 3, 7, 2), - V0Parameter(1, 11, 679, 44, 1, 14, 1), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 520, 8, 5, 5, 2), V0Parameter(1, 11, 549, 11, 4, 5, 2), - V0Parameter(1, 11, 600, 15, 3, 6, 2), - V0Parameter(1, 11, 628, 44, 1, 13, 1), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 506, 11, 4, 5, 2), V0Parameter(1, 11, 559, 15, 3, 5, 2), - V0Parameter(1, 11, 584, 44, 1, 12, 1), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 503, 15, 3, 9, 1), - V0Parameter(1, 11, 594, 23, 2, 12, 1), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0)}, - {V0Parameter(1, 11, 545, 22, 2, 11, 1), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0), V0Parameter(0, 0, 0, 0, 0, 0, 0), - V0Parameter(0, 0, 0, 0, 0, 0, 0)}}; - -V0Variances variances[NORM2_MAX][P_MAX] = { - {V0Variances(-7.186489389863581, -8.186489389863581), - V0Variances(-7.124998639947563, -8.124998639947563), - V0Variances(-7.058035238345106, -8.058035238345106), - V0Variances(-8.771444355069676, -9.771444355069676), - V0Variances(-8.673618068377664, -9.673618068377664), - V0Variances(-14.282552902365403, -15.282552902365403), - V0Variances(-13.46973411689919, -14.46973411689919)}, - {V0Variances(-6.183439290095372, -8.183439290095372), - V0Variances(-6.118015550365563, -8.118015550365563), - V0Variances(-7.822589795549476, -9.822589795549476), - V0Variances(-7.758317735238947, -9.758317735238947), - V0Variances(-13.330153794041763, -15.330153794041763), - V0Variances(-13.282552902365403, -15.282552902365403), - V0Variances(-12.468865546631157, -14.468865546631157)}, - {V0Variances(-5.175869991676414, -8.175869991676414), - V0Variances(-6.8804361404287775, -9.880436140428777), - V0Variances(-6.808508030310776, -9.808508030310776), - V0Variances(-12.38562757351147, -15.38562757351147), - V0Variances(-12.330153794041763, -15.330153794041763), - V0Variances(-12.281573475846372, -15.281573475846372), - V0Variances(-11.46109512118327, -14.46109512118327)}, - {V0Variances(-4.101526889142427, -8.101526889142427), - V0Variances(-5.869316883340595, -9.869316883340595), - V0Variances(-11.432333043294854, -15.432333043294854), - V0Variances(-11.384497819920412, -15.384497819920412), - V0Variances(-11.329107604561145, -15.329107604561145), - V0Variances(-11.279618603320834, -15.279618603320834), - V0Variances(-17.494100927447505, -21.494100927447505)}, - {V0Variances(-4.926710762046184, -9.926710762046184), - V0Variances(-10.475838324353795, -15.475838324353795), - V0Variances(-10.432333043294854, -15.432333043294854), - V0Variances(-10.384497819920412, -15.384497819920412), - V0Variances(-10.328062930199778, -15.328062930199778), - V0Variances(-10.270886650450088, -15.270886650450088), - V0Variances(-16.68961710493341, -21.68961710493341)}, - {V0Variances(-9.565782859612767, -15.565782859612767), - V0Variances(-9.475838324353795, -15.475838324353795), - V0Variances(-9.432333043294854, -15.432333043294854), - V0Variances(-9.38336983295023, -15.38336983295023), - V0Variances(-9.319759557706192, -15.319759557706192), - V0Variances(-16.395716120060897, -22.395716120060897), - V0Variances(-15.689617104933411, -21.68961710493341)}, - {V0Variances(-8.565782859612767, -15.565782859612767), - V0Variances(-8.475838324353795, -15.475838324353795), - V0Variances(-8.431127783999514, -15.431127783999514), - V0Variances(-8.37999641672993, -15.37999641672993), - V0Variances(-8.297406155773494, -15.297406155773494), - V0Variances(-15.395716120060897, -22.395716120060897), - V0Variances(-14.689617104933411, -21.68961710493341)}, - {V0Variances(-7.565782859612767, -15.565782859612767), - V0Variances(-7.4745582041945084, -15.474558204194508), - V0Variances(-7.427524042014056, -15.427524042014056), - V0Variances(-7.349249402293815, -15.349249402293815), - V0Variances(-14.443317011737257, -22.443317011737257), - V0Variances(-14.395716120060897, -22.395716120060897), - V0Variances(-13.688722687558503, -21.688722687558503)}, - {V0Variances(-6.564332914359866, -15.564332914359866), - V0Variances(-6.4681914592406144, -15.468191459240614), - V0Variances(-6.403948495328606, -15.403948495328606), - V0Variances(-13.498790791206964, -22.498790791206964), - V0Variances(-13.669055725537874, -22.669055725537874), - V0Variances(-13.621454833861513, -22.621454833861513), - V0Variances(-12.686937172982404, -21.686937172982404)}, - {V0Variances(-5.558562103418524, -15.558562103418524), - V0Variances(-5.486161899775176, -15.486161899775176), - V0Variances(-12.568787329094782, -22.568787329094782), - V0Variances(-12.498790791206964, -22.498790791206964), - V0Variances(-12.443317011737257, -22.443317011737257), - V0Variances(-12.394736693541866, -22.394736693541866), - V0Variances(-11.678956602726522, -21.678956602726522)}, - {V0Variances(-4.509944741401199, -15.509944741401199), - V0Variances(-11.629855880338809, -22.62985588033881), - V0Variances(-11.568787329094782, -22.568787329094782), - V0Variances(-11.498790791206964, -22.498790791206964), - V0Variances(-11.442270822256638, -22.44227082225664), - V0Variances(-11.393758595059204, -22.393758595059204), - V0Variances(-10.66671526012685, -21.66671526012685)}, - {V0Variances(-10.678946077308261, -22.67894607730826), - V0Variances(-10.629855880338809, -22.62985588033881), - V0Variances(-10.568787329094782, -22.568787329094782), - V0Variances(-10.497661037615906, -22.497661037615906), - V0Variances(-10.441226147895271, -22.44122614789527), - V0Variances(-10.614626611620722, -22.61462661162072), - V0Variances(-13.208593433538631, -25.20859343353863)}, - {V0Variances(-9.678946077308261, -22.67894607730826), - V0Variances(-9.629855880338809, -22.62985588033881), - V0Variances(-9.794526042895399, -22.7945260428954), - V0Variances(-9.497661037615906, -22.497661037615906), - V0Variances(-9.662801228084582, -22.662801228084582), - V0Variances(-9.601161066897156, -22.601161066897156), - V0Variances(-12.205026813068883, -25.205026813068883)}, - {V0Variances(-8.904684791108878, -22.904684791108878), - V0Variances(-8.628501236709816, -22.628501236709816), - V0Variances(-8.567542553081935, -22.567542553081935), - V0Variances(-8.72114553858065, -22.72114553858065), - V0Variances(-8.658661489202302, -22.658661489202302), - V0Variances(-12.345258643192501, -26.3452586431925), - V0Variances(-11.225779955449333, -25.225779955449333)}, - {V0Variances(-7.904684791108878, -22.904684791108878), - V0Variances(-7.628501236709816, -22.628501236709816), - V0Variances(-7.789559775289774, -22.789559775289774), - V0Variances(-7.707766221116806, -22.707766221116806), - V0Variances(-11.41941378254576, -26.41941378254576), - V0Variances(-11.34428054470984, -26.34428054470984), - V0Variances(-12.519191629935513, -27.519191629935513)}, - {V0Variances(-6.67749613205536, -22.67749613205536), - V0Variances(-6.851538271245765, -22.851538271245765), - V0Variances(-6.756903095664896, -22.756903095664896), - V0Variances(-10.476971625054944, -26.476971625054944), - V0Variances(-10.418329864204402, -26.418329864204402), - V0Variances(-10.335536831345415, -26.335536831345415), - V0Variances(-11.512705481362637, -27.512705481362637)}, - {V0Variances(-5.900343669558978, -22.90034366955898), - V0Variances(-5.839504391264846, -22.839504391264846), - V0Variances(-9.518064502732571, -26.51806450273257), - V0Variances(-9.475797747626736, -26.475797747626736), - V0Variances(-9.42376581698619, -26.42376581698619), - V0Variances(-11.303434156979073, -28.303434156979073), - V0Variances(-11.872866817284404, -28.872866817284404)}, - {V0Variances(-4.88182830408649, -22.88182830408649), - V0Variances(-8.579023186360452, -26.579023186360452), - V0Variances(-8.51682187103777, -26.51682187103777), - V0Variances(-8.47228753378785, -26.47228753378785), - V0Variances(-8.395935903330987, -26.395935903330987), - V0Variances(-10.292421003814077, -28.292421003814077), - V0Variances(-10.844682043562514, -28.844682043562514)}, - {V0Variances(-7.629468026958897, -26.629468026958897), - V0Variances(-7.579023186360452, -26.579023186360452), - V0Variances(-7.511872640504656, -26.511872640504656), - V0Variances(-7.4481829872665415, -26.44818298726654), - V0Variances(-9.358990169408344, -28.358990169408344), - V0Variances(-11.010736901553656, -30.010736901553656), - V0Variances(-10.711045072243067, -29.711045072243067)}, - {V0Variances(-6.628018081705996, -26.628018081705996), - V0Variances(-6.5749744525111495, -26.57497445251115), - V0Variances(-6.502074900467036, -26.502074900467036), - V0Variances(-8.40743607320482, -28.40743607320482), - V0Variances(-8.325978101743345, -28.325978101743345), - V0Variances(-9.990891151357076, -29.990891151357076), - V0Variances(-11.223620949776631, -31.22362094977663)}, - {V0Variances(-5.622247270764653, -26.622247270764653), - V0Variances(-5.5537568193509514, -26.55375681935095), - V0Variances(-7.4481874655765665, -28.448187465576567), - V0Variances(-7.378875433754644, -28.378875433754644), - V0Variances(-9.065046290710335, -30.065046290710335), - V0Variances(-9.877481950164196, -30.877481950164196), - V0Variances(-11.301978588237922, -32.30197858823793)}, - {V0Variances(-4.613676704353956, -26.613676704353956), - V0Variances(-6.511289277583067, -28.511289277583067), - V0Variances(-6.434749612580873, -28.434749612580873), - V0Variances(-8.114426826794364, -30.114426826794364), - V0Variances(-8.948491452600415, -30.948491452600415), - V0Variances(-9.745555749635251, -31.74555574963525), - V0Variances(-10.928045032645045, -32.928045032645045)}, - {V0Variances(-5.558562095836564, -28.558562095836564), - V0Variances(-5.486161892193216, -28.486161892193216), - V0Variances(-7.163697010897138, -30.16369701089714), - V0Variances(-7.997718751680708, -30.997718751680708), - V0Variances(-8.824140459442134, -31.824140459442134), - V0Variances(-10.447705621458134, -33.447705621458134), - V0Variances(-10.746667774832375, -33.746667774832375)}, - {V0Variances(-4.547157154382525, -28.547157154382525), - V0Variances(-6.216580824528357, -30.216580824528357), - V0Variances(-6.097210078262428, -30.097210078262428), - V0Variances(-7.873215469988139, -31.87321546998814), - V0Variances(-8.652284365210448, -32.65228436521045), - V0Variances(-9.889772718943199, -33.8897727189432), V0Variances(0, 0)}, - {V0Variances(-5.27365058987057, -30.27365058987057), - V0Variances(-5.164941786909992, -30.164941786909992), - V0Variances(-6.9233956231626195, -31.92339562316262), - V0Variances(-7.176178241418846, -32.17617824141885), - V0Variances(-8.490876045927656, -33.490876045927656), - V0Variances(-9.868458351961102, -34.86845835196111), V0Variances(0, 0)}, - {V0Variances(-4.228727281179324, -30.228727281179324), - V0Variances(-5.977329267849463, -31.977329267849463), - V0Variances(-6.232917187263332, -32.232917187263325), - V0Variances(-7.541325149103926, -33.541325149103926), - V0Variances(-8.245745876084897, -34.2457458760849), V0Variances(0, 0), - V0Variances(0, 0)}, - {V0Variances(-4.068209399541431, -31.06820939954143), - V0Variances(-5.284289051019407, -32.2842890510194), - V0Variances(-6.59434596780909, -33.59434596780909), - V0Variances(-7.306650734407292, -34.30665073440729), - V0Variances(-8.356857016386314, -35.356857016386314), V0Variances(0, 0), - V0Variances(0, 0)}, - {V0Variances(-4.349247565658466, -32.349247565658466), - V0Variances(-5.646300370431092, -33.64630037043109), - V0Variances(-6.361052340155609, -34.36105234015561), - V0Variances(-7.4131805240628665, -35.41318052406287), V0Variances(0, 0), - V0Variances(0, 0), V0Variances(0, 0)}, - {V0Variances(-4.705134752586538, -33.70513475258654), - V0Variances(-5.412109448981951, -34.41210944898195), - V0Variances(-6.465578619068673, -35.46557861906867), V0Variances(0, 0), - V0Variances(0, 0), V0Variances(0, 0), V0Variances(0, 0)}, - {V0Variances(-4.488254390500785, -34.488254390500785), - V0Variances(-5.062180911982956, -35.062180911982956), V0Variances(0, 0), - V0Variances(0, 0), V0Variances(0, 0), V0Variances(0, 0), - V0Variances(0, 0)}, - {V0Variances(-4.085183120157467, -35.08518312015747), V0Variances(0, 0), - V0Variances(0, 0), V0Variances(0, 0), V0Variances(0, 0), V0Variances(0, 0), - V0Variances(0, 0)}}; - -extern "C" V0Parameter *get_parameters(int norm, int p) { - // - 1 is an offset as norm and p are in [1, ...] and not [0, ...] - return ¶meters[norm - 1][p - 1]; -} - -extern "C" V0Variances *get_variances(int norm, int p) { - // - 1 is an offset as norm and p are in [1, ...] and not [0, ...] - return &variances[norm - 1][p - 1]; -} diff --git a/backends/concrete-cuda/implementation/src/boolean_gates.cu b/backends/concrete-cuda/implementation/src/boolean_gates.cu index 55f2d0e7c..9d44c04ee 100644 --- a/backends/concrete-cuda/implementation/src/boolean_gates.cu +++ b/backends/concrete-cuda/implementation/src/boolean_gates.cu @@ -6,8 +6,8 @@ #include "keyswitch.h" #include "linear_algebra.h" -constexpr int PLAINTEXT_TRUE{1 << (32 - 3)}; -constexpr int PLAINTEXT_FALSE{7 << (32 - 3)}; +constexpr uint32_t PLAINTEXT_TRUE{1 << (32 - 3)}; +constexpr uint32_t PLAINTEXT_FALSE{static_cast(7 << (32 - 3))}; extern "C" void cuda_boolean_not_32(void *v_stream, uint32_t gpu_index, void *lwe_array_out, void *lwe_array_in, diff --git a/backends/concrete-cuda/implementation/src/device.cu b/backends/concrete-cuda/implementation/src/device.cu index 976893481..667e01607 100644 --- a/backends/concrete-cuda/implementation/src/device.cu +++ b/backends/concrete-cuda/implementation/src/device.cu @@ -1,6 +1,5 @@ #include "device.h" #include -#include #include /// Unsafe function to create a CUDA stream, must check first that GPU exists @@ -37,6 +36,9 @@ void *cuda_malloc_async(uint64_t size, cudaStream_t *stream, cudaSetDevice(gpu_index); void *ptr; +#ifndef CUDART_VERSION +#error CUDART_VERSION Undefined! +#elif (CUDART_VERSION >= 11020) int support_async_alloc; check_cuda_error(cudaDeviceGetAttribute( &support_async_alloc, cudaDevAttrMemoryPoolsSupported, gpu_index)); @@ -46,6 +48,9 @@ void *cuda_malloc_async(uint64_t size, cudaStream_t *stream, } else { check_cuda_error(cudaMalloc((void **)&ptr, size)); } +#else + check_cuda_error(cudaMalloc((void **)&ptr, size)); +#endif return ptr; } @@ -184,6 +189,9 @@ int cuda_drop(void *ptr, uint32_t gpu_index) { int cuda_drop_async(void *ptr, cudaStream_t *stream, uint32_t gpu_index) { cudaSetDevice(gpu_index); +#ifndef CUDART_VERSION +#error CUDART_VERSION Undefined! +#elif (CUDART_VERSION >= 11020) int support_async_alloc; check_cuda_error(cudaDeviceGetAttribute( &support_async_alloc, cudaDevAttrMemoryPoolsSupported, gpu_index)); @@ -193,6 +201,9 @@ int cuda_drop_async(void *ptr, cudaStream_t *stream, uint32_t gpu_index) { } else { check_cuda_error(cudaFree(ptr)); } +#else + check_cuda_error(cudaFree(ptr)); +#endif return 0; } diff --git a/backends/concrete-cuda/implementation/test/test_bit_extraction.cpp b/backends/concrete-cuda/implementation/test/test_bit_extraction.cpp index 6762ca1e3..902953d63 100644 --- a/backends/concrete-cuda/implementation/test/test_bit_extraction.cpp +++ b/backends/concrete-cuda/implementation/test/test_bit_extraction.cpp @@ -7,8 +7,8 @@ #include #include -const unsigned REPETITIONS = 5; -const unsigned SAMPLES = 100; +const unsigned REPETITIONS = 2; +const unsigned SAMPLES = 50; typedef struct { int lwe_dimension; @@ -90,8 +90,10 @@ public: input_lwe_dimension = glwe_dimension * polynomial_size; output_lwe_dimension = lwe_dimension; // Generate the keys - generate_lwe_secret_keys(&lwe_sk_in_array, input_lwe_dimension, csprng, REPETITIONS); - generate_lwe_secret_keys(&lwe_sk_out_array, output_lwe_dimension, csprng, REPETITIONS); + generate_lwe_secret_keys(&lwe_sk_in_array, input_lwe_dimension, csprng, + REPETITIONS); + generate_lwe_secret_keys(&lwe_sk_out_array, output_lwe_dimension, csprng, + REPETITIONS); generate_lwe_keyswitch_keys( stream, gpu_index, &d_ksk_array, lwe_sk_in_array, lwe_sk_out_array, input_lwe_dimension, output_lwe_dimension, ks_level, ks_base_log, @@ -100,8 +102,9 @@ public: stream, gpu_index, &d_fourier_bsk_array, lwe_sk_out_array, lwe_sk_in_array, output_lwe_dimension, glwe_dimension, polynomial_size, pbs_level, pbs_base_log, csprng, glwe_modular_variance, REPETITIONS); - plaintexts = generate_plaintexts( - number_of_bits_of_message_including_padding, delta, number_of_inputs, REPETITIONS, SAMPLES); + plaintexts = + generate_plaintexts(number_of_bits_of_message_including_padding, delta, + number_of_inputs, REPETITIONS, SAMPLES); d_lwe_out_ct_array = (uint64_t *)cuda_malloc_async( (output_lwe_dimension + 1) * number_of_bits_to_extract * @@ -148,29 +151,27 @@ TEST_P(BitExtractionTestPrimitives_u64, bit_extraction) { void *v_stream = (void *)stream; int bsk_size = (glwe_dimension + 1) * (glwe_dimension + 1) * pbs_level * polynomial_size * (output_lwe_dimension + 1); - int ksk_size = - ks_level * input_lwe_dimension * (output_lwe_dimension + 1); + int ksk_size = ks_level * input_lwe_dimension * (output_lwe_dimension + 1); for (uint r = 0; r < REPETITIONS; r++) { double *d_fourier_bsk = d_fourier_bsk_array + (ptrdiff_t)(bsk_size * r); uint64_t *d_ksk = d_ksk_array + (ptrdiff_t)(ksk_size * r); uint64_t *lwe_in_sk = lwe_sk_in_array + (ptrdiff_t)(input_lwe_dimension * r); - uint64_t *lwe_sk_out = lwe_sk_out_array + (ptrdiff_t)(r * output_lwe_dimension); + uint64_t *lwe_sk_out = + lwe_sk_out_array + (ptrdiff_t)(r * output_lwe_dimension); for (uint s = 0; s < SAMPLES; s++) { for (int i = 0; i < number_of_inputs; i++) { uint64_t plaintext = plaintexts[r * SAMPLES * number_of_inputs + s * number_of_inputs + i]; uint64_t *lwe_in_ct = - lwe_in_ct_array + - (ptrdiff_t)( - i * (input_lwe_dimension + 1)); + lwe_in_ct_array + (ptrdiff_t)(i * (input_lwe_dimension + 1)); concrete_cpu_encrypt_lwe_ciphertext_u64( lwe_in_sk, lwe_in_ct, plaintext, input_lwe_dimension, lwe_modular_variance, csprng, &CONCRETE_CSPRNG_VTABLE); } cuda_memcpy_async_to_gpu(d_lwe_in_ct_array, lwe_in_ct_array, - (input_lwe_dimension + 1) * - number_of_inputs * sizeof(uint64_t), + (input_lwe_dimension + 1) * number_of_inputs * + sizeof(uint64_t), stream, gpu_index); // Execute bit extract @@ -184,14 +185,15 @@ TEST_P(BitExtractionTestPrimitives_u64, bit_extraction) { // Copy result back cuda_memcpy_async_to_cpu(lwe_out_ct_array, d_lwe_out_ct_array, - (output_lwe_dimension + 1) * number_of_bits_to_extract * + (output_lwe_dimension + 1) * + number_of_bits_to_extract * number_of_inputs * sizeof(uint64_t), stream, gpu_index); cuda_synchronize_stream(v_stream); for (int j = 0; j < number_of_inputs; j++) { uint64_t *result_array = - lwe_out_ct_array + - (ptrdiff_t)(j * number_of_bits_to_extract * (output_lwe_dimension + 1)); + lwe_out_ct_array + (ptrdiff_t)(j * number_of_bits_to_extract * + (output_lwe_dimension + 1)); uint64_t plaintext = plaintexts[r * SAMPLES * number_of_inputs + s * number_of_inputs + j]; for (int i = 0; i < number_of_bits_to_extract; i++) { @@ -245,4 +247,4 @@ printParamName(::testing::TestParamInfo p) { INSTANTIATE_TEST_CASE_P(BitExtractionInstantiation, BitExtractionTestPrimitives_u64, bit_extract_params_u64, - printParamName); \ No newline at end of file + printParamName); diff --git a/backends/concrete-cuda/implementation/test/test_bootstrap.cpp b/backends/concrete-cuda/implementation/test/test_bootstrap.cpp index cb6e744d8..f9000de65 100644 --- a/backends/concrete-cuda/implementation/test/test_bootstrap.cpp +++ b/backends/concrete-cuda/implementation/test/test_bootstrap.cpp @@ -8,9 +8,6 @@ #include #include -const unsigned REPETITIONS = 5; -const unsigned SAMPLES = 100; - typedef struct { int lwe_dimension; int glwe_dimension; @@ -22,6 +19,8 @@ typedef struct { int message_modulus; int carry_modulus; int number_of_inputs; + int repetitions; + int samples; } BootstrapTestParams; class BootstrapTestPrimitives_u64 @@ -38,6 +37,8 @@ protected: int carry_modulus; int payload_modulus; int number_of_inputs; + int repetitions; + int samples; uint64_t delta; Csprng *csprng; cudaStream_t *stream; @@ -68,6 +69,8 @@ public: message_modulus = (int)GetParam().message_modulus; carry_modulus = (int)GetParam().carry_modulus; number_of_inputs = (int)GetParam().number_of_inputs; + repetitions = (int)GetParam().repetitions; + samples = (int)GetParam().samples; payload_modulus = message_modulus * carry_modulus; // Value of the shift we multiply our messages by @@ -81,15 +84,17 @@ public: csprng, Uint128{.little_endian_bytes = {*seed}}); // Generate the keys - generate_lwe_secret_keys(&lwe_sk_in_array, lwe_dimension, csprng, REPETITIONS); + generate_lwe_secret_keys(&lwe_sk_in_array, lwe_dimension, csprng, + repetitions); generate_lwe_secret_keys(&lwe_sk_out_array, - glwe_dimension * polynomial_size, csprng, REPETITIONS); + glwe_dimension * polynomial_size, csprng, + repetitions); generate_lwe_bootstrap_keys( stream, gpu_index, &d_fourier_bsk_array, lwe_sk_in_array, lwe_sk_out_array, lwe_dimension, glwe_dimension, polynomial_size, - pbs_level, pbs_base_log, csprng, glwe_modular_variance, REPETITIONS); - plaintexts = generate_plaintexts(payload_modulus, delta, number_of_inputs, REPETITIONS, - SAMPLES); + pbs_level, pbs_base_log, csprng, glwe_modular_variance, repetitions); + plaintexts = generate_plaintexts(payload_modulus, delta, number_of_inputs, + repetitions, samples); // Create the LUT uint64_t *lut_pbs_identity = generate_identity_lut_pbs( @@ -117,21 +122,21 @@ public: number_of_inputs * sizeof(uint64_t), stream, gpu_index); d_lwe_ct_in_array = (uint64_t *)cuda_malloc_async( - (lwe_dimension + 1) * number_of_inputs * REPETITIONS * SAMPLES * + (lwe_dimension + 1) * number_of_inputs * repetitions * samples * sizeof(uint64_t), stream, gpu_index); uint64_t *lwe_ct_in_array = (uint64_t *)malloc((lwe_dimension + 1) * number_of_inputs * - REPETITIONS * SAMPLES * sizeof(uint64_t)); + repetitions * samples * sizeof(uint64_t)); // Create the input/output ciphertexts - for (uint r = 0; r < REPETITIONS; r++) { + for (int r = 0; r < repetitions; r++) { uint64_t *lwe_sk_in = lwe_sk_in_array + (ptrdiff_t)(r * lwe_dimension); - for (uint s = 0; s < SAMPLES; s++) { + for (int s = 0; s < samples; s++) { for (int i = 0; i < number_of_inputs; i++) { - uint64_t plaintext = plaintexts[r * SAMPLES * number_of_inputs + + uint64_t plaintext = plaintexts[r * samples * number_of_inputs + s * number_of_inputs + i]; uint64_t *lwe_ct_in = - lwe_ct_in_array + (ptrdiff_t)((r * SAMPLES * number_of_inputs + + lwe_ct_in_array + (ptrdiff_t)((r * samples * number_of_inputs + s * number_of_inputs + i) * (lwe_dimension + 1)); concrete_cpu_encrypt_lwe_ciphertext_u64( @@ -142,7 +147,7 @@ public: } cuda_synchronize_stream(v_stream); cuda_memcpy_async_to_gpu(d_lwe_ct_in_array, lwe_ct_in_array, - REPETITIONS * SAMPLES * number_of_inputs * + repetitions * samples * number_of_inputs * (lwe_dimension + 1) * sizeof(uint64_t), stream, gpu_index); free(lwe_ct_in_array); @@ -174,14 +179,14 @@ TEST_P(BootstrapTestPrimitives_u64, amortized_bootstrap) { int bsk_size = (glwe_dimension + 1) * (glwe_dimension + 1) * pbs_level * polynomial_size * (lwe_dimension + 1); // Here execute the PBS - for (uint r = 0; r < REPETITIONS; r++) { + for (int r = 0; r < repetitions; r++) { double *d_fourier_bsk = d_fourier_bsk_array + (ptrdiff_t)(bsk_size * r); uint64_t *lwe_sk_out = lwe_sk_out_array + (ptrdiff_t)(r * glwe_dimension * polynomial_size); - for (uint s = 0; s < SAMPLES; s++) { + for (int s = 0; s < samples; s++) { uint64_t *d_lwe_ct_in = d_lwe_ct_in_array + - (ptrdiff_t)((r * SAMPLES * number_of_inputs + s * number_of_inputs) * + (ptrdiff_t)((r * samples * number_of_inputs + s * number_of_inputs) * (lwe_dimension + 1)); // Execute PBS cuda_bootstrap_amortized_lwe_ciphertext_vector_64( @@ -200,7 +205,7 @@ TEST_P(BootstrapTestPrimitives_u64, amortized_bootstrap) { uint64_t *result = lwe_ct_out_array + (ptrdiff_t)(j * (glwe_dimension * polynomial_size + 1)); - uint64_t plaintext = plaintexts[r * SAMPLES * number_of_inputs + + uint64_t plaintext = plaintexts[r * samples * number_of_inputs + s * number_of_inputs + j]; uint64_t decrypted = 0; concrete_cpu_decrypt_lwe_ciphertext_u64( @@ -227,8 +232,8 @@ TEST_P(BootstrapTestPrimitives_u64, amortized_bootstrap) { TEST_P(BootstrapTestPrimitives_u64, low_latency_bootstrap) { int number_of_sm = 0; cudaDeviceGetAttribute(&number_of_sm, cudaDevAttrMultiProcessorCount, 0); - if(number_of_inputs > number_of_sm * 4 / (glwe_dimension + 1) / pbs_level) - GTEST_SKIP() << "The Low Latency PBS does not support this configuration"; + if (number_of_inputs > number_of_sm * 4 / (glwe_dimension + 1) / pbs_level) + GTEST_SKIP() << "The Low Latency PBS does not support this configuration"; uint64_t *lwe_ct_out_array = (uint64_t *)malloc((glwe_dimension * polynomial_size + 1) * number_of_inputs * sizeof(uint64_t)); @@ -239,14 +244,14 @@ TEST_P(BootstrapTestPrimitives_u64, low_latency_bootstrap) { int bsk_size = (glwe_dimension + 1) * (glwe_dimension + 1) * pbs_level * polynomial_size * (lwe_dimension + 1); // Here execute the PBS - for (uint r = 0; r < REPETITIONS; r++) { + for (int r = 0; r < repetitions; r++) { double *d_fourier_bsk = d_fourier_bsk_array + (ptrdiff_t)(bsk_size * r); uint64_t *lwe_sk_out = lwe_sk_out_array + (ptrdiff_t)(r * glwe_dimension * polynomial_size); - for (uint s = 0; s < SAMPLES; s++) { + for (int s = 0; s < samples; s++) { uint64_t *d_lwe_ct_in = d_lwe_ct_in_array + - (ptrdiff_t)((r * SAMPLES * number_of_inputs + s * number_of_inputs) * + (ptrdiff_t)((r * samples * number_of_inputs + s * number_of_inputs) * (lwe_dimension + 1)); // Execute PBS cuda_bootstrap_low_latency_lwe_ciphertext_vector_64( @@ -265,7 +270,7 @@ TEST_P(BootstrapTestPrimitives_u64, low_latency_bootstrap) { uint64_t *result = lwe_ct_out_array + (ptrdiff_t)(j * (glwe_dimension * polynomial_size + 1)); - uint64_t plaintext = plaintexts[r * SAMPLES * number_of_inputs + + uint64_t plaintext = plaintexts[r * samples * number_of_inputs + s * number_of_inputs + j]; uint64_t decrypted = 0; concrete_cpu_decrypt_lwe_ciphertext_u64( @@ -294,56 +299,28 @@ TEST_P(BootstrapTestPrimitives_u64, low_latency_bootstrap) { ::testing::internal::ParamGenerator pbs_params_u64 = ::testing::Values( // n, k, N, lwe_variance, glwe_variance, pbs_base_log, pbs_level, - // message_modulus, carry_modulus, number_of_inputs - // 1 bit message 0 bit carry parameters + // message_modulus, carry_modulus, number_of_inputs, repetitions, + // samples (BootstrapTestParams){567, 5, 256, 0.000007069849454709433, 0.00000000000000029403601535432533, 15, 1, 2, 1, - 1}, - (BootstrapTestParams){567, 5, 256, 0.000007069849454709433, - 0.00000000000000029403601535432533, 15, 1, 2, 1, - 10}, - // 2 bit message 3 bit carry parameters - (BootstrapTestParams){623, 6, 256, 0.000007069849454709433, - 0.00000000000000029403601535432533, 9, 3, 3, 4, - 1}, - (BootstrapTestParams){623, 6, 256, 0.000007069849454709433, - 0.00000000000000029403601535432533, 9, 3, 3, 4, - 10}, - // 3 bits message 0 bit carry parameters + 5, 2, 50}, + (BootstrapTestParams){623, 6, 256, 7.52316384526264e-37, + 7.52316384526264e-37, 9, 3, 2, 2, 5, 2, 50}, (BootstrapTestParams){694, 3, 512, 0.000007069849454709433, - 0.00000000000000029403601535432533, 18, 1, 4, 1, - 1}, - (BootstrapTestParams){694, 3, 512, 0.000007069849454709433, - 0.00000000000000029403601535432533, 18, 1, 4, 1, - 10}, - // 4 bits message 0 bit carry parameters + 0.00000000000000029403601535432533, 18, 1, 2, 1, + 5, 2, 50}, (BootstrapTestParams){769, 2, 1024, 0.000007069849454709433, - 0.00000000000000029403601535432533, 23, 1, 5, 1, - 1}, - (BootstrapTestParams){769, 2, 1024, 0.000007069849454709433, - 0.00000000000000029403601535432533, 23, 1, 5, 1, - 10}, - // 5 bits message 0 bit carry parameters + 0.00000000000000029403601535432533, 23, 1, 2, 1, + 5, 2, 50}, (BootstrapTestParams){754, 1, 2048, 0.000007069849454709433, - 0.00000000000000029403601535432533, 23, 1, 6, 1, - 1}, - (BootstrapTestParams){754, 1, 2048, 0.000007069849454709433, - 0.00000000000000029403601535432533, 23, 1, 6, 1, - 10}, - // 6 bits message 0 bit carry parameters + 0.00000000000000029403601535432533, 23, 1, 4, 1, + 5, 2, 50}, (BootstrapTestParams){847, 1, 4096, 0.000007069849454709433, - 0.00000000000000029403601535432533, 1, 22, 7, 1, - 1}, - (BootstrapTestParams){847, 1, 4096, 0.000007069849454709433, - 0.00000000000000029403601535432533, 1, 22, 7, 1, - 10}, - // 7 bits message 0 bit carry parameters + 0.00000000000000029403601535432533, 2, 12, 2, 1, + 2, 1, 50}, (BootstrapTestParams){881, 1, 8192, 0.000007069849454709433, - 0.00000000000000029403601535432533, 1, 22, 8, 1, - 1}, - (BootstrapTestParams){881, 1, 8192, 0.000007069849454709433, - 0.00000000000000029403601535432533, 1, 22, 8, 1, - 2}); + 0.00000000000000029403601535432533, 22, 1, 2, 1, + 2, 1, 25}); std::string printParamName(::testing::TestParamInfo p) { BootstrapTestParams params = p.param; @@ -357,4 +334,4 @@ std::string printParamName(::testing::TestParamInfo p) { } INSTANTIATE_TEST_CASE_P(BootstrapInstantiation, BootstrapTestPrimitives_u64, - pbs_params_u64, printParamName); \ No newline at end of file + pbs_params_u64, printParamName); diff --git a/backends/concrete-cuda/implementation/test/test_circuit_bootstrap.cpp b/backends/concrete-cuda/implementation/test/test_circuit_bootstrap.cpp index 2b83b7c58..4592bada6 100644 --- a/backends/concrete-cuda/implementation/test/test_circuit_bootstrap.cpp +++ b/backends/concrete-cuda/implementation/test/test_circuit_bootstrap.cpp @@ -7,8 +7,8 @@ #include #include -const unsigned REPETITIONS = 5; -const unsigned SAMPLES = 100; +const unsigned REPETITIONS = 2; +const unsigned SAMPLES = 50; typedef struct { int lwe_dimension; diff --git a/backends/concrete-cuda/implementation/test/test_cmux_tree.cpp b/backends/concrete-cuda/implementation/test/test_cmux_tree.cpp index 3b0a8359a..192d568a2 100644 --- a/backends/concrete-cuda/implementation/test/test_cmux_tree.cpp +++ b/backends/concrete-cuda/implementation/test/test_cmux_tree.cpp @@ -5,11 +5,10 @@ #include "gtest/gtest.h" #include #include -#include #include const unsigned REPETITIONS = 5; -const unsigned SAMPLES = 100; +const unsigned SAMPLES = 50; typedef struct { int glwe_dimension; diff --git a/backends/concrete-cuda/implementation/test/test_keyswitch.cpp b/backends/concrete-cuda/implementation/test/test_keyswitch.cpp index a4770a297..4d555e24d 100644 --- a/backends/concrete-cuda/implementation/test/test_keyswitch.cpp +++ b/backends/concrete-cuda/implementation/test/test_keyswitch.cpp @@ -8,8 +8,8 @@ #include #include -const unsigned REPETITIONS = 5; -const unsigned SAMPLES = 100; +const unsigned REPETITIONS = 2; +const unsigned SAMPLES = 50; typedef struct { int input_lwe_dimension; @@ -169,18 +169,12 @@ TEST_P(KeyswitchTestPrimitives_u64, keyswitch) { ::testing::Values( // n, k*N, noise_variance, ks_base_log, ks_level, // message_modulus, carry_modulus - // 1 bit message 0 bit carry parameters (KeyswitchTestParams){567, 1280, 2.9802322387695312e-08, 3, 3, 2, 1}, - // 3 bits message 0 bit carry parameters - (KeyswitchTestParams){694, 1536, 2.9802322387695312e-08, 4, 3, 4, 1}, - // 4 bits message 0 bit carry parameters - (KeyswitchTestParams){769, 2048, 2.9802322387695312e-08, 4, 3, 5, 1}, - // 5 bits message 0 bit carry parameters - (KeyswitchTestParams){754, 2048, 2.9802322387695312e-08, 3, 5, 6, 1}, - // 6 bits message 0 bit carry parameters - (KeyswitchTestParams){847, 4096, 2.9802322387695312e-08, 4, 4, 7, 1}, - // 7 bits message 0 bit carry parameters - (KeyswitchTestParams){881, 8192, 2.9802322387695312e-08, 3, 6, 8, 1}); + (KeyswitchTestParams){694, 1536, 2.9802322387695312e-08, 4, 3, 2, 1}, + (KeyswitchTestParams){769, 2048, 2.9802322387695312e-08, 4, 3, 2, 1}, + (KeyswitchTestParams){754, 2048, 2.9802322387695312e-08, 3, 5, 2, 1}, + (KeyswitchTestParams){847, 4096, 2.9802322387695312e-08, 4, 4, 2, 1}, + (KeyswitchTestParams){881, 8192, 2.9802322387695312e-08, 3, 6, 2, 1}); std::string printParamName(::testing::TestParamInfo p) { KeyswitchTestParams params = p.param; diff --git a/backends/concrete-cuda/implementation/test/test_wop_bootstrap.cpp b/backends/concrete-cuda/implementation/test/test_wop_bootstrap.cpp index 57212623f..494c0f3b6 100644 --- a/backends/concrete-cuda/implementation/test/test_wop_bootstrap.cpp +++ b/backends/concrete-cuda/implementation/test/test_wop_bootstrap.cpp @@ -8,7 +8,7 @@ #include #include -const unsigned REPETITIONS = 5; +const unsigned REPETITIONS = 2; const unsigned SAMPLES = 10; typedef struct { @@ -102,12 +102,14 @@ public: input_lwe_dimension = glwe_dimension * polynomial_size; // Generate the keys - generate_lwe_secret_keys(&lwe_sk_in_array, input_lwe_dimension, csprng, REPETITIONS); - generate_lwe_secret_keys(&lwe_sk_out_array, lwe_dimension, csprng, REPETITIONS); - generate_lwe_keyswitch_keys(stream, gpu_index, &d_ksk_array, - lwe_sk_in_array, lwe_sk_out_array, - input_lwe_dimension, lwe_dimension, ks_level, - ks_base_log, csprng, lwe_modular_variance, REPETITIONS); + generate_lwe_secret_keys(&lwe_sk_in_array, input_lwe_dimension, csprng, + REPETITIONS); + generate_lwe_secret_keys(&lwe_sk_out_array, lwe_dimension, csprng, + REPETITIONS); + generate_lwe_keyswitch_keys( + stream, gpu_index, &d_ksk_array, lwe_sk_in_array, lwe_sk_out_array, + input_lwe_dimension, lwe_dimension, ks_level, ks_base_log, csprng, + lwe_modular_variance, REPETITIONS); generate_lwe_bootstrap_keys( stream, gpu_index, &d_fourier_bsk_array, lwe_sk_out_array, lwe_sk_in_array, lwe_dimension, glwe_dimension, polynomial_size, @@ -223,11 +225,13 @@ TEST_P(WopBootstrapTestPrimitives_u64, wop_pbs) { cuda_get_max_shared_memory(gpu_index)); //// Copy result back - cuda_memcpy_async_to_cpu(lwe_out_ct_array, d_lwe_ct_out_array, - (input_lwe_dimension + 1) * tau * sizeof(uint64_t), stream, gpu_index); - cuda_synchronize_stream(v_stream); + cuda_memcpy_async_to_cpu(lwe_out_ct_array, d_lwe_ct_out_array, + (input_lwe_dimension + 1) * tau * + sizeof(uint64_t), + stream, gpu_index); + cuda_synchronize_stream(v_stream); - for (int i = 0; i < tau; i++) { + for (int i = 0; i < tau; i++) { uint64_t plaintext = plaintexts[r * SAMPLES * tau + s * tau + i]; uint64_t *result_ct = lwe_out_ct_array + (ptrdiff_t)(i * (input_lwe_dimension + 1)); @@ -251,17 +255,19 @@ TEST_P(WopBootstrapTestPrimitives_u64, wop_pbs) { // n, k, N, lwe_variance, glwe_variance, pbs_base_log, pbs_level, // ks_base_log, ks_level, tau (WopBootstrapTestParams){481, 2, 512, 7.52316384526264e-37, - 7.52316384526264e-37, 4, - 9, 1, 9, 4, 9, 6, 4, 1} -// (WopBootstrapTestParams){481, 2, 512, 7.52316384526264e-37, -// 7.52316384526264e-37, 4, 9, 1, 9, 4, 9, 6, 4, -// 2} , -// (WopBootstrapTestParams){481, 2, 1024, 7.52316384526264e-37, -// 7.52316384526264e-37, 4, -// 9, 1, 9, 4, 9, 6, 4, 1}, -// (WopBootstrapTestParams){481, 2, 1024, 7.52316384526264e-37, -// 7.52316384526264e-37, 4, -// 9, 1, 9, 4, 9, 6, 4, 2} + 7.52316384526264e-37, 4, 9, 1, 9, 4, 9, 6, 4, + 1} + // (WopBootstrapTestParams){481, 2, 512, 7.52316384526264e-37, + // 7.52316384526264e-37, 4, 9, 1, 9, 4, + // 9, 6, 4, 2} , + // (WopBootstrapTestParams){481, 2, 1024, 7.52316384526264e-37, + // 7.52316384526264e-37, + // 4, 9, 1, 9, 4, 9, + // 6, 4, 1}, + // (WopBootstrapTestParams){481, 2, 1024, 7.52316384526264e-37, + // 7.52316384526264e-37, + // 4, 9, 1, 9, 4, 9, + // 6, 4, 2} ); std::string printParamName(::testing::TestParamInfo p) { diff --git a/backends/concrete-cuda/implementation/test/utils.cpp b/backends/concrete-cuda/implementation/test/utils.cpp index 3dff50b3b..4ff0036c3 100644 --- a/backends/concrete-cuda/implementation/test/utils.cpp +++ b/backends/concrete-cuda/implementation/test/utils.cpp @@ -12,8 +12,8 @@ // The payload_modulus is the message modulus times the carry modulus // (so the total message modulus) uint64_t *generate_plaintexts(uint64_t payload_modulus, uint64_t delta, - int number_of_inputs, const unsigned repetitions, const unsigned - samples) { + int number_of_inputs, const unsigned repetitions, + const unsigned samples) { uint64_t *plaintext_array = (uint64_t *)malloc( repetitions * samples * number_of_inputs * sizeof(uint64_t)); @@ -121,7 +121,8 @@ uint64_t *generate_identity_lut_cmux_tree(int polynomial_size, int num_lut, // Generate repetitions LWE secret keys void generate_lwe_secret_keys(uint64_t **lwe_sk_array, int lwe_dimension, Csprng *csprng, const unsigned repetitions) { - *lwe_sk_array = (uint64_t *)malloc(lwe_dimension * repetitions * sizeof(uint64_t)); + *lwe_sk_array = + (uint64_t *)malloc(lwe_dimension * repetitions * sizeof(uint64_t)); int shift = 0; for (uint r = 0; r < repetitions; r++) { // Generate the lwe secret key for each repetition @@ -134,7 +135,8 @@ void generate_lwe_secret_keys(uint64_t **lwe_sk_array, int lwe_dimension, // Generate repetitions GLWE secret keys void generate_glwe_secret_keys(uint64_t **glwe_sk_array, int glwe_dimension, - int polynomial_size, Csprng *csprng, const unsigned repetitions) { + int polynomial_size, Csprng *csprng, + const unsigned repetitions) { int glwe_sk_array_size = glwe_dimension * polynomial_size * repetitions; *glwe_sk_array = (uint64_t *)malloc(glwe_sk_array_size * sizeof(uint64_t)); int shift = 0; @@ -148,13 +150,11 @@ void generate_glwe_secret_keys(uint64_t **glwe_sk_array, int glwe_dimension, } // Generate repetitions LWE bootstrap keys -void generate_lwe_bootstrap_keys(cudaStream_t *stream, int gpu_index, - double **d_fourier_bsk_array, - uint64_t *lwe_sk_in_array, - uint64_t *lwe_sk_out_array, int lwe_dimension, - int glwe_dimension, int polynomial_size, - int pbs_level, int pbs_base_log, - Csprng *csprng, double variance, const unsigned repetitions) { +void generate_lwe_bootstrap_keys( + cudaStream_t *stream, int gpu_index, double **d_fourier_bsk_array, + uint64_t *lwe_sk_in_array, uint64_t *lwe_sk_out_array, int lwe_dimension, + int glwe_dimension, int polynomial_size, int pbs_level, int pbs_base_log, + Csprng *csprng, double variance, const unsigned repetitions) { void *v_stream = (void *)stream; int bsk_size = (glwe_dimension + 1) * (glwe_dimension + 1) * pbs_level * polynomial_size * (lwe_dimension + 1); @@ -190,11 +190,14 @@ void generate_lwe_bootstrap_keys(cudaStream_t *stream, int gpu_index, } // Generate repetitions keyswitch keys -void generate_lwe_keyswitch_keys( - cudaStream_t *stream, int gpu_index, uint64_t **d_ksk_array, - uint64_t *lwe_sk_in_array, uint64_t *lwe_sk_out_array, - int input_lwe_dimension, int output_lwe_dimension, int ksk_level, - int ksk_base_log, Csprng *csprng, double variance, const unsigned repetitions) { +void generate_lwe_keyswitch_keys(cudaStream_t *stream, int gpu_index, + uint64_t **d_ksk_array, + uint64_t *lwe_sk_in_array, + uint64_t *lwe_sk_out_array, + int input_lwe_dimension, + int output_lwe_dimension, int ksk_level, + int ksk_base_log, Csprng *csprng, + double variance, const unsigned repetitions) { int ksk_size = ksk_level * (output_lwe_dimension + 1) * input_lwe_dimension; int ksk_array_size = ksk_size * repetitions;