1. GPUアクセラレーションにおけるプライベート情報検索の技術的背景
プライベート情報検索(PIR)技術は理論から実用段階へと進化しています。アップルのプライベートビジョンサーチシステムが示したように、従来のPIRスキームは商業価値を証明しました。現代の格子ベース準同型暗号(HE)はPIRに数学的基盤を提供しますが、計算複雑度がボトルネックに。GB級データベースでの単一クエリに数秒かかるのは、データベース全体に対する暗号演算が必要なためです。
GPUアクセラレーションはこの課題解決の新機軸を提供します。NVIDIA RTX 5090の例では、31.5 TOPSの整数演算スループットと1660GB/sのメモリ帯域幅が、PIRの中核演算である数論変換(NTT)と行列乗算(GEMM)に適しています。しかし実装には3つの課題があります：
-
計算特性の異質性 PIRプロトコルはExpandQuery、RowSel、ColTorの3段階からなり、それぞれ異なる計算パターンを示します。ExpandQueryとColTorはNTT主体、RowSelは並列GEMM演算です。
-
メモリ壁問題 クライアントのバッチ処理でスループットが向上しますが、RTX 5090の96MB L2キャッシュを超えるとDRAMアクセスが激増し、性能が急落します。
-
データ配置の競合 HE演算にはNTT最適化配置(p-major)が必要ですが、RowSelのGEMMにはtiling最適化配置(major/k-major)が必要で、性能損失が50%にもなります。
実装ヒント：SEALやHElibなどの既存HEライブラリを単純適用するのはNG。PIRonGPUなどのオープンソース実装はバッチ処理最適化が不十分で、スループットが理論値の1%未満になるケースがあります。
2. PIRプロトコルの三段階計算特性
2.1 クエリ拡張フェーズ(ExpandQuery)
ExpandQueryはクライアントの単一暗号文をN個に拡張します。各ノードで実行されるSubs演算の核心：
struct CipherText Subs(struct CipherText ct, struct EvalKey evk) {
// 1. 数論変換
float* poly_data = inverseNTT(ct.a);
// 2. 数字分解(Dcp)
float* decomposed = (float*)malloc(5 * sizeof(float));
for(int i=0; i<5; i++) {
decomposed[i] = (poly_data >> (i*22)) & 0x3FFFFF; // 基数z=2^22分解
}
// 3. evkとの累積乗算
float result = 0;
for(int i=0; i<5; i++) {
result += forwardNTT(decomposed[i]) * evk.coeff[i];
}
return createCT(result);
}
メモリアクセス特性：
- 段階的指数増加：k層目で2^k個の暗号文処理
- 一時ワークセット：Dcpで5倍(z=2^22, ℓ=5)のデータ膨張
- キー共有：同層ノードでeval_key(evk)を共有
2.2 行選択フェーズ(RowSel)
RowSelは暗号化行列乗算の実装：
void matrixMul(float* M_out, float* Min0, float* Min1, int p, int m, int k, int n) {
#pragma omp parallel for collapse(2)
for(int i=0; i<p; i++) {
for(int j=0; j<n; j++) {
float sum = 0;
for(int l=0; l<k; l++) {
sum += Min0[i*m*k + l] * Min1[l*n + j];
}
M_out[i*m*n + j] = sum;
}
}
}
バッチ処理32クエリ時の次元規模：
- p: 4N (多項式点数)
- m: 64 (2×32)
- k: D0 (データベース行数)
- n: D1 (データベース列数)
算術強度が0.9→13.8 Ops/Byteに向上し、メモリバンド幅制限から計算制限へ転換。
2.3 列トーナメントフェーズ(ColTor)
ColTorはトーナメント形式で目標列を抽出します。ExpandQueryと対称的ですが、2つの相違点：
- 双多項式処理：暗号文のa/b部分を同時に処理しワークセットが2倍
- RGSW暗号文使用：外部積(External Product)演算が必要
3. フェーズ認識型ハイブリッド実行モデル
3.1 キャッシュ容量壁現象
バッチ処理規模32時のワークセット増加：
| 木の深さ | 単一クエリワークセット | バッチ32ワークセット |
|---|---|---|
| 1 | 64KB | 2MB |
| 5 | 1MB | 32MB |
| 10 | 32MB | 1GB |
RTX 5090での測定ではL2キャッシュの60%(約57.6MB)超えでDRAMトラフィックが3倍以上に。
3.2 オペレーションレベルとフェーズレベルカーネル比較
オペレーションレベル：
- 個別演算(NTT/Dcp)を独立カーネル化
- 中間結果をDRAMに書き戻し
フェーズレベル：
- 全演算を単一カーネルに融合
- 中間データをレジスタ/SMEM保持
測定性能比較(バッチ32)：
| 拡張タイプ | カーネルタイプ | DRAMトラフィック | 実行時間 |
|---|---|---|---|
| シャローエクスパンション | オペレーションレベル | 12GB | 38ms |
| フェーズレベル | 15GB | 52ms | |
| ディープエクスパンション | オペレーションレベル | 142GB | 210ms |
| フェーズレベル | 71GB | 134ms |
3.3 動的切り替え戦略
void executeOptimalKernel(int stage_depth, int batch_size) {
size_t working_set = calculateMemoryUsage(stage_depth, batch_size);
if (working_set < L2_CACHE_SAFE_ZONE) { // 80MBの安全マージン
launchOperationLevelKernels();
} else {
launchStageLevelKernel();
}
}
4. RowSel向け配置最適化技術
4.1 データ配置競合分析
従来のHEライブラリのNTT最適化配置：
[p][m][k] : p-major (4N点連続)
効率的なGEMMには：
[m][p][k] : m-major (出力行列)
[k][p][m] : k-major (入力行列)
RTX 5090での測定結果：
| 配置タイプ | 計算効率 | L2キャッシュヒット率 |
|---|---|---|
| p-major | 41% | 68% |
| m-major | 83% | 92% |
| k-major | 79% | 89% |
4.2 転置配置GEMM設計
- オンライン転置パイプライン：
__global__ void optimizedGEMM(float* A, float* B, float* C) {
__shared__ float tileA[TILE_DIM][TILE_DIM+1]; // バンク競合回避
__shared__ float tileB[TILE_DIM][TILE_DIM+1];
for (int block=0; block<K; block += TILE_DIM) {
// 協調的データ転置読み込み
cooperativeLoadTranspose(A, tileA);
cooperativeLoadTranspose(B, tileB);
__syncthreads();
// 転置行列演算
computeTile(tileA, tileB, C);
}
}
- ハイブリッド精度演算：
- 入力：FP32 (HE精度要件)
- 中間：FP64 (整数溢れ防止)
- 出力：FP32
RTX 5090で1.7TFLOPS達成(理論ピークの85%)
5. マルチGPU拡張スキーム
5.1 データベースシャーディング
2次元シャーディングによる行/列スケーリング：
GPU0: DB[0:D0/2][0:D1/2] GPU1: DB[0:D0/2][D1/2:D1]
GPU2: DB[D0/2:D0][0:D1/2] GPU3: DB[D0/2:D0][D1/2:D1]
NVLinkによる高速データ交換で1/4データ保持。
5.2 GPU間同期最適化
CUDAグラフによるトーナメント構造最適化：
cudaGraph_t buildReductionGraph() {
cudaGraph_t graph;
cudaGraphCreate(&graph, 0);
addLocalTournament(graph, gpu0);
addMergeNode(graph, gpu0, gpu1);
addFinalComputation(graph, gpu0);
return graph;
}
6. 実測性能と最適化推奨
2GBデータベースでの結果：
| 最適化項目 | スループット改善 | 遅延削減 |
|---|---|---|
| ハイブリッド実行モデル | 4.2× | 61% |
| 転置GEMM | 3.1× | 68% |
| マルチGPU拡張(4×) | 3.8× | 73% |
| 総合最適化 | 305.7× | 96% |
導入推奨事項：
- 最適バッチ処理量：
int optimalBatch = (0.6 * L2_CACHE_SIZE) / (DECOMPOSE_FACTOR * STAGE1_DATA_SIZE);
- カーネル選択ヒューリスティック：
#define STAGE_THRESHOLD 80MB
bool useStageLevel(int depth, int batch) {
return (64KB * (1 << depth) * batch) > STAGE_THRESHOLD;
}
- メモリ割当戦略：
- cudaMallocAsyncによるパイプラインバッファ確保
- cudaMemAdviseReadMostlyによる読み取り最適化
パフォーマンストラップ：
- 256バイトのメモリアラインメント
- デフォルトストリームでの同期回避
- 32スレッド超時のSMEMパディング