Generate a public key from a private key with OpenCL for secp256k1

I'ld like to calculate a public key from a private key. It must work using OpenCL. It seems hashcat has a nice library but I'm not sure how to use it.

Other libararies are also possible. secp256k1-cl seems to be very difficult to use https://github.com/hhanh00/secp256k1-cl/blob/master/secp256k1.cl

See below an example of how it might work based on inc_ecc_secp256k1.cl.

Should I use parse_public on the known base point G to get the coordinates in a secp256k1_t object? Afterwards use point_mul with my random secret passed via k and the base point G passed via basepoint_precalculated to get the pubkey in r?

// 8+1 to make room for the parity
#define KEY_LENGTH_WITH_PARITY 9
// (32*8 == 256)
#define PRIVATE_KEY_LENGTH 8

// the base point G in compressed form
#define SECP256K1_G0 0x00000002
#define SECP256K1_G1 0x79be667e
#define SECP256K1_G2 0xf9dcbbac
#define SECP256K1_G3 0x55a06295
#define SECP256K1_G4 0xce870b07
#define SECP256K1_G5 0x029bfcdb
#define SECP256K1_G6 0x2dce28d9
#define SECP256K1_G7 0x59f2815b
#define SECP256K1_G8 0x16f81798

__kernel void generateKeysKernel(__global u32 *r, __global const u32 *k)
{
    u32 basepoint_g[KEY_LENGTH_WITH_PARITY];
    u32 u32r_local[KEY_LENGTH_WITH_PARITY];
    u32 u32k_local[PRIVATE_KEY_LENGTH];
    secp256k1_t basepoint_precalculated;

    basepoint_g[0] = SECP256K1_G0;
    basepoint_g[1] = SECP256K1_G1;
    basepoint_g[2] = SECP256K1_G2;
    basepoint_g[3] = SECP256K1_G3;
    basepoint_g[4] = SECP256K1_G4;
    basepoint_g[5] = SECP256K1_G5;
    basepoint_g[6] = SECP256K1_G6;
    basepoint_g[7] = SECP256K1_G7;
    basepoint_g[8] = SECP256K1_G8;

    // global to local
    u32k_local[0] = k[0];
    u32k_local[1] = k[1];
    u32k_local[2] = k[2];
    u32k_local[3] = k[3];
    u32k_local[4] = k[4];
    u32k_local[5] = k[5];
    u32k_local[6] = k[6];
    u32k_local[7] = k[7];

    parse_public(&basepoint_precalculated, basepoint_g);
    point_mul(u32r_local, u32k_local, &basepoint_precalculated);

    // local to global
    r[0] = u32r_local[0];
    r[1] = u32r_local[1];
    r[2] = u32r_local[2];
    r[3] = u32r_local[3];
    r[4] = u32r_local[4];
    r[5] = u32r_local[5];
    r[6] = u32r_local[6];
    r[7] = u32r_local[7];
    r[8] = u32r_local[8];
}

Solution

I've found a working solution

// finite field Fp
// p = FFFFFFFF FFFFFFFF FFFFFFFF FFFFFFFF FFFFFFFF FFFFFFFF FFFFFFFE FFFFFC2F
#define SECP256K1_P0 0xfffffc2f
#define SECP256K1_P1 0xfffffffe
#define SECP256K1_P2 0xffffffff
#define SECP256K1_P3 0xffffffff
#define SECP256K1_P4 0xffffffff
#define SECP256K1_P5 0xffffffff
#define SECP256K1_P6 0xffffffff
#define SECP256K1_P7 0xffffffff

// prime order N
// n = FFFFFFFF FFFFFFFF FFFFFFFF FFFFFFFE BAAEDCE6 AF48A03B BFD25E8C D0364141
#define SECP256K1_N0 0xd0364141
#define SECP256K1_N1 0xbfd25e8c
#define SECP256K1_N2 0xaf48a03b
#define SECP256K1_N3 0xbaaedce6
#define SECP256K1_N4 0xfffffffe
#define SECP256K1_N5 0xffffffff
#define SECP256K1_N6 0xffffffff
#define SECP256K1_N7 0xffffffff

// the base point G in compressed form for transform_public
// G = 02 79BE667E F9DCBBAC 55A06295 CE870B07 029BFCDB 2DCE28D9 59F2815B 16F81798
#define SECP256K1_G_PARITY 0x00000002
#define SECP256K1_G0 0x16f81798
#define SECP256K1_G1 0x59f2815b
#define SECP256K1_G2 0x2dce28d9
#define SECP256K1_G3 0x029bfcdb
#define SECP256K1_G4 0xce870b07
#define SECP256K1_G5 0x55a06295
#define SECP256K1_G6 0xf9dcbbac
#define SECP256K1_G7 0x79be667e

// the base point G in compressed form for parse_public
// parity and reversed byte/char (8 bit) byte order
// G = 02 79BE667E F9DCBBAC 55A06295 CE870B07 029BFCDB 2DCE28D9 59F2815B 16F81798
#define SECP256K1_G_STRING0 0x66be7902
#define SECP256K1_G_STRING1 0xbbdcf97e
#define SECP256K1_G_STRING2 0x62a055ac
#define SECP256K1_G_STRING3 0x0b87ce95
#define SECP256K1_G_STRING4 0xfc9b0207
#define SECP256K1_G_STRING5 0x28ce2ddb
#define SECP256K1_G_STRING6 0x81f259d9
#define SECP256K1_G_STRING7 0x17f8165b
#define SECP256K1_G_STRING8 0x00000098

#define SECP256K1_PRE_COMPUTED_XY_SIZE 96
#define SECP256K1_NAF_SIZE 33 // 32+1, we need one extra slot

#define PUBLIC_KEY_LENGTH_WITHOUT_PARITY 8
// 8+1 to make room for the parity
#define PUBLIC_KEY_LENGTH_WITH_PARITY 9

// (32*8 == 256)
#define PRIVATE_KEY_LENGTH 8

__kernel void generateKeysKernel_parse_public(__global u32 *r, __global const u32 *k)
{
    u32 g_local[PUBLIC_KEY_LENGTH_WITH_PARITY];
    u32 r_local[PUBLIC_KEY_LENGTH_WITH_PARITY];
    u32 k_local[PRIVATE_KEY_LENGTH];
    secp256k1_t g_xy_local;

    g_local[0] = SECP256K1_G_STRING0;
    g_local[1] = SECP256K1_G_STRING1;
    g_local[2] = SECP256K1_G_STRING2;
    g_local[3] = SECP256K1_G_STRING3;
    g_local[4] = SECP256K1_G_STRING4;
    g_local[5] = SECP256K1_G_STRING5;
    g_local[6] = SECP256K1_G_STRING6;
    g_local[7] = SECP256K1_G_STRING7;
    g_local[8] = SECP256K1_G_STRING8;

    // global to local
    k_local[0] = k[0];
    k_local[1] = k[1];
    k_local[2] = k[2];
    k_local[3] = k[3];
    k_local[4] = k[4];
    k_local[5] = k[5];
    k_local[6] = k[6];
    k_local[7] = k[7];
    
    /*
    // example private key (in)
    // hex: 68e23530deb6d5011ab56d8ad9f7b4a3b424f1112f08606357497495929f72dc
    // decimal: 47440210799387980664936216788675555637818488436833759923669526136462528967388
    // WiF
    // to generate the public key (out)
    // 025d99d81d9e731e0d7eebd1c858b1155da7981b1f0a16d322a361f8b589ad2e3b
    // hex:
    k_local[7] = 0x68e23530;
    k_local[6] = 0xdeb6d501;
    k_local[5] = 0x1ab56d8a;
    k_local[4] = 0xd9f7b4a3;
    k_local[3] = 0xb424f111;
    k_local[2] = 0x2f086063;
    k_local[1] = 0x57497495;
    k_local[0] = 0x929f72dc;
    */

    parse_public(&g_xy_local, g_local);
    point_mul(r_local, k_local, &g_xy_local);

    // local to global
    r[0] = r_local[0];
    r[1] = r_local[1];
    r[2] = r_local[2];
    r[3] = r_local[3];
    r[4] = r_local[4];
    r[5] = r_local[5];
    r[6] = r_local[6];
    r[7] = r_local[7];
    r[8] = r_local[8];
}