Генерация публичного ключа из приватного с помощью GPU OpenCL secp256k1

Я хотел бы вычислить публичный ключ, выполняя вычисления с помощью gpu. Код работает, но сравнивая вычисления с эталонными становится понятно, что ошибка либо в реализации функций вычислений в secp256k1 (что наверняка невозможно), либо в правильности передачи значений на gpu. Прошу помощи и подсказок, только учусь и нуждаюсь в объяснениях. Библиотеки взяты отсюда: https://github.com/bernardladenthin/BitcoinAddressFinder/blob/b71531a24af1037c71ba7a13d30ffb0c90956444/src/main/resources/inc_ecc_secp256k1.cl https://github.com/bernardladenthin/BitcoinAddressFinder/blob/b71531a24af1037c71ba7a13d30ffb0c90956444/src/main/resources/inc_ecc_secp256k1.h

содержимое kernel.cl

#include "inc_ecc_secp256k1.cl"

#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
#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
#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
#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
#define PUBLIC_KEY_LENGTH_WITHOUT_PARITY 8
#define PUBLIC_KEY_LENGTH_WITH_PARITY 9
#define PRIVATE_KEY_LENGTH 8

// Эта функция выполняет следующее:
// Инициализирует массивы для хранения значений базовой точки, публичного ключа и приватного ключа.
// Копирует значения базовой точки G и приватного ключа k в локальные массивы.
// Вызывает функцию parse_public для преобразования значений базовой точки.
// Вызывает функцию point_mul для выполнения умножения точки на приватный ключ, чтобы получить публичный ключ.
// Копирует результат вычислений в выходной массив.
// Отладочный буфер для возврата промежуточных данных
// debug_data[0:7]   = Приватный ключ
// debug_data[8:15]  = Координата X
// debug_data[16]    = gpu_parity

__kernel void generateKeysKernel_parse_public(__global u32 *r, __global const u32 *k, __global u32 *debug_data)
{
    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;

    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];
    
    // Загружаем приватный ключ и сохраняем в отладочный буфер
    for (int i = 0; i < PRIVATE_KEY_LENGTH; ++i) {
        k_local[i] = k[i];
        debug_data[i] = k[i]; // Отладка: сохраняем приватный ключ
    }

    // Парсим базовую точку и умножаем её на приватный ключ
    parse_public(&g_xy_local, g_local);
    point_mul(r_local, k_local, &g_xy_local);

    // Устанавливаем gpu_parity в зависимости от чётности Y
    r_local[8] = (u32)((g_xy_local.xy[8] & 1) == 0 ? 0x02 : 0x03);

    // 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];
    
    // Отладка: сохраняем координату X в debug_data
    for (int i = 0; i < PUBLIC_KEY_LENGTH_WITHOUT_PARITY; ++i) {
        debug_data[8 + i] = r_local[i];
    }

    // Отладка: сохраняем gpu_parity (последний элемент массива)
    debug_data[16] = r_local[8]; // Сохраняем для анализа в Python

    // Записываем результат в выходной массив r
    for (int i = 0; i < PUBLIC_KEY_LENGTH_WITH_PARITY; ++i) {
        r[i] = r_local[i];
    }
}

и собственно сам код python:

import os
import numpy as np
import pyopencl as cl
import hashlib

def bswap32(v):
    return ((v >> 24) & 0xff) | (((v >> 16) & 0xff) << 8) | (((v >> 8) & 0xff) << 16) | ((v & 0xff) << 24)

def assemble_big_endian(words):
    """Собирает число из списка 32-битных слов в big-endian порядке"""
    result = 0
    for word in words:
        result = (result << 32) | bswap32(word)
    return result

def main():
    kernel_filename = "kernel.cl"
    if not os.path.exists(kernel_filename):
        print(f"Файл {kernel_filename} не найден в текущей директории.")
        return
    with open(kernel_filename, "r", encoding="utf-8") as f:
        kernel_source_raw = f.read()

    # Пролог ядра
    definitions = """
    #define u32 uint
    #define u64 ulong
    #define DECLSPEC
    #define PRIVATE_AS __private
    #define GLOBAL_AS __private
    #define SECP256K1_TMPS_TYPE __private
    #define __ENDIAN_LITTLE__ 1
    """
    kernel_source = definitions + "\n" + kernel_source_raw

    # Абсолютный путь к inc‑файлам
    inc_dir = os.path.abspath(".")
    print("Абсолютный путь к директории с inc‑файлами:", inc_dir)

    # Ввод приватного ключа
    priv_key_str = input("Введите приватный ключ (число от 1 до 2^256): ")
    try:
        priv_key_int = int(priv_key_str)
    except ValueError:
        print("Некорректный ввод. Нужно ввести число.")
        return
    if not (1 <= priv_key_int < 2**256):
        print("Приватный ключ вне допустимого диапазона.")
        return

    print("[Шаг 1] Приватный ключ (BigUint):", priv_key_int)
    priv_key_hex = format(priv_key_int, "064x")
    print("[Шаг 2] Приватный ключ (32 байта, hex):", priv_key_hex)

    # Формируем массив из 8 uint32 (big-endian, без переворота)
    k_list = [int(priv_key_hex[i:i+8], 16) for i in range(0, 64, 8)]
    k = np.array(k_list, dtype=np.uint32)
    print("Отправляем в GPU приватный ключ (uint32, big-endian):", k.tolist())

    # Настраиваем OpenCL
    try:
        ctx = cl.create_some_context()
    except Exception as e:
        print("Ошибка создания OpenCL-контекста:", e)
        return
    queue = cl.CommandQueue(ctx)
    mf = cl.mem_flags

    # Буферы
    k_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=k)
    r = np.zeros(9, dtype=np.uint32)  # 8 слов для X + 1 для gpu_parity
    r_buf = cl.Buffer(ctx, mf.WRITE_ONLY, r.nbytes)

    # Отладочный буфер (8 слов приватного ключа + 8 слов X + 1 слово gpu_parity = 17 слов)
    debug_data = np.zeros(17, dtype=np.uint32)
    debug_buf = cl.Buffer(ctx, mf.WRITE_ONLY, debug_data.nbytes)

    # Компилируем программу
    try:
        prg = cl.Program(ctx, kernel_source).build(options=["-I" + inc_dir])
    except Exception as e:
        print("Ошибка компиляции OpenCL-программы:", e)
        return

    # Запускаем ядро generateKeysKernel_parse_public
    kernel = prg.generateKeysKernel_parse_public
    kernel.set_args(r_buf, k_buf, debug_buf)
    cl.enqueue_nd_range_kernel(queue, kernel, (1,), None)
    cl.enqueue_copy(queue, r, r_buf)
    cl.enqueue_copy(queue, debug_data, debug_buf)
    queue.finish()

    # Выводим содержимое отладочного буфера
    print("\n[Отладка GPU]:")
    print("Приватный ключ (uint32, внутри ядра):", debug_data[0:8].tolist())
    print("Координата X (uint32, внутри ядра):", debug_data[8:16].tolist())
    print("gpu_parity (uint32, внутри ядра):", hex(debug_data[16]))

    # Обрабатываем результат из GPU
    gpu_words = r[:8].tolist()
    x_gpu = assemble_big_endian(gpu_words)
    gpu_parity = r[8] & 0xff
    print("GPU результат (массив r):", r.tolist())
    print("Интерпретированное значение X (от GPU):", hex(x_gpu))
    print("Значение gpu_parity (ожидается 0x02 или 0x03):", hex(gpu_parity))

    # Вычисляем Y
    p = 0xFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFFEFFFFFC2F
    y2 = (pow(x_gpu, 3, p) + 7) % p
    def mod_sqrt(a, p):
        return pow(a, (p + 1) // 4, p)
    y = mod_sqrt(y2, p)
    if (y % 2 == 0 and gpu_parity == 0x03) or (y % 2 == 1 and gpu_parity == 0x02):
        y = p - y

    # Публичные ключи
    x_bytes = x_gpu.to_bytes(32, byteorder="big")
    y_bytes = y.to_bytes(32, byteorder="big")
    uncompressed_pub_key = b'\x04' + x_bytes + y_bytes
    print("[Шаг 3] Несжатый публичный ключ (65 байт, hex):", uncompressed_pub_key.hex())

    prefix = b'\x02' if (y % 2 == 0) else b'\x03'
    compressed_pub_key = prefix + x_bytes
    print("[Шаг 4] Сжатый публичный ключ (33 байта, hex):", compressed_pub_key.hex())

    sha256_hash = hashlib.sha256(compressed_pub_key).hexdigest()
    print("[Шаг 5] SHA-256 сжатого публичного ключа:", sha256_hash)

    ripemd160_hash = hashlib.new("ripemd160", bytes.fromhex(sha256_hash)).hexdigest()
    print("[Шаг 6] RIPEMD-160 (HASH160):", ripemd160_hash)

if __name__ == "__main__":
    main()