Traduzido do artigo original em inglês. Ler em inglês

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.
Quase todo o tempo de geração de provas vai para duas operações: Multi-Scalar Multiplication (MSM) e Number Theoretic Transform (NTT).
A MSM calcula a soma de pontos de curva elíptica escalados: . 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.
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.
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.
__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.
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.
Os dois sequencers precisam provar um bloco que contém 500 transações com as seguintes operações:
Isso gera um circuito com aproximadamente restrições, o que é típico para um bloco zkEVM de tamanho médio.
A Alice roda o prover dela em uma CPU de ponta - 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)
}
}Resultados da Alice:
O Bob usa uma RTX 4090 com aceleração ICICLE:
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:
| 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× |
Finalidade da rede:
// 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 precisaria de 14 CPUs para igualar o setup de uma única GPU do Bob, o que leva a:
A diferença de performance começa, no fundo, pela forma como os dois servidores acessam a memória:
// 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.
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.
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 })
}
}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.
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,
}
}
}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.
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.
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
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
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.
Receba os artigos mais recentes no seu e-mail.