⚡ cuda-labs

Aprendendo CUDA de verdade — de funções de hash a um pipeline secp256k1 completo que transforma chave privada em endereço Bitcoin a mais de 1 bilhão de chaves por segundo numa RTX 4060.

Começou como um punhado de exemplos de hashing na GPU e foi crescendo, um conceito por vez, até virar o núcleo de um buscador de chaves no estilo BitCrack / VanitySearch / Keyhunt — só que construído do zero, para aprender, com cada peça conferida contra uma referência independente. É um projeto pessoal e educativo: todos os testes buscam apenas chaves que já conhecemos (veja Escopo & ética).

---

✨ Destaques

• Hashing em lote na GPU — SHA-256, RIPEMD-160 e a cascata HASH160, cada um processando milhares de mensagens em paralelo.
• Pipeline secp256k1 completo — aritmética de corpo → curva elíptica → multiplicação escalar → inversão em lote → chave pública → endereço, tudo fundido num único kernel.
• Rápido de verdade — ~1.250 Mkeys/s sustentados (1,25 Gkey/s), no mesmo nível de referências consagradas do ramo.
• CLI Python bonita — barra de progresso ao vivo, tabela de puzzles, presets, retomada automática de onde parou, suíte de sanidade e um auto-tuner de GPU embutido.
• Correção levada a sério — cada camada é validada contra um oráculo independente em Python, então um bug sutil de sinal ou endianness não passa despercebido.

---

🧠 Como funciona (a ideia, sem matemática)

Descobrir o endereço a partir de uma chave privada é um caminho de mão única — fácil de ir, impossível de voltar. A busca simplesmente percorre um intervalo de chaves e refaz esse caminho para cada uma, comparando o resultado com o endereço-alvo:

chave privada k
     │
     ▼   multiplicação na curva elíptica
   k · G ─────────────►  ponto (x, y) na curva secp256k1
     │
     ▼   compressão (33 bytes)
 chave pública
     │
     ▼   SHA-256  →  RIPEMD-160
  HASH160 (20 bytes) ───────────►  == alvo ?  🎯

O truque de desempenho é fazer tudo dentro de um só kernel, para milhares de chaves ao mesmo tempo, sem nunca voltar para a CPU no meio do caminho.

---

🚀 Começando rápido

Pré-requisitos (já instalados nesta máquina): GPU NVIDIA RTX 4060, CUDA Toolkit 13.3, Visual Studio 2026 (fornece o cl.exe), CMake ≥ 3.24 e Ninja. Detalhes e por que essas versões em Aprofundando.

Sem GPU NVIDIA? O CUDA é opcional — o projeto também compila e roda 100% na CPU. Veja Sem GPU? Modo CPU. Aí só o cl.exe (ou qualquer compilador C++), o CMake e o Ninja são necessários.

1. Prepare o terminal (o nvcc precisa do cl.exe do Visual Studio no PATH):

& "C:\Program Files\Microsoft Visual Studio\18\Community\Common7\Tools\Launch-VsDevShell.ps1" -Arch amd64 -HostArch amd64 -SkipAutomaticLocation

No VSCode, basta abrir um terminal com o perfil "PowerShell CUDA (VS Dev)" — ele já carrega esse ambiente.

2. Compile:

cmake -B build -G Ninja
cmake --build build

Ou, equivalente, pelos presets: cmake --preset default e cmake --build --preset default. Há também --preset profile (instrumentação para o Nsight, em build-profile/) e --preset no-cuda (força a configuração só-CPU mesmo com o nvcc instalado).

O configure imprime cuda-labs: build type = Release. Release é o padrão — sem isso o CMake escolheria Debug sozinho no MSVC, o que não quebra nada mas custa ~6–8× no backend CPU. Para depurar, passe -DCMAKE_BUILD_TYPE=Debug explicitamente.

3. Valide o build — os 13 drivers de conformidade, via CTest:

ctest --test-dir build --output-on-failure

4. Veja funcionando — a CLI recuperando a chave do puzzle #1:

Set-Location .\python
python main.py --preset puzzle-1-demo

---

🖥️ Usando a CLI

A CLI vive em python/ e fala com o backend CUDA por uma DLL (sem etapa de compilação do lado Python). Instale as dependências uma vez e rode a partir de python/:

Set-Location .\python
python -m pip install -r requirements.txt      # rich

Buscar uma chave

Escolha o alvo por número de puzzle (--wallet) ou por endereço/hash160 (--target). Ao achar, a chave aparece em destaque (hex, decimal e WIF):

┌────────────────── keyhunt - run summary ──────────────────┐
│          Device  NVIDIA GeForce RTX 4060 (sm_89, 24 SMs)  │
│          Target  1BgGZ9tcN4rm9KBzDn7KprQz87SZ26SAMH       │
│         HASH160  751e76e8199196d454941c45d1b3a323f1433bd6 │
│          Puzzle  #1                                       │
│           Range  [0x1, 0x1]                               │
│          Tuning  library defaults                         │
│ Already scanned  0.0000%  (0 keys)                        │
└───────────────────────────────────────────────────────────┘

window 1  100%|██████████| 1.00/1.00 [3.1 Mk/s, ETA 0s]

┌─────────────────────────────── ** PRIVATE KEY FOUND ** ───────────────────────────────┐
│           Address  1BgGZ9tcN4rm9KBzDn7KprQz87SZ26SAMH                                 │
│ Private key (hex)  0x0000000000000000000000000000000000000000000000000000000000000001 │
│ Private key (dec)  1                                                                  │
│  WIF (compressed)  KwDiBf89QgGbjEhKnhXJuH7LrciVrZi3qYjgd9M7rFU73sVHnoWn               │
└───────────────────────────────────────────────────────────────────────────────────────┘

Numa busca longa a barra mostra %, velocidade instantânea, média e ETA ao vivo. O progresso é salvo por alvo: se você parar com Ctrl-C e rodar de novo, ele retoma de onde parou em vez de reescanear o mesmo trecho.

Com --rng, cada janela cai num ponto aleatório do intervalo. O --rng-mode extend é uma variação que continua a partir do fim de um trecho já escaneado — as janelas se fundem no registro de progresso em vez de fragmentá-lo, mantendo o JSON com tamanho ~estável.

Quando a chave é achada, ela também fica marcada no arquivo de progresso do alvo, no campo "found" (null enquanto ninguém achou; a chave em hex de 64 dígitos depois). É o que faz uma execução posterior no mesmo alvo mostrar Key found: yes - recovered in an earlier run em vez de no — os intervalos escaneados sozinhos não conseguem dizer isso, porque um acerto interrompe a janela no meio e deixa um intervalo parcial idêntico ao de um Ctrl-C. O registro definitivo da chave (hex, decimal e WIF) continua sendo o data/results.json.

{
 "address": "1BgGZ9tcN4rm9KBzDn7KprQz87SZ26SAMH",
 "hash160": "751e76e8199196d454941c45d1b3a323f1433bd6",
 "found": "0000000000000000000000000000000000000000000000000000000000000001",
 "updated": "2026-08-06T22:44:51+00:00",
 "intervals": [["1","1"]]
}

Arquivos antigos (sem o campo) continuam carregando normalmente — a ausência é lida como null. O --join também preserva o campo: se qualquer arquivo de entrada tiver a chave, ela sobrevive na saída.

Listar os puzzles

python main.py --list-puzzles

                              Bitcoin puzzles
┌─────┬─────────────────────────────┬───────────────┬──────────────┐
│   # │ Address                     │         Range │ Size         │
├─────┼─────────────────────────────┼───────────────┼──────────────┤
│   1 │ 1BgGZ9tcN4rm9KBzDn7KprQz87… │    [0x1, 0x1] │ 2^0 (1)      │
│   2 │ 1CUNEBjYrCn2y1SdiUMohaKUi4… │    [0x2, 0x3] │ 2^1 (2)      │
│   3 │ 19ZewH8Kk1PDbSNdJ97FP4EiCj… │    [0x4, 0x7] │ 2^2 (4)      │
│   4 │ 1EhqbyUMvvs7BfL8goY6qcPbD6… │    [0x8, 0xf] │ 2^3 (8)      │
│   5 │ 1E6NuFjCi27W5zoXg8TRdcSRq8… │  [0x10, 0x1f] │ 2^4 (16)     │
│   6 │ 1PitScNLyp2HCygzadCh7FveTn… │  [0x20, 0x3f] │ 2^5 (32)     │
└─────┴─────────────────────────────┴───────────────┴──────────────┘

(… e assim por diante até o puzzle #160.)

Ajustar a GPU automaticamente (--tune)

Faz um grid search de lanes × stride × gpu-batch, mede a velocidade de cada combinação e mostra a melhor — com aquecimento, mediana entre waves e ordem alternada para não ser enganado por variação térmica, tudo automático:

python main.py --tune --wallet 71

                 tuning results (top 4 of 4)
┌───┬───────┬────────┬───────────┬─────────┬────────┬───────┐
│ # │ lanes │ stride │ gpu-batch │ Mkeys/s │ passes │ clock │
├───┼───────┼────────┼───────────┼─────────┼────────┼───────┤
│ 1 │ 32768 │   6144 │       384 │  1074.3 │      1 │  2865 │
│ 2 │ 32768 │   6144 │       256 │  1018.0 │      1 │  2865 │
│ 3 │ 32768 │   2048 │       384 │   979.5 │      1 │  2865 │
│ 4 │ 32768 │   2048 │       256 │   966.8 │      1 │  2865 │
└───┴───────┴────────┴───────────┴─────────┴────────┴───────┘

Os resultados também vão para benchmark_results/ (CSV + JSON). Depois é só passar o combo vencedor para a busca com --lanes / --stride / --gpu-batch.

Conferir que tudo funciona (--test)

Recupera 42 chaves de puzzles já resolvidos (puzzles 1 a 42) e confere cada WIF contra uma implementação independente — um teste de ponta a ponta do backend, sem efeitos colaterais:

┌────┬───────────┬───────────┬─────┬─────────┬────────┐
│  # │ key (hex) │ recovered │ wif │ Mkeys/s │ result │
├────┼───────────┼───────────┼─────┼─────────┼────────┤
│  1 │         1 │ ok        │ ok  │     291 │ PASS   │
│  2 │         3 │ ok        │ ok  │     320 │ PASS   │
│  3 │         7 │ ok        │ ok  │     410 │ PASS   │
│  4 │         8 │ ok        │ ok  │     417 │ PASS   │
│  5 │        15 │ ok        │ ok  │     318 │ PASS   │
└────┴───────────┴───────────┴─────┴─────────┴────────┘
                       (… 42 vetores no total)
┌─ sanity result ─┐
│ 42/42 passed    │
└─────────────────┘

Rodar as suítes de testes

São duas, independentes. A do C++/CUDA roda pelo CTest, da raiz do projeto — são os 13 drivers de conformidade, cada um cruzando kernel × core × gold vector:

ctest --test-dir build --output-on-failure

A do frontend Python é independente da placa — só unittest da stdlib, sem venv e sem DLL. De python/:

python -m unittest discover -s tests -v

São 1046 testes (~7 s); os que exigem uma DLL pulam sozinhos quando ela não foi compilada.

Olhe a CONTAGEM, não só o OK. Um arquivo de teste que falha no import é reportado como um erro enquanto todos os testes dele somem em silêncio. Já aconteceu duas vezes aqui (112 e 170 testes desaparecidos). Se o número vier abaixo de 1046, tem arquivo não carregando.

Esse único comando cobre os dois backends: os testes de paridade são gerados uma vez por backend (KnownKeyRecovery_cuda, KnownKeyRecovery_cpu, ...) e cada um dirige a sua própria DLL no mesmo processo. Não é preciso definir KEYHUNT_BACKEND — quem só tem um backend compilado vê os do outro pularem, dizendo qual biblioteca falta.

Escolher a GPU (--gpu)

Num PC com mais de uma placa, --list-gpus mostra os índices e --gpu N escolhe em qual rodar. Para usar duas ao mesmo tempo, rode duas instâncias em terminais separados — cada uma na sua:

python main.py --list-gpus                 # lista as GPUs e seus índices
python main.py --preset wallet-71 --gpu 0  # 1a instância -> GPU 0
python main.py --preset wallet-71 --gpu 1  # 2a instância -> GPU 1 (outro terminal)

⚠️ Duas instâncias no mesmo alvo precisam de --progress-suffix (logo abaixo), senão uma sobrescreve o arquivo de progresso da outra.

Duas instâncias no mesmo alvo (--progress-suffix)

O progresso fica em data/progress/<hash160>.json e é reescrito por inteiro a cada gravação, a partir da memória do processo. Então duas instâncias no mesmo alvo apagam o trabalho uma da outra: a última a gravar vence. --progress-suffix NOME dá a cada uma o seu próprio arquivo:

python main.py --preset wallet-71 --gpu 0 --progress-suffix gpu0   # -> <hash160>.gpu0.json
python main.py --preset wallet-71 --gpu 1 --progress-suffix gpu1   # -> <hash160>.gpu1.json
python main.py --preset wallet-71 --backend cpu --progress-suffix cpu

# no fim, consolide tudo no arquivo principal:
python main.py --join data/progress/<hash160>.gpu0.json data/progress/<hash160>.gpu1.json

• Nada de perda de dados: cada processo só lê e escreve o arquivo dele. O --join sem --out grava em data/progress/<hash160>.json e mescla o que já estiver lá — é idempotente, dá para rodar quantas vezes quiser.
• Uma instância não enxerga a outra durante a corrida. O preço é retrabalho quando as duas planejam a mesma janela, nunca uma chave pulada. Com --rng (janelas sorteadas num intervalo de 2^70) a chance é desprezível; no modo sequencial, dê faixas disjuntas a cada uma com --from / --range-size.
• Depois do --join, os arquivos por instância continuam com o conteúdo antigo. Se reusar o mesmo sufixo, aquela instância recomeça do progresso dela, não do consolidado — o que é seguro (re-varre, não pula). Apague-os depois de consolidar se quiser evitar isso.
• Sem a flag, nada muda: continua o <hash160>.json compartilhado de sempre. Aceita letras, dígitos, _ e - (1–32 caracteres) — qualquer outra coisa é recusada, não corrigida em silêncio, porque o valor decide qual arquivo a execução lê. Vale como chave de preset ("progress_suffix": "gpu0") e aparece no painel de resumo e no cabeçalho do Telegram — é assim que se distinguem duas instâncias reportando pro mesmo chat.

Se o ganho é throughput, cheque o hardware primeiro: duas instâncias na mesma GPU não aceleram nada — o kernel já satura a placa. Isso vale para duas GPUs, GPU + --backend cpu, ou duas máquinas.

Desligar o PC no fim (--shutdown)

Para varreduras longas sem ninguém olhando: --shutdown N desliga a máquina N segundos depois que a execução termina — mas só se ela terminar sozinha.

python main.py --preset wallet-71-biggest -bc 40 --shutdown 30
python main.py --preset wallet-71 -tl 8h --shutdown 60   # roda 8h e desliga

Como a execução terminou	Desliga?
O plano acabou (--range-count cumprido)	✅ sim
O intervalo foi 100% varrido	✅ sim
--time-limit estourou	✅ sim
Você apertou Ctrl-C	❌ não
A chave foi encontrada	❌ não — essa é a hora de estar no console
Erro, ou nem chegou a começar	❌ não

É exatamente essa condicional que faz a flag valer mais do que um shutdown /s /t 0 na linha seguinte de um .bat: aquele dispara de qualquer jeito, inclusive depois de um Ctrl-C ou de um erro na largada.

O mínimo é 5 segundos — a contagem regressiva é a sua única chance de cancelar (shutdown /a no Windows, shutdown -c fora dele). Um valor menor é recusado com uma mensagem, não arredondado em silêncio. Quando dispara, a linha warn shutting down in 30s - run 'shutdown /a' to cancel aparece no console.

O progresso já está no disco antes do disparo (o flush do final da varredura vem primeiro), então cancelar o desligamento não perde nada. Aceita também o formato de duração do --time-limit (--shutdown 1m30s), vale como chave de preset ("shutdown": 30), e funciona nos dois backends. Vale só para a busca: --tune e --test ignoram a flag, como já fazem com --telegram e --time-limit.

Referência rápida de comandos

Quero…	Comando
Buscar num puzzle	python main.py --wallet 71 --range-size 1e9 --rng
Buscar num endereço	python main.py --target 1BgGZ9… --from 0x1 --range-size 0x2000
Usar um preset	python main.py --preset wallet-71
Listar puzzles	python main.py --list-puzzles
Listar / escolher a GPU	python main.py --list-gpus · --gpu 1
Duas instâncias no mesmo alvo	--gpu 0 --progress-suffix gpu0 / --gpu 1 --progress-suffix gpu1, depois --join
Achar a melhor config de GPU	python main.py --tune --wallet 71
Rodar a noite toda e desligar	python main.py --preset wallet-71 -tl 8h --shutdown 30
Testar o backend	python main.py --test
Rodar a suíte Python	python -m unittest discover -s tests -v
Rodar a suíte C++/CUDA	ctest --test-dir build --output-on-failure (da raiz)

---

🧩 Sem GPU? Modo CPU (funciona sem NVIDIA/CUDA)

O projeto roda mesmo sem GPU NVIDIA e sem o CUDA Toolkit. Existe um segundo backend, escrito em C++ puro (src/search/cpu/), que reaproveita exatamente os mesmos cores matemáticos (via o caminho host que já era conferido contra os gold vectors) e expõe a mesma ABI C. Ou seja: toda a CLI — progresso por alvo, retomada, --join, --time-limit, presets, monitor Telegram — funciona sem alteração; só o motor por baixo muda.

• Build: o CMake detecta o CUDA automaticamente. Com CUDA presente, compila os dois backends (keyhunt_c.dll + keyhunt_cpu.dll); sem CUDA, configura como projeto C++ puro e compila só o backend CPU (keyhunt_cpu.dll + os labs de CPU). Nenhuma edição necessária.

• Escolher o backend: --backend {auto,cpu,cuda} ou a variável de ambiente KEYHUNT_BACKEND, nesta ordem de precedência:

  --backend > preset > KEYHUNT_BACKEND > auto

  O padrão auto usa a GPU se houver uma com CUDA, senão cai para a CPU. Exportar KEYHUNT_BACKEND vale para todos os comandos (o run, o --list-gpus, o --test); uma flag ou um preset sobrepõem a variável, mas nada a sobrescreve por conta própria.

python main.py --list-gpus                        # na CPU, mostra "CPU (<modelo>, N threads)"
python main.py --backend cpu --test               # suíte de sanidade no backend CPU (42/42)
python main.py --backend cpu --preset puzzle-1-demo
$env:KEYHUNT_BACKEND = "cpu"; python main.py --preset puzzle-1-demo   # idem, pela variável

Expectativa honesta: a CPU é ~100× mais lenta que a RTX 4060 — medido aqui: 10,7 Mk/s num Ryzen 7 5800X (16 threads) contra ~1.250 Mk/s na 4060; em CPUs mais modestas a distância passa de 1000×. A GPU vence pelo paralelismo massivo (dezenas de milhares de lanes); a CPU tem núcleos × threads. O backend CPU usa std::thread + a mesma inversão de Montgomery em lote — é "tão otimizado quanto um bom código de CPU", não uma throughput parecida com a da GPU. Mesma disciplina de correção (gold vectors + oráculo) e mesmo escopo (só chaves conhecidas).

⚠️ Esses números valem para um build Release, que é o padrão desde que o CMakeLists.txt passou a defini-lo. Um build Debug custa ~6–8× no backend CPU (o código de device do CUDA é pouco afetado). O configure imprime cuda-labs: build type = ... — confira se algo parecer lento.

---

📈 Performance

A busca saiu de ~60 Mkeys/s na primeira versão para ~1.250 Mkeys/s (1,25 Gk/s — mais de um bilhão de chaves por segundo) — cerca de 21× mais rápido — ao longo de três ciclos de otimização guiados por profiler mais uma varredura conjunta com --tune, chegando ao mesmo nível de referências como a VanitySearch.

O número vale na geometria tunada lanes=49152 stride=8192 gpu_batch=768 (a dos presets wallet-71*). A configuração anterior, 32768/2048/384, entregava ~1.050 Mk/s: era um ótimo local, achado varrendo um botão de cada vez. Foi a varredura conjunta dos três que encontrou o ponto melhor.

Cada ganho foi medido, não adivinhado — várias ideias "óbvias" na verdade deixaram o kernel mais lento e foram descartadas. Os três saltos que ficaram:

Otimização	O que muda	Ganho
Incremento afim + inversão em lote	troca a soma de pontos cara por +G barato, amortizando a inversão entre milhares de chaves	~1,4×
Coalescência de memória (SoA)	acessos à memória alinhados, sem desperdício de banda	~2,4×
Simetria ±jG	uma inversão gera dois pontos, dando ao compilador trabalho paralelo para esconder latência	~1,6×

As RTX 4060 tem limite de 115 W, então o clock e a velocidade oscilam com a temperatura — por isso o --tune mede em janela quente com mediana. Números aqui são da RTX 4060.

---

🔍 Como sei que está correto

Um erro de sinal ou de endianness na matemática da curva não trava o programa — ele só gera silenciosamente um endereço errado. Comparar GPU contra CPU não pega isso, porque as duas rodam exatamente o mesmo código.

A rede de segurança é um oráculo independente em Python (tools/secp256k1_oracle.py, big-ints puros + hashlib) que calcula os valores esperados por um caminho totalmente diferente. Esses valores viram gold vectors embutidos em cada teste, e toda camada é conferida contra eles. O marco de confiança: hash160(1) = 751e76e8…, que é exatamente o HASH160 público do puzzle #1 — um número que dá para verificar em qualquer lugar.

Os gold vectors são versionados junto com o código; o build em CUDA nunca roda Python. Só é preciso regerar se você mexer no oráculo:

python tools\gen_vectors.py

---

🔒 Escopo & ética

Este é um exercício de aprendizado e correção — não uma ferramenta para atacar carteiras.

• Todos os testes ponta a ponta buscam apenas chaves que já conhecemos: puzzles que a gente mesmo gera (escolhe k, calcula o endereço, confirma que o kernel reencontra k) e o puzzle público #1 (chave 1, sem saldo).
• O espaço de 256 bits não é força-brutável — é a própria base da segurança do Bitcoin. Nada aqui mira, ou poderia mirar, endereços com fundos.
• As referências (BitCrack, VanitySearch, Keyhunt) entram só como comparação de arquitetura.

---

🗂️ Estrutura do projeto

cuda-labs/
├── src/                todo o C++/CUDA — é também o ÚNICO include root
│   ├── common/         portability/cuda_compat.h (shim host) + host/ + cuda/ (CUDA_CHECK, benchmark)
│   ├── hashing/        algoritmos de hash (SHA-256, RIPEMD-160, HASH160)   → hashing    (CUDA)
│   ├── secp256k1/      u256 → corpo → curva → escalar → batchinv → pubkey → addr → secp256k1 (CUDA)
│   ├── search/         os dois backends de busca, lado a lado
│   │   ├── cuda/       kernel fundido tunado (~1.250 Mk/s)                → search_cuda (CUDA)
│   │   └── cpu/        varredura em threads, C++ puro, sem CUDA           → search_cpu
│   └── capi/           ABI C (extern "C"): keyhunt_c.dll + keyhunt_cpu.dll, mesma ABI
├── tests/              drivers de conformidade — cruzam GPU × CPU × gold; vectors/ tem os gold
├── python/             frontend: pykeyhunt (ctypes, escolhe o backend) + keyhunt_cli + testes
└── tools/              oráculo de referência em Python + gerador de gold vectors

Um único include root. Todo include interno é escrito qualificado a partir de src/ — #include "secp256k1/field/field.cuh", #include "common/portability/cuda_compat.h" — então a linha do include já diz onde a dependência mora, sem include-path por diretório.

Os cores em secp256k1//hashing/ são __host__ __device__ e compilam nos dois backends: sob o nvcc (GPU) e como C++ puro (CPU, via src/common/portability/cuda_compat.h, que anula os qualificadores CUDA no host). keyhunt_c.dll e keyhunt_cpu.dll exportam a mesma ABI.

Os drivers de conformidade ficam em build/tests/ — cada um imprime PASS/FAIL por vetor e sai com código 0 (passou) ou 1 (falhou). Estão registrados no CTest, então rodar todos é:

ctest --test-dir build --output-on-failure

Um driver isolado continua funcionando como executável (.\build\tests\test-keyhunt.exe), e ctest -R field roda só os que casam com o nome.

---

🔧 Aprofundando (para devs)

Versões da toolchain (e por quê)

• CUDA 13.3 é o toolkit no PATH. O 13.0, também instalado, rejeita o compilador do VS 2026 (MSVC 19.50) — não use.
• Visual Studio 2026 fornece o cl.exe que o nvcc usa como host compiler. Ele não fica no PATH de um PowerShell comum; daí o Launch-VsDevShell.ps1 do Começando rápido. Para confirmar: Get-Command cl deve apontar para …\MSVC\14.50.…\bin\HostX64\x64\cl.exe.
• O Launch-VsDevShell.ps1 pode imprimir um aviso vswhere.exe não reconhecido — pode ignorar.

Arquitetura de GPU — native (um build roda em qualquer placa da máquina)

O build usa CMAKE_CUDA_ARCHITECTURES=native: detecta todas as GPUs instaladas e gera um cubin nativo por arquitetura no mesmo binário (ex.: RTX 4060 sm_89 + RTX 50xx sm_120 lado a lado). É cubin puro, sem PTX embutido — necessário porque o driver instalado é mais antigo que o toolkit 13.3 e rejeitaria PTX em tempo de execução ("the provided PTX was compiled with an unsupported toolchain"). Trocou/adicionou uma GPU? Basta recompilar — sem editar a arquitetura.

CMAKE_CUDA_ARCHITECTURES é cache: um build/ já existente mantém a arquitetura da primeira configuração. Para o native reavaliar (ex.: depois de plugar outra placa), apague build/CMakeCache.txt ou configure num diretório novo.

Integração Python (a DLL e a API)

A DLL keyhunt_c.dll (CUDA) faz parte do build normal (cmake --build build já a gera em build/capi/, junto com keyhunt_cpu.dll). O lado Python é puro ctypes, sem compilar nada. A DLL CUDA depende do cudart, que o loader (python/pykeyhunt/_ffi.py) encontra via CUDA_PATH\bin. O loader escolhe o backend por KEYHUNT_BACKEND (auto/cuda/cpu; padrão auto = GPU se houver, senão CPU) — veja pykeyhunt.backend(). KEYHUNT_DLL aponta uma biblioteca específica e ganha de tudo. A variável é lida uma vez, no import: para trocar de backend dentro do mesmo processo, importe o módulo sob outro nome.

Além da CLI, dá para usar a API diretamente:

from pykeyhunt import Search

with Search("1BgGZ9tcN4rm9KBzDn7KprQz87SZ26SAMH", base_key=1, count=1 << 12) as s:
    for p in s:                       # cada iteração avança a busca uma "wave"
        print(f"{p.fraction:6.1%}  {p.rate_kps/1e6:5.1f} Mkeys/s")
        if p.found:
            print("chave privada:", p.found_key_int)   # -> 1
            break

A suíte de testes cobre o binding e todo o frontend -- de python/, rode python -m unittest discover -s tests -v. São 1046 testes (~7 s): os unitários rodam em qualquer máquina, e os de integração rodam contra cada backend compilado (pulando os das DLLs ausentes).

Adicionar um novo lab / algoritmo

Código de algoritmo mora numa lib sob src/ (src/hashing/ para hashes, src/secp256k1/ para a curva, src/search/ para os backends de busca); tests/ só tem drivers.

1. Core em src/secp256k1/<nome>/<nome>.cuh — um core inline __host__ __device__ (para inlinar no kernel; sem -rdc). Kernels + launchers vão no <nome>.cu, que entra no CMakeLists.txt da lib.
2. Gold — adicione um gen_<nome>() em tools/gen_vectors.py, rode o gerador e versione o tests/vectors/gold_<nome>.h.
3. Driver tests/<área>/<nome>/main.cu — cruza kernel × core × gold e retorna ≠0 em falha. Registre com add_conformance_test(test-<nome> <área>/<nome>/main.cu <libs...>) no tests/CMakeLists.txt; isso já cria o executável em build/tests/ e o inscreve no CTest.

Includes. Escreva sempre qualificado a partir de src/ (#include "secp256k1/field/field.cuh"); o gold é #include "vectors/gold_<nome>.h".

Por que o core fica no header? Para o compilador inlinar o algoritmo dentro do kernel — é isso que garante o throughput. Um split clássico exigiria separable compilation (-rdc) e não inlinaria entre unidades de tradução.

A documentação interna detalhada (decisões, gotchas, histórico de otimização) fica em .agents/.

Build limpo

O diretório build/ é descartável e ignorado pelo git:

Remove-Item -Recurse -Force build
cmake -B build -G Ninja ; if ($?) { cmake --build build }

---

_{Projeto pessoal de estudo de CUDA. Referências de arquitetura: BitCrack, VanitySearch, Keyhunt. Sem afiliação; sem uso contra endereços com fundos.}