github/concrete
1
1
Fork 0
mirror of https://github.com/zama-ai/concrete.git synced 2026-02-08 11:35:02 -05:00
Code Issues Packages Projects Releases Wiki Activity

chore(concrete_cuda): fix tests, reduce test time

Browse Source 
- 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
This commit is contained in:
Agnes Leroy
2023-03-15 10:55:08 +01:00
committed by Agnès Leroy
parent 4d82d177ab
commit 629e22aaac
14 changed files with 142 additions and 549 deletions
Download Patch File Download Diff File Expand all files Collapse all files

16
.github/workflows/concrete_cuda_test.yml vendored
View File

@@ -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

3
backends/concrete-cuda/implementation/CMakeLists.txt
View File

@@ -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)

6
backends/concrete-cuda/implementation/include/device.h
View File

@@ -4,10 +4,10 @@
#pragma once
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
extern "C" {
cudaStream_t *cuda_create_stream(uint32_t gpu_index);

4
backends/concrete-cuda/implementation/parameters/CMakeLists.txt
View File

@@ -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)

380
backends/concrete-cuda/implementation/parameters/parameters.cpp
View File

@@ -1,380 +0,0 @@
#include <iostream>
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 &parameters[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];
}

4
backends/concrete-cuda/implementation/src/boolean_gates.cu
View File

@@ -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<uint32_t>(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,

13
backends/concrete-cuda/implementation/src/device.cu
View File

@@ -1,6 +1,5 @@
#include "device.h"
#include <cstdint>
#include <cstdio>
#include <cuda_runtime.h>
/// 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;
}

38
backends/concrete-cuda/implementation/test/test_bit_extraction.cpp
View File

@@ -7,8 +7,8 @@
#include <stdio.h>
#include <stdlib.h>
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<BitExtractionTestParams> p) {
INSTANTIATE_TEST_CASE_P(BitExtractionInstantiation,
BitExtractionTestPrimitives_u64, bit_extract_params_u64,
printParamName);
printParamName);

115
backends/concrete-cuda/implementation/test/test_bootstrap.cpp
View File

@@ -8,9 +8,6 @@
#include <stdio.h>
#include <stdlib.h>
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<BootstrapTestParams> 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<BootstrapTestParams> p) {
BootstrapTestParams params = p.param;
@@ -357,4 +334,4 @@ std::string printParamName(::testing::TestParamInfo<BootstrapTestParams> p) {
}
INSTANTIATE_TEST_CASE_P(BootstrapInstantiation, BootstrapTestPrimitives_u64,
pbs_params_u64, printParamName);
pbs_params_u64, printParamName);

4
backends/concrete-cuda/implementation/test/test_circuit_bootstrap.cpp
View File

@@ -7,8 +7,8 @@
#include <stdio.h>
#include <stdlib.h>
const unsigned REPETITIONS = 5;
const unsigned SAMPLES = 100;
const unsigned REPETITIONS = 2;
const unsigned SAMPLES = 50;
typedef struct {
int lwe_dimension;

3
backends/concrete-cuda/implementation/test/test_cmux_tree.cpp
View File

@@ -5,11 +5,10 @@
#include "gtest/gtest.h"
#include <cstdint>
#include <functional>
#include <stdio.h>
#include <stdlib.h>
const unsigned REPETITIONS = 5;
const unsigned SAMPLES = 100;
const unsigned SAMPLES = 50;
typedef struct {
int glwe_dimension;

20
backends/concrete-cuda/implementation/test/test_keyswitch.cpp
View File

@@ -8,8 +8,8 @@
#include <stdio.h>
#include <stdlib.h>
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<KeyswitchTestParams> p) {
KeyswitchTestParams params = p.param;

50
backends/concrete-cuda/implementation/test/test_wop_bootstrap.cpp
View File

@@ -8,7 +8,7 @@
#include <stdio.h>
#include <stdlib.h>
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<WopBootstrapTestParams> p) {

35
backends/concrete-cuda/implementation/test/utils.cpp
View File

@@ -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;