# GPU 加速的零知识证明生成

> 零知识证明让计算可以被验证，同时又不会暴露私密交易，为隐私保护交易和可扩展的链上验证提供了强大的基础构件。

作者 Shodipo Ayomide，发布于 2025年9月7日。网页版本：https://blog.shodipoayomide.com/zh/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`，它们所依赖的核心运算都与 GPU 加速完美契合。

## 到底是什么这么耗时？

几乎所有的证明时间都花在了两种运算上：多标量乘法（Multi-Scalar Multiplication，MSM）和数论变换（Number Theoretic Transform，NTT）。

### 什么是多标量乘法？

MSM 计算的是一组经过标量缩放的椭圆曲线点之和：$S = \sum_{i=1}^{n} a_i \cdot P_i$。在 `SNARKs` 中，单单这一个运算就要吃掉高达 `60-80%` 的证明时间。它本身需要进行大量的随机内存访问来获取椭圆曲线点，这让内存带宽成为了这里的主要性能瓶颈。Pippenger 算法通过把标量分拣到不同的桶（bucket）中，来实现这项任务的并行化。

```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 与 GPU 的证明性能对比

假设 Alice 和 Bob 都在为各自的 `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 的结果：

- **证明时间：** 每个区块 5.25 分钟
- **内存占用：** 峰值 45GB
- **吞吐量：** 每分钟约证明 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 的结果：

- **证明时间：** 每个区块 23 秒
- **内存占用：** 峰值 18GB
- **吞吐量：** 每分钟约证明 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
    }
}
```

每日证明能力（Alice 对比 Bob）：

- Alice 每天可以证明 274 个区块，也就是 137,000 笔交易
- Bob 每天可以证明 3,756 个区块，也就是 1,878,000 笔交易

#### Alice 要怎样才能跟上

Alice 需要 14 台 CPU 服务器，才能赶上 Bob 单块 GPU 的配置，这会导致：

- 硬件成本高出 14 倍 - \$112,000 对比 \$8,000
- 功耗高出 14 倍 - 14kW 对比 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
}
```

这个案例研究简单地说明了，为什么 GPU 加速已经成为生产环境 `zkEVM` 部署的必需品。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 加速

由于 `EVM` 指令的复杂性和庞大的状态表示，zkEVM 面临着独特的挑战。证明过程需要为每一个 `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 开发的开源 GPU 库，用于加速 ZK 原语。[GitHub](https://github.com/ingonyama-zk/icicle)
  * **Arkworks：** 一个基于 Rust 的生态系统，用于设计和实现 ZK 证明系统。[官网](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)