FSS_SSE_GPU

FSS_SSE_GPU 是一个通用的 CUDA 密码原语框架，用于帮助基于 FSS、RSS 的方案将底层计算迁移到 GPU。

当前仓库提供 DPF/DCF 批量求值、完整域 EvalAll、FSS key 序列化、RSS 本地算子、EvalPred、范围谓词组合和 GPU reduction。框架负责单台服务器上的 GPU 计算与设备内存，不实现方案的客户端、数据库、RPC 或跨服务器调度。

环境要求

• Linux 或 WSL2；
• 支持 CUDA 的 NVIDIA GPU；
• CUDA Toolkit 12.x 或 13.x；
• CMake 3.22 及以上；
• 支持 C++20/CUDA C++20 的主机编译器。当前验证使用 GCC 13。

项目所需 FSS 头文件已经放在 include/fss_sse_gpu/vendor，基础构建不需要下载其他第三方代码。

部署、构建与测试

推荐使用 Release 构建进行性能测试：

git clone git@github.com:YY404NF/FSS_SSE_GPU.git
cd FSS_SSE_GPU

cmake -S . -B build-release \
  -DCMAKE_BUILD_TYPE=Release \
  -DCMAKE_CXX_COMPILER=/usr/bin/g++-13 \
  -DCMAKE_CUDA_HOST_COMPILER=/usr/bin/g++-13 \
  -DCMAKE_CUDA_ARCHITECTURES=86

cmake --build build-release -j 4
ctest --test-dir build-release --output-on-failure

CMAKE_CUDA_ARCHITECTURES=86 对应 RTX 30 系列。其他 GPU 应改成对应的 CUDA compute capability，也可以省略该参数让 CMake 自动选择。

前三个 CTest 用例分别覆盖通用原语、三方 DPF/EvalPred 语义和 Ring32 范围判断；默认启用的第四个用例覆盖 Ring32 多谓词 RSS 的 CPU/GPU 逐层校验。下列性能程序需要单独运行：

./build-release/fss_sse_gpu_eval_all_benchmark
./build-release/fss_sse_gpu_eval_pred_benchmark
./build-release/fss_sse_gpu_rss_async_benchmark

分组传输测试模拟多个基础谓词的 RSS 结果：

./build-release/fss_sse_gpu_rss_group_transport_benchmark

每组包含一批 share，测试逐组处理与整批处理在 GPU 准备、D2H/H2D 和 share 组装上的差异。计时不包含真实网络等待。

Ring32 多谓词 RSS 连续执行测试（需要 OpenMP；可通过 -DFSS_SSE_GPU_BUILD_RSS_QUERY_TEST=OFF 关闭此测试目标）：

bash tests/run_rss_query.sh results/my-rss-query-run

脚本默认使用 CUDA 13.3、GCC 13 和 sm_86；如果本机版本不同，可以通过 FSS_CUDA_HOME、FSS_CXX_COMPILER、FSS_CUDA_HOST_COMPILER、 FSS_CUDA_ARCHITECTURES、FSS_CUDA_FLAGS 和 FSS_RSS_QUERY_BUILD_DIR 覆盖。比如当前 CUDA 13.1 环境可以运行：

FSS_CUDA_HOME=/usr/local/cuda-13.1 \
FSS_CXX_COMPILER=/usr/bin/g++ \
FSS_CUDA_HOST_COMPILER=/usr/bin/g++ \
FSS_CUDA_FLAGS=--allow-unsupported-compiler \
FSS_RSS_QUERY_BUILD_DIR="$PWD/build-20261017" \
bash tests/run_rss_query.sh results/my-rss-query-run

结果目录必须没有同名实验文件。脚本依次运行 CTest、CPU 单线程/8 线程与两条 GPU 路径、单独的阶段计时及 Compute Sanitizer，并保存原始 CSV、汇总、环境和源码哈希。 输入是预构造的 FilterPred 输出 share，测试覆盖批量 Reshare 和多层乘法合取； CPU/GPU 均使用 Ring32。同一进程模拟三方，GPU 使用一张卡；不包含 FSS 求值、 真实网络、PRF、聚合或完整 Waldo 查询。详见 测试说明。

primitives/rss_cpu.h 提供不依赖 CUDA 的 Ring32 CPU 参考接口： PrepareReshareOutgoing、PrepareMulOutgoing、AssembleSharePair。 CPU 和 GPU 共用 SharePair<T> 布局，默认类型仍为 uint64_t；新增 CPU 接口使用 SharePair<uint32_t>。OpenMP 只用于测试程序，公共参考头文件不依赖 OpenMP。

安装头文件和 CMake target：

cmake --install build-release --prefix "$PWD/install"

其他 CMake 项目可以直接使用源码目录：

add_subdirectory(path/to/FSS_SSE_GPU)
target_link_libraries(your_cuda_target PRIVATE fss_sse_gpu)

也可以使用安装后的 package：

find_package(FSS_SSE_GPU CONFIG REQUIRED)
target_link_libraries(your_cuda_target PRIVATE FSS_SSE_GPU::fss_sse_gpu)

目录结构

include/fss_sse_gpu/
  runtime/       GPU buffer、stream、错误检查和输出环
  primitives/    DPF、DCF、FSS key 序列化和 RSS 本地算子
  composition/   通用 EvalPred 和 GPU 聚合
  vendor/fss/    项目内部使用的 FSS 实现

tests/
  correctness.cu          基础原语正确性测试
  three_party_dpf.cu      三方 DPF Gen + EvalPred 测试和 DPF 性能测试
  range_predicate.cu      两个 DCF 前缀结果组合范围谓词的边界测试
  eval_pred_benchmark.cu  EvalPred CPU/GPU 性能对比测试
  eval_all_benchmark.cu   DCF 独立 Eval 与 EvalAll 性能对比

主要接口

模块	接口	作用
输出环	runtime::Ring32	32 位 FSS 输出环，可用于半诚实的值和 filter share
输出环	runtime::Ring64	默认 64 位输出环
认证数据	runtime::AuthenticatedShare	模 2^128 的本地 RSS 值和 MAC 算术类型
DPF	dpf::Gen	生成两方 DPF key，调用方必须提供安全随机 seed
DPF	dpf::PreparedKey	将 correction words 上传并常驻 GPU
DPF	dpf::EvalBatchAsync	使用已上传 key 批量执行异步 GPU Eval
DPF	dpf::EvalAllAsync	逐层展开 FSS 树并复用公共前缀，评估完整输入域
DCF	dcf::Gen、dcf::PreparedKey、dcf::EvalBatchAsync	DCF 对应的 key 生成和 GPU Eval
DCF	dcf::EvalAllAsync	DCF 完整输入域的 GPU EvalAll
序列化	SerializeDpfKeyShare、DeserializeDpfKeyShare	将 DPF key 转换为可发送的 byte / char* 数据
序列化	SerializeDcfKeyShare、DeserializeDcfKeyShare	将 DCF key 转换为可发送的 byte / char* 数据
RSS	AddLocalBatchAsync、SubLocalBatchAsync	异步提交本地 RSS share pair 加减，不主动等待 GPU 完成
RSS	PrepareReshareOutgoingBatchAsync	异步计算两份输入之和并加入零分享，生成待发送 buffer
RSS	PrepareMulOutgoingBatchAsync	异步计算 RSS 乘法中的本地乘积和，生成待发送 buffer
RSS	AssembleSharePairBatchAsync	异步将收到的 share 与本地 share 组成新的 RSS share pair
RSS	不带 Async 后缀的同名接口	保留原有同步行为，兼容已有调用
谓词计算	predicate::EvalPredBatch	将 FSS 输出与 domain-major one-hot 表 share 逐项乘加，输出每条记录的 predicate share
范围谓词	predicate::EvalRangeFromLtBatch	将两个小于型 DCF Eval share 相减，得到 [lower, upper) 的范围 share
范围谓词	predicate::EvalRangeFromLtAllAsync	使用两个 DCF EvalAll 结果生成完整输入域的范围 share
聚合	aggregation::ReduceSum	对 32/64/128 位 GPU buffer 做求和聚合

连续处理输入深度相同的查询时，可以在首次查询前创建 PreparedKey 和 EvalAllWorkspace。后续查询通过 PreparedKey::Reset 上传新 key，并继续把同一个 EvalAllWorkspace 传给 EvalAll；底层 GPU 内存保持不变，不会为每个谓词重复执行 cudaMalloc 和 cudaFree。这里复用的只是内存空间，不会沿用上一次查询的 key、输入或输出。输入深度变化时，需要创建对应尺寸的对象并重新分配缓冲区。

DPF、DCF 的输出环通过模板参数选择：

using Ring = fss_sse_gpu::runtime::Ring32;

auto keys = fss_sse_gpu::primitives::dpf::Gen<16, Ring>(
    alpha, fss_sse_gpu::runtime::Share32{1}, secure_seed_pair);

fss_sse_gpu::primitives::dpf::PreparedKey<16, Ring> prepared(keys.party0);
fss_sse_gpu::primitives::dpf::EvalBatchAsync<16, Ring>(
    prepared, d_inputs, d_outputs, count);

secure_seed_pair 应由接入方案使用 CSPRNG 生成。DeterministicTestSeedPair 仅用于可重复测试，不能用于实际查询密钥。

范围谓词使用两个 beta=1 的小于型 DCF key：

Range_[L,U)(x) = LT_U(x) - LT_L(x)

闭区间 [L,R] 可在 R 不是输入域最大值时转换为 [L,R+1)。当右端点等于输入域最大值时，需要由方案层提供正确共享的常数 1，再计算 1-LT_L(x)。

参数边界

半诚实值可以在 Z_(2^32) 中计算。需要认证的值和 MAC 可以使用原生 Z_(2^128) 本地算术，但协议层仍需负责 MAC key、随机性、通信和最终校验。

• Ring32 用于当前半诚实 DPF/DCF、RSS 和三方单谓词测试。
• AuthenticatedShare 使用无符号 128 位回绕，实现恶意安全路径所需的本地 RSS、EvalPred 和聚合算术。
• FSS 的输入深度由特征域大小决定；PRG seed 和 key 参数应由接入协议的安全目标确定。

当前内置 FSS 后端使用 128 位 seed block，但已验证的 FSS 输出环仅为 32 位和 64 位。AuthenticatedShare 提供 128 位本地算术，不代表 FSS Gen/Eval 已支持 128 位 payload；这类 key 必须单独扩展并验证。

协议边界

半诚实单谓词可以只做服务器本地计算：

GPU FSS Eval
  -> GPU EvalPred(FSS share, one-hot table share)
  -> GPU ReduceSum（count）
  -> 客户端重构三组结果

这一流程不需要服务器间 RSSReshare。多个谓词需要把中间 filter 转为 RSS share 并进行乘法组合时，框架再把重分享和乘法拆成“GPU 本地计算”和“外部网络通信”两部分：

GPU: 计算 outgoing share buffer
  -> 方案网络层: 发送给相邻服务器并接收 share
  -> GPU: 组装新的 RSS share pair

框架不会在单个进程中收集全部服务器 share，也不提供通过明文重构来模拟 RSS 乘法的接口。中间结果可以继续保存在 device buffer 中，只有网络层需要的数据和最终结果才由接入方案决定是否传出 GPU。

RSS 的 *Async 接口只向 BatchConfig::stream 提交 kernel，调用方负责在读取通信 buffer 或结果之前等待对应 stream。多个本地步骤存在数据依赖时，可以提交到同一个 stream，并在最后统一等待：

rss::PrepareReshareOutgoingBatchAsync(
    d_first, d_second, d_zero, d_outgoing, count, config);
rss::AssembleSharePairBatchAsync(
    d_received, d_outgoing, d_pair, count, config);
rss::AddLocalBatchAsync(d_pair, d_y, d_result, count, config);
cudaStreamSynchronize(config.stream);

测试范围

正确性测试覆盖：

• in_bits=8/16/20/32 的 DPF、DCF；
• 32 位和 64 位 FSS 输出环；
• FSS key 序列化、反序列化前后的 Eval 一致性；
• correction words 常驻 GPU 后的异步 Eval；
• DPF/DCF GPU EvalAll 与逐点 EvalBatch 的 share 完全一致；
• 32/64/128 位 EvalPredBatch 乘表及 count 聚合；
• 64 位 RSS 与 128 位认证 RSS 的加减、乘法本地步骤；
• 32/64/128 位 RSS 异步接口、非默认 stream、零长度输入和连续调用；
• 三组两方 DPF key、三份 RSS one-hot 表 share 和六次本地 EvalPred；
• Ring32 DCF 范围组合的普通、单点、空区间和边界情况；
• 三方单谓词 filter 与 count 的客户端重构；
• CPU 与 GPU 对相同 key、相同输入得到完全一致的 share。

三方测试覆盖半诚实单谓词的本地计算和客户端重构语义。实际网络接口、范围查询一致性检查和恶意安全校验仍应在接入协议的集成测试中验证。

eval_pred_benchmark 固定使用 2^8 个特征值，对 2^10, 2^12, ..., 2^20 条记录测试 Ring32 半诚实路径。输入 FSS share 和 one-hot 表 share 在计时前准备完成，GPU 计时不包含 H2D/D2H、FSS Eval、key 生成和客户端重构。

eval_all_benchmark 对比单线程逐点 Eval、单线程 EvalAll、GPU 逐点批量 Eval 和 GPU EvalAll。EvalAll 通过复用树上的公共前缀减少 PRG 调用，与 CPU 多线程属于两项独立优化。

算子层 CPU/GPU 性能

测试环境：AMD Ryzen 7 5800H、NVIDIA GeForce RTX 3070 Laptop GPU 8 GiB、CUDA 13.3、GCC 13.3、Release、sm_86。CPU reference 为单线程；GPU 的 prepared/EvalAll/EvalPred 时间假设输入、key、表和输出 buffer 已经在显存中。除特别标出的 key upload 路径外，不包含 H2D/D2H、网络或方案层处理。

Ring32 DPF 双 key Eval

每个配置模拟一台服务器持有两份 DPF key。GPU prepared 复用已上传的 correction words。

输入数	CPU 双 key (ms)	GPU 每轮上传 key (ms)	GPU prepared (ms)	CPU / GPU prepared
1,024	3.796	0.445	0.077	49.1x
4,096	14.150	0.477	0.079	178.5x
65,536	225.691	1.081	0.470	480.3x
1,048,576	3596.754	34.118	9.431	381.4x

Ring32 DCF 完整域 EvalAll

GPU batch 让每个输入独立遍历 FSS 树；GPU EvalAll 逐层展开并复用公共前缀。小输入域中逐层 kernel 启动占主导，达到 2^18 后 EvalAll 开始优于独立 Eval。

输入域	CPU 逐点 (ms)	CPU EvalAll (ms)	GPU batch (ms)	GPU EvalAll (ms)	GPU EvalAll 收益
2^10	0.984	0.116	0.023	0.121	0.19x
2^14	22.792	1.811	0.068	0.156	0.43x
2^16	109.228	7.968	0.216	0.231	0.94x
2^18	508.453	32.981	1.063	0.385	2.76x
2^20	2262.394	127.790	4.108	0.556	7.39x

Ring32 EvalPred

固定 2^8 个特征值，对 domain-major one-hot 表执行乘加。GPU 时间只包含 EvalPredBatch kernel。

记录数	表大小 (MiB)	CPU (ms)	GPU (ms)	加速比
2^10	1	0.051	0.017	2.9x
2^12	4	0.214	0.027	7.8x
2^14	16	1.419	0.052	27.4x
2^16	64	4.227	0.187	22.6x
2^18	256	17.872	0.774	23.1x
2^20	1,024	73.323	2.659	27.6x

原始输出和完整环境记录见 results/。这些结果用于说明算子级并行能力，不代表完整查询流程的端到端加速比。

Ring32 RSS 连续调用

rss_async_benchmark 比较三个本地步骤逐步等待 GPU 完成，以及三个步骤提交到同一个非默认 stream 后统一等待。输入和输出已在 GPU，计时包含 kernel launch 和同步，不包含显存分配、H2D/D2H 或网络通信。

数据量	每步等待（ms）	最后统一等待（ms）	前者 / 后者
1,000	0.140	0.055	2.54x
10,000	0.128	0.050	2.57x
100,000	0.136	0.067	2.02x
1,000,000	0.297	0.224	1.32x

三个步骤依次为准备重分享 outgoing buffer、组装 share pair 和本地加法。测试预先准备收到的 share，只测量通信两侧的 GPU 本地接口能否连续提交及其同步成本，不能解释为完整 RSS 重分享的加速比。原始数据见 results/2026-09-04_rtx3070_rss_async.csv。