// gpu_scan_secp256k1_full.cu
// Full GPU-based secp256k1 -> SHA256 -> RIPEMD160 scanner using actual CUDA ECC and hash implementations
// Dependencies: secp256k1.cuh, sha256.cuh, ripemd160.cuh, ptx.cuh

#include <stdio.h>
#include <stdint.h>
#include <cuda.h>
#include <cuda_runtime.h>

#include "secp256k1.cuh"
#include "sha256.cuh"
#include "ripemd160.cuh"
#include "ptx.cuh"

__device__ bool check_prefix(const uint8_t* h160) {
    return h160[0] == 0xf6 && h160[1] == 0xf5 && h160[2] == 0x43;
}

__global__ void kernel_scan(uint64_t base, uint64_t total_keys) {
    uint64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= total_keys) return;

    uint64_t priv = base + idx;

    uint256_t seckey;
    set_uint256_from_uint64(seckey, priv);

    secp256k1_pubkey_t pubkey;
    secp256k1_scalar_multiply(pubkey, seckey);

    uint8_t pubkey33[33];
    secp256k1_compress_pubkey(pubkey33, pubkey);

    uint8_t sha_digest[32];
    sha256(pubkey33, 33, sha_digest);

    uint8_t h160[20];
    ripemd160(sha_digest, 32, h160);

    if (check_prefix(h160)) {
        printf("[MATCH] priv=0x%016llx h160=%.2x%.2x%.2x...\n",
               (unsigned long long)priv, h160[0], h160[1], h160[2]);
    }
}

int main(int argc, char** argv) {
    if (argc != 4) {
        printf("Usage: ./gpu_scan <start> <count> <threads_per_block>\n");
        return 1;
    }

    uint64_t start = strtoull(argv[1], nullptr, 0);
    uint64_t count = strtoull(argv[2], nullptr, 0);
    int threads = atoi(argv[3]);

    uint64_t blocks = (count + threads - 1) / threads;

    cudaEvent_t t_start, t_end;
    cudaEventCreate(&t_start);
    cudaEventCreate(&t_end);
    cudaEventRecord(t_start);

    kernel_scan<<<blocks, threads>>>(start, count);
    cudaDeviceSynchronize();

    cudaEventRecord(t_end);
    cudaEventSynchronize(t_end);

    float ms = 0;
    cudaEventElapsedTime(&ms, t_start, t_end);
    printf("\nScanned %llu keys in %.2f s (%.2f MH/s)\n",
           (unsigned long long)count, ms / 1000.0, count / (ms * 1e3));

    return 0;
}