# Geração de provas zero-knowledge acelerada por GPU

> As provas zero-knowledge permitem computação verificável sem expor transações privadas, criando blocos de construção poderosos para transações que preservam a privacidade e para verificação onchain escalável.

Por Shodipo Ayomide, publicado em 7 de set. de 2025. Versão web: https://blog.shodipoayomide.com/pt/gpu-accelerated-zero-knowledge-proving

![Geração de provas zero-knowledge acelerada por GPU](https://blog.shodipoayomide.com/banners/GPU-Accelerated-Zero-Knowledge-Proving.png)

As provas zero-knowledge permitem computação verificável sem expor transações privadas, criando blocos de construção poderosos para transações que preservam a privacidade e para verificação onchain escalável. Só que a intensidade computacional da geração de provas sempre foi um problema grande. Um tradeoff bem conhecido nos sistemas ZK é o "prover faminto por recursos": conseguir uma prova pequena e rápida de verificar exige muita computação, mas muitas vezes é bem lento, o que impede o uso em tempo real.

As GPUs modernas mudam essa equação de forma drástica. Aproveitando o seu design paralelo, elas transformam o que era um problema fundamental em um problema de engenharia que dá para resolver. Os sistemas ZK de hoje usam principalmente duas construções: `zk-SNARKs` e `zk-STARKs`, e ambas dependem de operações centrais que se encaixam perfeitamente na aceleração por GPU.

## O que realmente demora tanto?

Quase todo o tempo de geração de provas vai para duas operações: Multi-Scalar Multiplication (MSM) e Number Theoretic Transform (NTT).

### O que é Multi-Scalar Multiplication?

A MSM calcula a soma de pontos de curva elíptica escalados: $S = \sum_{i=1}^{n} a_i \cdot P_i$. Essa única operação consome nada menos que `60-80%` do tempo de geração de provas em `SNARKs`. Ela exige por si só um número alto de acessos aleatórios à memória para buscar os pontos da curva elíptica, o que faz da largura de banda de memória o principal problema de performance aqui. O algoritmo de Pippenger paraleliza essa tarefa organizando os escalares em buckets.

```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)
}
```

Então, o que acontece na implementação em GPU é que ela divide isso entre milhares de threads, chegando a mais de `700 million` operações de ponto por segundo em uma `RTX 4090`.

### O que é Number Theoretic Transform?

A NTT, por outro lado, multiplica polinômios rápido e, ao contrário dos saltos aleatórios na memória da MSM, a NTT segue um padrão previsível, exatamente aquilo para o que as GPUs foram feitas.

```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);
}
```

Esse padrão sequencial permite que implementações de NTT com boa performance alcancem `85-95%` da largura de banda de memória.

## Estudo de caso: performance de geração de provas em CPU vs GPU

Pense na Alice e no Bob, os dois operando `zkEVM Sequencers` para as suas respectivas chains `Layer-2`. Eles processam lotes de transações idênticos, mas usam infraestruturas de geração de provas diferentes.

### Cenário: geração de prova de um bloco da zkEVM

Os dois sequencers precisam provar um bloco que contém `500` transações com as seguintes operações:

- **1.200** transferências ERC-20
- **300** swaps no Uniswap V3
- **200** atualizações de storage
- **100** deploys de contratos

Isso gera um circuito com aproximadamente $2^{26}$ restrições, o que é típico para um bloco `zkEVM` de tamanho médio.

### A abordagem da Alice com CPU

A Alice roda o prover dela em uma `CPU` de ponta - `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)
    }
}
```

Resultados da Alice:

- **Tempo de geração da prova:** 5,25 minutos por bloco
- **Uso de memória:** pico de 45GB
- **Throughput:** ~95 transações/minuto provadas
- **Custo de hardware:** servidor de US$ 8.000

### A abordagem do Bob acelerada por GPU

O Bob usa uma `RTX 4090` com aceleração `ICICLE`:

```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)
    }
}
```

Resultados do Bob:

- **Tempo de geração da prova:** 23 segundos por bloco
- **Uso de memória:** pico de 18GB
- **Throughput:** ~1.300 transações/minuto provadas
- **Custo de hardware:** US$ 1.800 (RTX 4090 + servidor básico)

### Comparando os resultados

| Operação | Alice (CPU) | Bob (GPU) | Aceleração |
|-----------|-------------|-----------|---------|
| Compilação do circuito | 45s | 12s | 3,8× |
| Operações MSM | 180s | 8s | **22,5×** |
| Operações NTT | 90s | 3s | **30×** |
| **Prova total do bloco** | **315s** | **23s** | **13,7×** |

### Por que isso muda tudo

Finalidade da rede:

- Rede da Alice: finalidade de 5,25 minutos para confirmações na L1
- Rede do Bob: finalidade de 23 segundos para confirmações na L1

#### Custo por transação

```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
    }
}
```

Capacidade diária de geração de provas (Alice vs Bob):

- A Alice consegue provar: 274 blocos/dia, ou seja, 137.000 transações
- O Bob consegue provar: 3.756 blocos/dia, ou seja, 1.878.000 transações

#### O que a Alice precisa para competir

A Alice precisaria de 14 CPUs para igualar o setup de uma única GPU do Bob, o que leva a:

- Custo de hardware 14× mais caro - US\$ 112.000 vs US\$ 8.000
- Consumo de energia 14× maior - 14kW vs 1kW
- Overhead de coordenação distribuída complexa

### Como a memória é lida

A diferença de performance começa, no fundo, pela forma como os dois servidores acessam a memória:

```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
}
```

Este estudo de caso, por exemplo, simplesmente mostra por que a aceleração por GPU se tornou essencial para deploys de `zkEVM` em produção. A infra do Bob consegue atender `13.7×` mais usuários enquanto mantém uma finalidade mais rápida e custos operacionais menores.

## Padrões de aceleração em SNARK vs. STARK

### Sistemas SNARK

`Groth16` e `PLONK` são pesados em MSM. O `Groth16`, por exemplo, gasta cerca de `80%` do seu tempo de geração de provas em operações do tipo curva elíptica. Já o `PLONK` distribui a carga de forma mais equilibrada entre MSM `~60%` e operações polinomiais `~30%`, o que exige uma otimização equilibrada para as duas.

```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 })
    }
}
```

### Sistemas STARK

Os `STARKs` são dominados por aritmética polinomial `NTTs` e hashing, por exemplo `Poseidon` ou `Blake3`. Como os padrões de acesso à memória nos `STARKs` são regulares, isso muitas vezes leva a um aproveitamento melhor da largura de banda de memória da GPU, em comparação com os padrões de acesso mais aleatórios das operações MSM em `SNARK`.

```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,
        }
    }
}
```

## Aceleração da zkEVM

As zkEVMs trazem desafios únicos por causa da complexidade das instruções da `EVM` e das grandes representações de estado. O processo de geração de provas envolve gerar restrições para cada opcode da `EVM` e para cada transição de estado.

### Provando a 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)
    }
}
```

Cada opcode, de **ADD** a **SSTORE**, precisa ser traduzido em um conjunto de restrições polinomiais.

### Provando a árvore de estado

As zkEVMs em geral sempre precisam provar padrões corretos de acesso à memória e atualizações de storage. Isso normalmente envolve operações de `Merkle tree` para a verificação do state root

```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);
}
```

A aceleração por GPU geralmente se torna importante para as extensas operações de hashing exigidas nas verificações de transição de estado

## Conclusão

As GPUs não só aceleraram a geração de provas ZK, elas basicamente mudaram a forma como a geração de provas ZK é feita, resolvendo o problema do "prover faminto por recursos", Agora temos Layer 2s finalizando em segundos e DeFi privado rodando em produção, coisas que antes não eram totalmente possíveis. À medida que a complexidade dos circuitos aumenta com aplicações como as `zkEVMs`, a aceleração por GPU deixa de ser só uma questão de otimização, ela passa a ser um requisito fundamental para deploys em produção.

### Links essenciais

  * **ICICLE:** Uma biblioteca open-source para GPU da Ingonyama para acelerar primitivas ZK. [GitHub](https://github.com/ingonyama-zk/icicle)
  * **Arkworks:** Um ecossistema baseado em Rust para projetar e implementar sistemas de prova ZK. [Site](https://arkworks.rs/)
  * **Halo2:** Um sistema de geração de provas baseado em PLONK, desenvolvido pela Zcash Foundation e muito usado em projetos como a PSE zkEVM. [GitHub](https://github.com/zcash/halo2)
  * **ZPrize:** Uma competição do setor focada em acelerar a performance de ZK por meio de software e hardware. [Site](https://www.zprize.io/)
  * "zk-SNARKs: A Primer", de Vitalik Buterin: Uma excelente explicação de alto nível sobre como os zk-SNARKs são construídos. [Artigo](https://vitalik.eth.limo/general/2021/01/26/snarks.html)
  * "Anatomy of a STARK", da StarkWare: Um mergulho profundo nos componentes matemáticos dos zk-STARKs.[Artigo](https://medium.com/starkware/stark-math-the-journey-begins-51bd2b063c71)
  * "Introducing ICICLE: An Open-Source GPU Library for Zero Knowledge Acceleration", da Ingonyama [Artigo](https://www.ingonyama.com/post/introducing-icicle-an-open-source-gpu-library-for-zero-knowledge-acceleration)
  * "Blitzar, A novel opensource zero-knowledge proof for SQL operations", da Space and Time. O Blitzar oferece um framework mais rápido e mais robusto para construir provas ZK aceleradas por GPU [Github](https://github.com/spaceandtimefdn/blitzar)