Shodipo Ayomide の写真
すべての記事

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

·
2025年9月7日
·
5 分で読めます
·
••• 回閲覧
·
••• いいね

英語の原文から翻訳した記事です。 英語で読む

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

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

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

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

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

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

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

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 が得意とするところです。

__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 件のトランザクションからなるブロックを証明しなければなりません：

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

Alice の CPU アプローチ

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

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 の結果：

Bob の GPU 高速化アプローチ

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

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 の結果：

結果を比較する

処理Alice（CPU）Bob（GPU）高速化率
回路のコンパイル45s12s3.8×
MSM 演算180s8s22.5×
NTT 演算90s3s30×
ブロックのプルーフ全体315s23s13.7×

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

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

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

// 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 が対抗するために必要なもの

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

メモリの読み取り方

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

// 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% の間で負荷をより均等に分散させているため、両方をバランスよく最適化する必要があります。

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 のメモリ帯域幅をより有効に活用できることが多いです。

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 を証明する

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 の操作が伴います

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 による高速化はもはや単なる最適化ではなく、本番環境へのデプロイに欠かせない基本要件になっています。

必読リンク

•••
この記事をシェア：

最新情報を受け取る

最新の記事をメールでお届けします。