# GPU によるゼロ知識証明生成の高速化

> ゼロ知識証明は、プライベートなトランザクションを明かすことなく検証可能な計算を実現し、プライバシーを保護するトランザクションやスケーラブルなオンチェーン検証のための強力なビルディングブロックとなります。

著者 Shodipo Ayomide、2025年9月7日 公開。Web 版：https://blog.shodipoayomide.com/ja/gpu-accelerated-zero-knowledge-proving

![GPU によるゼロ知識証明生成の高速化](https://blog.shodipoayomide.com/banners/GPU-Accelerated-Zero-Knowledge-Proving.png)

ゼロ知識証明は、プライベートなトランザクションを明かすことなく検証可能な計算を実現し、プライバシーを保護するトランザクションやスケーラブルなオンチェーン検証のための強力なビルディングブロックとなります。しかし、プルーフ生成の計算負荷の高さは、長らく大きな課題となってきました。ZK システムでよく知られたトレードオフが「リソースを食うプルーバー」です。小さくて素早く検証できるプルーフを実現するには膨大な計算が必要で、しかもかなり遅くなりがちなため、リアルタイムでの利用を妨げてしまいます。

現代の GPU は、この方程式を劇的に変えます。GPU の並列設計を活用することで、根本的な課題だったものを、解決可能なエンジニアリングの問題へと変えてくれるのです。今日の ZK システムは主に `zk-SNARKs` と `zk-STARKs` という 2 つの構成を使っていて、どちらも GPU による高速化にぴったり合うコア演算に依存しています。

## 実際に何にそんなに時間がかかっているのか？

プルーフ生成時間のほぼすべては、2 つの演算に費やされています。マルチスカラー乗算（MSM）と数論変換（NTT）です。

### マルチスカラー乗算とは？

MSM は、スカラー倍された楕円曲線上の点の総和 $S = \sum_{i=1}^{n} a_i \cdot P_i$ を計算します。この 1 つの演算だけで、`SNARKs` のプルーフ生成時間のなんと `60-80%` を占めています。楕円曲線上の点を取得するために大量のランダムなメモリアクセスが必要になるため、ここではメモリ帯域幅が主なパフォーマンス上の問題になります。Pippenger アルゴリズムは、スカラーをバケットに振り分けることでこのタスクを並列化します。

```rust
fn pippenger_msm<G: CurveGroup>(
    bases: &[G::Affine],
    scalars: &[G::ScalarField],
    window_size: usize,
) -> G {
    let num_buckets = 1 << window_size;
    let mut buckets = vec![G::zero(); num_buckets];
    
    for (base, scalar) in bases.iter().zip(scalars.iter()) {
        let bucket_index = extract_window(scalar, window_size);
        buckets[bucket_index] += base;
    }
    
    // bucket aggregation
    buckets.into_iter()
           .enumerate()
           .map(|(i, bucket)| bucket * G::ScalarField::from(i))
           .fold(G::zero(), |acc, term| acc + term)
}
```

では GPU 実装で何が起きているかというと、この処理を数千のスレッドに分割することで、`RTX 4090` 上で毎秒 `700 million` 回を超える点演算を実現しているのです。

### 数論変換とは？

一方、NTT は多項式の乗算を高速に行います。MSM のランダムなメモリジャンプとは違い、NTT は予測可能なパターンに従うので、まさに GPU が得意とするところです。

```c
__global__ void ntt_radix2_kernel(
    uint64_t* coeffs,
    uint64_t* twiddles,
    uint32_t stage,
    uint32_t n
) {
    uint32_t tid = threadIdx.x + blockIdx.x * blockDim.x;
    uint32_t stride = 1 << stage;
    uint32_t m = n >> (stage + 1);
    
    if (tid >= m) return;
    
    uint32_t k = tid % stride;
    uint32_t j = tid / stride;
    uint32_t idx1 = j * (stride << 1) + k;
    uint32_t idx2 = idx1 + stride;
    
    uint64_t u = coeffs[idx1];
    uint64_t v = mod_mul(coeffs[idx2], twiddles[k], MODULUS);
    
    coeffs[idx1] = mod_add(u, v, MODULUS);
    coeffs[idx2] = mod_sub(u, v, MODULUS);
}
```

この逐次的なパターンのおかげで、高性能な NTT 実装はメモリ帯域幅の `85-95%` を引き出すことができます。

## ケーススタディ：CPU vs GPU のプルーフ生成性能

Alice と Bob を考えてみましょう。2 人はそれぞれ自分の `Layer-2` チェーンで `zkEVM Sequencers` を運用しています。処理するトランザクションバッチはまったく同じですが、使っているプルーフ生成のインフラが異なります。

### シナリオ：zkEVM のブロックのプルーフ生成

どちらのシーケンサーも、次の操作を含む `500` 件のトランザクションからなるブロックを証明しなければなりません：

- **1,200** 件の ERC-20 送金
- **300** 件の Uniswap V3 スワップ
- **200** 件のストレージ更新
- **100** 件のコントラクトデプロイ

これにより、約 $2^{26}$ 個の制約を持つ回路が生成されます。中規模の `zkEVM` ブロックとしては典型的な規模です。

### Alice の CPU アプローチ

Alice は、ハイエンドな `CPU` - `AMD EPYC 7763, 64 cores` でプルーバーを動かしています：

```rust
impl CpuZkEvmProver {
    fn prove_block(&self, block: &Block) -> ProofResult {
        let start = Instant::now();
        
        // circuit compilation
        let circuit = self.compile_block_circuit(block); // 45s
        
        // msm operations
        let witness_commitment = self.cpu_msm(&circuit.witnesses); // 180s
        
        // ntt operations
        let constraint_proof = self.cpu_ntt(&circuit.constraints); // 90s
        
        // final proof
        let proof = self.assemble_proof(witness_commitment, constraint_proof);
        
        ProofResult {
            proof,
            proving_time: start.elapsed(), // ~315s total
            memory_usage: 45_000_000_000, // 45GB peak
        }
    }
    
    fn cpu_msm(&self, witnesses: &[Fr]) -> G1Projective {
        witnesses.iter()
                .zip(self.setup.iter())
                .map(|(scalar, base)| base.mul(scalar))
                .fold(G1Projective::identity(), |acc, point| acc + point)
    }
}
```

Alice の結果：

- **プルーフ生成時間：** 1 ブロックあたり 5.25 分
- **メモリ使用量：** ピーク時 45GB
- **スループット：** 1 分あたり約 95 トランザクションを証明
- **ハードウェアコスト：** $8,000 のサーバー

### Bob の GPU 高速化アプローチ

Bob は、`ICICLE` による高速化を施した `RTX 4090` を使っています：

```rust
use icicle_bls12_381::{curve::G1Projective, msm};
use icicle_core::msm::MSMConfig;

impl GpuZkEvmProver {
    fn prove_block(&self, block: &Block) -> ProofResult {
        let start = Instant::now();
        
        // circuit compilation
        let circuit = self.compile_block_circuit_gpu(block); // 12s
        
        // gpu msm - 16,384 cores parallel
        let witness_commitment = self.gpu_msm(&circuit.witnesses); // 8s
        
        // gpu ntt
        let constraint_proof = self.gpu_ntt(&circuit.constraints); // 3 seconds
        
        // final proof
        let proof = self.assemble_proof(witness_commitment, constraint_proof);
        
        ProofResult {
            proof,
            proving_time: start.elapsed(), // ~23s total
            memory_usage: 18_000_000_000, // 18GB peak
        }
    }
    
    fn gpu_msm(&self, witnesses: &[Fr]) -> Result<G1Projective> {
        let cfg = MSMConfig {
            precompute_factor: 8,
            c: 16,
            ..Default::default()
        };
        
        msm::msm(witnesses, &self.gpu_setup, &cfg)
    }
}
```

Bob の結果：

- **プルーフ生成時間：** 1 ブロックあたり 23 秒
- **メモリ使用量：** ピーク時 18GB
- **スループット：** 1 分あたり約 1,300 トランザクションを証明
- **ハードウェアコスト：** $1,800（RTX 4090 + 基本的なサーバー）

### 結果を比較する

| 処理 | Alice（CPU） | Bob（GPU） | 高速化率 |
|-----------|-------------|-----------|---------|
| 回路のコンパイル | 45s | 12s | 3.8× |
| MSM 演算 | 180s | 8s | **22.5×** |
| NTT 演算 | 90s | 3s | **30×** |
| **ブロックのプルーフ全体** | **315s** | **23s** | **13.7×** |

### なぜこれがすべてを変えるのか

ネットワークのファイナリティ：

- Alice のネットワーク：L1 での承認まで 5.25 分のファイナリティ
- Bob のネットワーク：L1 での承認まで 23 秒のファイナリティ

#### トランザクションあたりのコスト

```rust
// one every 12s
const BLOCKS_PER_DAY: u32 = 7200;

impl ProvingEconomics {
    fn daily_capacity(&self, proving_time_seconds: u32) -> u32 {
        let daily_seconds = 24 * 60 * 60;
        daily_seconds / proving_time_seconds
    }
}
```

1 日あたりのプルーフ生成キャパシティ（Alice vs Bob）：

- Alice が証明できるのは：1 日あたり 274 ブロック、つまり 137,000 トランザクション
- Bob が証明できるのは：1 日あたり 3,756 ブロック、つまり 1,878,000 トランザクション

#### Alice が対抗するために必要なもの

Bob のたった 1 台の GPU 構成に追いつくには、Alice は 14 台の CPU が必要になり、その結果：

- 14 倍高いハードウェアコスト - \$112,000 vs \$8,000
- 14 倍もの消費電力 - 14kW vs 1kW
- 複雑な分散協調のオーバーヘッド

### メモリの読み取り方

パフォーマンスの差は、根本的には両者のサーバーがどのようにメモリにアクセスするかに由来します：

```rust
// alice's cpu msm - sequential, cache-bound
for (scalar, base) in witnesses.iter().zip(setup.iter()) {
    result += base * scalar; // cache miss every iteration
}

// bob's gpu msm - parallel, lightweight
__global__ void msm_kernel(scalar_t* scalars, point_t* bases, point_t* buckets) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int bucket_idx = extract_window(scalars[tid], WINDOW_SIZE);
    atomicAdd(&buckets[bucket_idx], bases[tid]); // coalesced memory access
}
```

たとえばこのケーススタディは、本番環境の `zkEVM` デプロイで GPU による高速化が不可欠になった理由をシンプルに示しています。Bob のインフラは、より速いファイナリティと低い運用コストを維持しながら、`13.7×` 多くのユーザーにサービスを提供できます。

## SNARK と STARK の高速化パターン

### SNARK システム

`Groth16` と `PLONK` は MSM が中心です。たとえば `Groth16` は、プルーフ生成時間の約 `80%` を楕円曲線系の演算に費やしています。一方 `PLONK` は、MSM `~60%` と多項式演算 `~30%` の間で負荷をより均等に分散させているため、両方をバランスよく最適化する必要があります。

```rust
impl<E: Pairing> Groth16Prover<E> {
    fn create_proof(&self, circuit: &Circuit<E::ScalarField>) -> Result<Proof<E>> {
        let (a, b, c) = circuit.to_matrices();
        
        let g_a = self.params.alpha_g1.msm(&a);
        let g_b = self.params.beta_g2.msm(&b);  
        let g_c = self.params.gamma_g1.msm(&c);
        
        Ok(Proof { a: g_a, b: g_b, c: g_c })
    }
}
```

### STARK システム

`STARKs` の処理の大半は、多項式演算の `NTTs` と、`Poseidon` や `Blake3` などのハッシュ計算が占めています。`STARKs` のメモリアクセスパターンは規則的なので、`SNARK` の MSM 演算に見られるランダムアクセス寄りのパターンと比べて、GPU のメモリ帯域幅をより有効に活用できることが多いです。

```rust
impl<F: PrimeField> StarkProver<F> {
    fn prove(&self, trace: &ExecutionTrace<F>) -> StarkProof<F> {
        // polynomial interpolation via ntt
        let trace_poly = self.interpolate(trace);
        
        let composition_poly = self.evaluate_constraints(&trace_poly);
        
        let fri_proof = self.fri_commit(&composition_poly);
        
        StarkProof {
            trace_commitment: self.commit(&trace_poly),
            composition_commitment: self.commit(&composition_poly),
            fri_proof,
        }
    }
}
```

## zkEVM の高速化

zkEVM には、`EVM` 命令の複雑さと大きなステート表現に起因する特有の課題があります。プルーフ生成の過程では、`EVM` の各オペコードとステート遷移ごとに制約を生成する必要があります。

### EVM を証明する

```rust
impl ZkEvmProver {
    fn prove_block(&self, block: &Block) -> ZkEvmProof {
        let mut circuit = EvmCircuit::new();
        
        for tx in &block.transactions {
            for (pc, opcode) in tx.trace.iter().enumerate() {
                match opcode {
                    OpCode::ADD => {
                        let (a, b) = self.stack.pop2();
                        let result = a + b;
                        circuit.add_constraint(a + b - result);
                        self.stack.push(result);
                    },
                    OpCode::SSTORE => {
                        let (key, value) = self.stack.pop2();
                        circuit.add_storage_constraint(key, value);
                        self.update_state_root(key, value);
                    },
                    // ...
                }
            }
        }
        
        self.prove_circuit(circuit)
    }
}
```

**ADD** から **SSTORE** まで、各オペコードは多項式制約のセットに変換されなければなりません。

### ステートツリーを証明する

zkEVM は基本的に、メモリアクセスパターンとストレージ更新が正しいことを常に証明しなければなりません。これには通常、ステートルートを検証するための `Merkle tree` の操作が伴います

```rust
fn constrain_merkle_update(
    circuit: &mut Circuit,
    old_root: Variable,
    new_root: Variable,
    key: Variable,
    old_value: Variable,
    new_value: Variable,
    path: &[Variable],
) {
    let mut current = old_value;
    let mut updated = new_value;
    
    for (i, &sibling) in path.iter().enumerate() {
        let bit = circuit.extract_bit(key, i);
        
        let hash_input = circuit.select(bit, (sibling, current), (current, sibling));
        current = circuit.poseidon_hash(hash_input.0, hash_input.1);
        
        let updated_input = circuit.select(bit, (sibling, updated), (updated, sibling));
        updated = circuit.poseidon_hash(updated_input.0, updated_input.1);
    }
    
    circuit.assert_equal(current, old_root);
    circuit.assert_equal(updated, new_root);
}
```

GPU による高速化は、ステート遷移の検証に必要な大量のハッシュ演算で特に重要になります

## まとめ

GPU は単に ZK のプルーフ生成を速くしただけではありません。「リソースを食うプルーバー」問題を解決したことで、ZK のプルーフ生成のあり方そのものを根本から変えたのです。今では数秒でファイナライズする Layer 2 や、本番環境で動くプライベートな DeFi が実現していますが、これらは以前は到底不可能だったものです。`zkEVMs` のようなアプリケーションによって回路の複雑さが増すにつれ、GPU による高速化はもはや単なる最適化ではなく、本番環境へのデプロイに欠かせない基本要件になっています。

### 必読リンク

  * **ICICLE：** Ingonyama による、ZK プリミティブを高速化するためのオープンソース GPU ライブラリ。[GitHub](https://github.com/ingonyama-zk/icicle)
  * **Arkworks：** ZK 証明システムを設計・実装するための Rust ベースのエコシステム。[ウェブサイト](https://arkworks.rs/)
  * **Halo2：** Zcash Foundation が開発した PLONK ベースの証明システムで、PSE zkEVM などのプロジェクトで広く使われています。[GitHub](https://github.com/zcash/halo2)
  * **ZPrize：** ソフトウェアとハードウェアを通じて ZK のパフォーマンス向上を目指す業界コンペティション。[ウェブサイト](https://www.zprize.io/)
  * Vitalik Buterin による「zk-SNARKs: A Primer」：zk-SNARK がどのように構築されるのかを高いレベルで解説した優れた記事。[記事](https://vitalik.eth.limo/general/2021/01/26/snarks.html)
  * StarkWare による「Anatomy of a STARK」：zk-STARK の数学的な構成要素を深く掘り下げた解説。[記事](https://medium.com/starkware/stark-math-the-journey-begins-51bd2b063c71)
  * Ingonyama による「Introducing ICICLE: An Open-Source GPU Library for Zero Knowledge Acceleration」[記事](https://www.ingonyama.com/post/introducing-icicle-an-open-source-gpu-library-for-zero-knowledge-acceleration)
  * Space and Time による「Blitzar, A novel opensource zero-knowledge proof for SQL operations」。Blitzar は、GPU で高速化された ZK プルーフを構築するための、より高速で堅牢なフレームワークを提供します [Github](https://github.com/spaceandtimefdn/blitzar)
