XBTGPUARC Programmupdate.

Words
2164
Reading
10 min
Listen
Play
1y

Es ist natürlich noch nicht Fertig. Für nun 6 Wochen Programmierabenteuer, kann sich der Fortschritt inzwischen aber vermutlich sehen lassen. Wie folgt derzeit der aktuellste Stand.

alu@king
 OS Garuda Linux x86_64
├ Kernel Linux 6.15.8-zen1-2-zen
├󰏖 Packages 1469 (pacman)[stable]
└ Shell fish 4.0.2

 DE KDE Plasma 6.4.3
├󰧨 Window Manager KWin (Wayland)
├󰧨 Login Manager sddm-autologin 0.21.0 (Wayland)
├󰉼 WM Theme plastik
├󰉼 Color Themes Breeze (Kvantum) [Qt], Sweet-Dark [GTK2/3/4]
├󰀻 System Icons breeze-dark [Qt], breeze-dark [GTK2/3/4]
├ System Fonts Fira Sans (10pt) [Qt], Fira Sans (10pt) [GTK2/3/4]
└ Terminal konsole 25.4.3

󰌢 PC Notebook (3.0)
├󰻠 CPU 12th Gen Intel(R) Core(TM) i7-12700H (3) @ 3.50 GHz
├󰍛 GPU Intel Arc A730M @ 1.15 GHz [Discrete]
├󰍛 GPU Intel Iris Xe Graphics @ 0.30 GHz [Integrated]
├󰍛 Vulkan 1.4.311 - Intel open-source Mesa driver [Mesa 25.1.7-arch1.1]
└󰍹 Display(s) 1920x1080 @ 144 Hz in 16" [Built-in]

╭─alu@king in ~
╰─λ cd ./Dev/XBTGPUARC/

╭─alu@king in repo: XBTGPUARC on  master [x!+?]
╰─λ make clean && make
rm -f xbtgpuarc cpu_miner *.o *.d
g++ -std=c17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -c main.cpp -o main.o
g
-std=c17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -c miner_loop.cpp -o miner_loop.o
miner_loop.cpp: In function »void miner_loop(const MiningJob&, const std::function<void(unsigned int, const std::array<unsigned char, 32>&, const MiningJob&)>&, const GpuResources&, int)«:
miner_loop.cpp:51:22: Warnung: Variable »LOCAL_INPUT_SIZE« wird nicht verwendet [-Wunused-variable]
51 | constexpr size_t LOCAL_INPUT_SIZE = 140;
| ^~~~~~~~~~~~~~~~
miner_loop.cpp:52:22: Warnung: Variable »LOCAL_HASH_SIZE« wird nicht verwendet [-Wunused-variable]
52 | constexpr size_t LOCAL_HASH_SIZE = 32;
| ^~~~~~~~~~~~~~~
g
-std=c17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -c opencl_utils.cpp -o opencl_utils.o
g
-std=c17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -c stratum_notify_listener.cpp -o stratum_notify_listener.o
[[A[[Ag
-std=c17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -c globals.cpp -o globals.o
g
-std=c++17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP -o xbtgpuarc main.o miner_loop.o opencl_utils.o stratum_notify_listener.o globals.o -lOpenCL -lboost_system -lboost_json -lpthread

╭─alu@king in repo: XBTGPUARC on  master [x!+?] took 16s
╰─λ ./run.sh
🚀 Starte XBTGPUARC mit Algo: zhash_144_5
👤 Worker: Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8.A730m
🎛 Platform: 1 | Device: 0 | Intensity: 64128
🌐 Pool: solo-btg.2miners.com:4040
🌍 Gefundene OpenCL-Plattformen: 3

[Plattform 0]
Name: rusticl
Vendor: Mesa/X.org
Version: OpenCL 3.0
⚠ Keine Geräte gefunden.

[Plattform 1]
Name: Intel(R) OpenCL Graphics
Vendor: Intel(R) Corporation
Version: OpenCL 3.0
[Device 0] Intel(R) Arc(TM) A730M Graphics

[Plattform 2]
Name: Intel(R) OpenCL Graphics
Vendor: Intel(R) Corporation
Version: OpenCL 3.0
[Device 0] Intel(R) Iris(R) Xe Graphics
🧠 Gerät: 11597.6 MiB VRAM | Nutzung: 33.3 MiB
📡 Verbunden mit solo-btg.2miners.com:4040
🌐 Nachricht:
{"id":1,"result":[null,"0000000000041ac0"],"error":null}
🌐 Nachricht:
{"id":2,"result":true,"error":null}
🌐 Nachricht:
{"id":null,"method":"mining.set_target","params":["0003c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3c3"]}
⚠ Merkle-Branch kein Array, wird übersprungen.
🎯 Job ID: 101084223893030
🧱 PrevHash: aca92cdb086c7a38bb2276a27ea503727111c778b7ad29920192cc901a000000
🌐 Nachricht:
{"id":null,"method":"mining.notify","params":["101084223893030","00000020","aca92cdb086c7a38bb2276a27ea503727111c778b7ad29920192cc901a000000","3edbf3bcc28253fac9d1c2f47923d588bae93609f599c821409c49fc19233da7","2fd70d0000000000000000000000000000000000000000000000000000000000","1d269368","7b823f1d",false,"144_5","BgoldPoW"]}
⚠ Warnung: Eingabepuffer hat nicht exakt 140 Bytes (108)
terminate called after throwing an instance of 'std::invalid_argument'
what(): stoul
./run.sh: Zeile 20: 16641 Abgebrochen (Speicherabzug geschrieben) ./xbtgpuarc --platform 1 --device 0 --algo zhash_144_5 --pool solo-btg.2miners.com --port 4040 --wallet Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8 --worker A730m --password x --intensity 64128

╭─alu@king in repo: XBTGPUARC on  master [x!+?] took 1s
[⚡] ×

XBTGPUARC Miner von Sascha Pürner und ChatGPT Open Source

kernels/zhash.cl:
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable

#define INPUT_SIZE 512
#define HASH_SIZE 32
#define NONCE_OFFSET 12

__constant uint IV[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a,
0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
};

__constant uchar sigma[10][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9,10,11,12,13,14,15 },
{14,10, 4, 8, 9,15,13, 6, 1,12, 0, 2,11, 7, 5, 3 },
{11, 8,12, 0, 5, 2,15,13,10,14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1,13,12,11,14, 2, 6, 5,10, 4, 0,15, 8 },
{ 9, 0, 5, 7, 2, 4,10,15,14, 1,11,12, 6, 8, 3,13 },
{ 2,12, 6,10, 0,11, 8, 3, 4,13, 7, 5,15,14, 1, 9 },
{12, 5, 1,15,14,13, 4,10, 0, 7, 6, 3, 9, 2, 8,11 },
{13,11, 7,14,12, 1, 3, 9, 5, 0,15, 4, 8, 6, 2,10 },
{ 6,15,14, 9,11, 3, 0, 8,12, 2,13, 7, 1, 4,10, 5 },
{10, 2, 8, 4, 7, 6, 1, 5,15,11, 9,14, 3,12,13, 0 }
};

inline uint rotr32(uint x, uint n) {
return (x >> n) | (x << (32 - n));
}

inline void blake2s_core(const uchar* input_data, uchar* out_data) {
uint m[16];
for (int i = 0; i < 16; ++i)
m[i] = (__private uint)&input_data[i * 4];

uint v[16];
for (int i = 0; i < 8; ++i) {
    v[i] = IV[i];
    v[i + 8] = IV[i];
}

v[12] ^= 32;

for (int r = 0; r < 10; ++r) {
    const __constant uchar* s = sigma[r];
    #define G(a,b,c,d,x,y) \
    a += b + x; d = rotr32(d ^ a, 16); \
    c += d;     b = rotr32(b ^ c, 12); \
    a += b + y; d = rotr32(d ^ a, 8);  \
    c += d;     b = rotr32(b ^ c, 7);

    G(v[0],v[4],v[8],v[12], m[s[0]], m[s[1]]);
    G(v[1],v[5],v[9],v[13], m[s[2]], m[s[3]]);
    G(v[2],v[6],v[10],v[14], m[s[4]], m[s[5]]);
    G(v[3],v[7],v[11],v[15], m[s[6]], m[s[7]]);
    G(v[0],v[5],v[10],v[15], m[s[8]], m[s[9]]);
    G(v[1],v[6],v[11],v[12], m[s[10]], m[s[11]]);
    G(v[2],v[7],v[8],v[13], m[s[12]], m[s[13]]);
    G(v[3],v[4],v[9],v[14], m[s[14]], m[s[15]]);
    #undef G
}

for (int i = 0; i < 8; ++i) {
    *(__private uint*)&out_data[i * 4] = v[i] ^ v[i + 8];
}

}

__kernel void zhash_144_5(
__global uchar* input_data,
__global uchar* output_hashes,
__global uint* output_nonces,
const uint start_nonce
) {
int gid = get_global_id(0);
uint nonce = start_nonce + gid;

uchar input[INPUT_SIZE];

// lade input_data + nonce
for (int i = 0; i < INPUT_SIZE; ++i)
    input[i] = input_data[gid * INPUT_SIZE + i];

input[NONCE_OFFSET + 0] = (uchar)((nonce >>  0) & 0xFF);
input[NONCE_OFFSET + 1] = (uchar)((nonce >>  8) & 0xFF);
input[NONCE_OFFSET + 2] = (uchar)((nonce >> 16) & 0xFF);
input[NONCE_OFFSET + 3] = (uchar)((nonce >> 24) & 0xFF);

uchar hash[HASH_SIZE];
blake2s_core(input, hash);

// Ergebnis schreiben
for (int i = 0; i < HASH_SIZE; ++i)
    output_hashes[gid * HASH_SIZE + i] = hash[i];

output_nonces[gid] = nonce;

}

cpu_miner.cpp:
#include
#include
#include
#include
#include
#include
#include
#include
#include

#include "globals.hpp"

constexpr int INPUT_SIZE = 512;
constexpr int HASH_SIZE = 32;
constexpr int NONCE_OFFSET = 12;
constexpr int BUCKET_COUNT = 32;

constexpr uint32_t IV[8] = {
0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a,
0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19
};

constexpr uint8_t sigma[10][16] = {
{ 0, 1, 2, 3, 4, 5, 6, 7, 8, 9,10,11,12,13,14,15 },
{14,10, 4, 8, 9,15,13, 6, 1,12, 0, 2,11, 7, 5, 3 },
{11, 8,12, 0, 5, 2,15,13,10,14, 3, 6, 7, 1, 9, 4 },
{ 7, 9, 3, 1,13,12,11,14, 2, 6, 5,10, 4, 0,15, 8 },
{ 9, 0, 5, 7, 2, 4,10,15,14, 1,11,12, 6, 8, 3,13 },
{ 2,12, 6,10, 0,11, 8, 3, 4,13, 7, 5,15,14, 1, 9 },
{12, 5, 1,15,14,13, 4,10, 0, 7, 6, 3, 9, 2, 8,11 },
{13,11, 7,14,12, 1, 3, 9, 5, 0,15, 4, 8, 6, 2,10 },
{ 6,15,14, 9,11, 3, 0, 8,12, 2,13, 7, 1, 4,10, 5 },
{10, 2, 8, 4, 7, 6, 1, 5,15,11, 9,14, 3,12,13, 0 }
};

inline uint32_t rotr32(uint32_t x, uint32_t n) {
return (x >> n) | (x << (32 - n));}

void blake2s_cpu(const uint8_t* input, size_t len, uint8_t* out) {
uint32_t m[16] = {0}; {
for (int i = 0; i < 16 && (size_t)(i * 4 + 3) < len; ++i)
{
m[i] = input[i4 + 0] |
(input[i
4 + 1] << 8) |
(input[i4 + 2] << 16) |
(input[i
4 + 3] << 24);
}

uint32_t v[16];
for (int i = 0; i < 8; ++i) {
    v[i] = IV[i];
    v[i + 8] = IV[i];
}

v[12] ^= len;

for (int r = 0; r < 10; ++r) {
    const uint8_t* s = sigma[r];
    #define G(a,b,c,d,x,y) \
    a += b + x; d = rotr32(d ^ a, 16); \
    c += d;     b = rotr32(b ^ c, 12); \
    a += b + y; d = rotr32(d ^ a, 8);  \
    c += d;     b = rotr32(b ^ c, 7);

    G(v[0],v[4],v[8],v[12], m[s[0]], m[s[1]]);
    G(v[1],v[5],v[9],v[13], m[s[2]], m[s[3]]);
    G(v[2],v[6],v[10],v[14], m[s[4]], m[s[5]]);
    G(v[3],v[7],v[11],v[15], m[s[6]], m[s[7]]);
    G(v[0],v[5],v[10],v[15], m[s[8]], m[s[9]]);
    G(v[1],v[6],v[11],v[12], m[s[10]], m[s[11]]);
    G(v[2],v[7],v[8],v[13], m[s[12]], m[s[13]]);
    G(v[3],v[4],v[9],v[14], m[s[14]], m[s[15]]);
    #undef G
}

for (int i = 0; i < 8; ++i) {
    uint32_t h = v[i] ^ v[i + 8];
    out[i * 4 + 0] = h & 0xFF;
    out[i * 4 + 1] = (h >> 8) & 0xFF;
    out[i * 4 + 2] = (h >> 16) & 0xFF;
    out[i * 4 + 3] = (h >> 24) & 0xFF;
}
}

}

struct MinerConfig {
uint32_t start_nonce;
uint32_t nonce_range;
int thread_count;
uint8_t* input_buffer;
uint8_t* output_hashes;
};

struct BucketResult {
std::atomic<uint32_t> min_nonce;
};

void submit_share(const std::string& nonce_hex, const std::string& ntime_hex); // Vorwärtsdeklaration

void cpu_miner_thread(MinerConfig cfg, BucketResult* buckets, int tid) {

uint32_t nonce_start = cfg.start_nonce + tid * (cfg.nonce_range / cfg.thread_count);
uint32_t nonce_end   = (tid == cfg.thread_count - 1)
? cfg.start_nonce + cfg.nonce_range
: nonce_start + (cfg.nonce_range / cfg.thread_count);

for (uint32_t nonce = nonce_start; nonce < nonce_end; ++nonce) {
    uint8_t* input     = &cfg.input_buffer[tid * INPUT_SIZE];
    uint8_t* hash_out  = &cfg.output_hashes[tid * HASH_SIZE];

    std::memset(input, 0, INPUT_SIZE);
    input[NONCE_OFFSET + 0] = (nonce >>  0) & 0xFF;
    input[NONCE_OFFSET + 1] = (nonce >>  8) & 0xFF;
    input[NONCE_OFFSET + 2] = (nonce >> 16) & 0xFF;
    input[NONCE_OFFSET + 3] = (nonce >> 24) & 0xFF;

    blake2s_cpu(input, INPUT_SIZE, hash_out);
    uint32_t bucket = hash_out[0] % BUCKET_COUNT;

    // Hash ≤ Target?
    bool is_valid = true;
    for (int i = 31; i >= 0; --i) {
        if (hash_out[i] < current_target[i]) break;
        if (hash_out[i] > current_target[i]) { is_valid = false; break; }
    }

    if (is_valid) {
        std::stringstream ss_nonce;
        ss_nonce << std::hex << std::setw(8) << std::setfill('0') << nonce;
        submit_share(ss_nonce.str(), "00000000"); // FIXME
    }

    uint32_t current = buckets[bucket].min_nonce.load();
    while (nonce < current &&
        !buckets[bucket].min_nonce.compare_exchange_weak(current, nonce)) {}
}

}

void run_cpu_miner(uint32_t start_nonce, uint32_t nonce_range, int threads) {
std::vector<std::thread> workers(threads);
std::vector buckets(BUCKET_COUNT);
std::vector<uint8_t> input_buffers(threads * INPUT_SIZE);
std::vector<uint8_t> output_hashes(threads * HASH_SIZE);

for (auto& b : buckets)
    b.min_nonce.store(UINT32_MAX);

MinerConfig cfg {
    .start_nonce    = start_nonce,
    .nonce_range    = nonce_range,
    .thread_count   = threads,
    .input_buffer   = input_buffers.data(),
    .output_hashes  = output_hashes.data()
};

for (int i = 0; i < threads; ++i)
    workers[i] = std::thread(cpu_miner_thread, cfg, buckets.data(), i);

for (auto& t : workers)
    t.join();

for (int i = 0; i < BUCKET_COUNT; ++i) {
    uint32_t n = buckets[i].min_nonce.load();
    if (n != UINT32_MAX)
        std::cout << "Bucket[" << i << "] Min Nonce: 0x" << std::hex << n << std::dec << '\n';
}

}

int main(int argc, char** argv) {
uint32_t start_nonce = 0x00000000;
uint32_t nonce_range = 0x00100000;
int thread_count = std::thread::hardware_concurrency();

if (argc > 1) start_nonce = std::stoul(argv[1], nullptr, 16);
if (argc > 2) nonce_range = std::stoul(argv[2], nullptr, 16);
if (argc > 3) thread_count = std::stoi(argv[3]);

std::cout << "CPU Miner StartNonce=0x" << std::hex << start_nonce
<< " Range=0x" << nonce_range
<< " Threads=" << std::dec << thread_count << '\n';

run_cpu_miner(start_nonce, nonce_range, thread_count);
return 0;

}

globals.cpp:

#include "globals.hpp"

int next_request_id = 1;
std::string current_job_id = "";
std::string worker_name = "";

std::array<uint8_t, 32> current_target = {};

globals.hpp:

#pragma once
#include
#include <CL/cl.h>
#include
#include

#define INPUT_SIZE 512
#define HASH_SIZE 32
#define NONCES_PER_THREAD 1
#define BUCKET_COUNT 32
#define HASH_ROUNDS_OUTPUT_SIZE 32

struct GpuResources {
cl_context context = nullptr;
cl_command_queue queue = nullptr;
cl_program program = nullptr;
cl_kernel kernel = nullptr;
cl_device_id device = nullptr;
cl_mem input_buffer = nullptr;
cl_mem output_buffer = nullptr;
cl_mem output_hashes_buffer = nullptr;
cl_mem pool_target_buffer = nullptr;
};

extern int next_request_id;
extern std::string current_job_id;
extern std::string worker_name;

extern std::array<uint8_t, 32> current_target;

main.cpp:

#include "stratum_notify_listener.hpp"
#include "globals.hpp"
#include "notify_parser.hpp"
#include "miner_loop.hpp"
#include "mining_job.hpp"
#include "opencl_utils.hpp"

#include <CL/cl.h>
#include
#include
#include
#include
#include
#include

// OpenCL-Geräte
void list_opencl_devices() {
cl_uint num_platforms = 0;
cl_int err = clGetPlatformIDs(0, nullptr, &num_platforms);
if (err != CL_SUCCESS) {
std::cerr << "❌ Fehler bei clGetPlatformIDs: " << err << "\n";
return;
}

std::vector<cl_platform_id> platforms(num_platforms);
clGetPlatformIDs(num_platforms, platforms.data(), nullptr);

std::cout << "🌍 Gefundene OpenCL-Plattformen: " << num_platforms << "\n";

for (cl_uint i = 0; i < num_platforms; ++i) {
    char name[128], vendor[128], version[128];
    clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(name), name, nullptr);
    clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, nullptr);
    clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(version), version, nullptr);

    std::cout << "\n[Plattform " << i << "]\n";
    std::cout << "  Name:    " << name << "\n";
    std::cout << "  Vendor:  " << vendor << "\n";
    std::cout << "  Version: " << version << "\n";

    cl_uint num_devices = 0;
    err = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
    if (err != CL_SUCCESS || num_devices == 0) {
        std::cout << "  ⚠️  Keine Geräte gefunden.\n";
        continue;
    }

    std::vector<cl_device_id> devices(num_devices);
    clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, num_devices, devices.data(), nullptr);

    for (cl_uint j = 0; j < num_devices; ++j) {
        char devname[128];
        clGetDeviceInfo(devices[j], CL_DEVICE_NAME, sizeof(devname), devname, nullptr);
        std::cout << "    [Device " << j << "] " << devname << "\n";
    }
}

}

int main(int argc, char** argv) {
// Default-Werte
int platform_index = 0;
int device_index = 0;
int intensity = 256;
std::string algo = "zhash_144_5";
std::string wallet = "Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8";
std::string worker = "A730m";
std::string password = "x";
std::string pool_host = "solo-btg.2miners.com";
int pool_port = 4040;

// 🧾 Argumente parsen
for (int i = 1; i < argc; ++i) {
    std::string arg = argv[i];
    if (arg == "--platform" && i + 1 < argc) platform_index = std::atoi(argv[++i]);
    else if (arg == "--device" && i + 1 < argc) device_index = std::atoi(argv[++i]);
    else if (arg == "--intensity" && i + 1 < argc) intensity = std::atoi(argv[++i]);
    else if (arg == "--algo" && i + 1 < argc) algo = argv[++i];
    else if (arg == "--wallet" && i + 1 < argc) wallet = argv[++i];
    else if (arg == "--worker" && i + 1 < argc) worker = argv[++i];
    else if (arg == "--password" && i + 1 < argc) password = argv[++i];
    else if (arg == "--pool" && i + 1 < argc) pool_host = argv[++i];
    else if (arg == "--port" && i + 1 < argc) pool_port = std::atoi(argv[++i]);
    else if (arg == "--help") {
        std::cout << "Usage: ./xbtgpuarc [options]\n"
        << "Options:\n"
        << "  --platform N         OpenCL Plattform-Index (default 0)\n"
        << "  --device N           OpenCL Geräte-Index (default 0)\n"
        << "  --intensity N        Threads pro Gerät (default 256)\n"
        << "  --algo NAME          Kernel/Algo-Name (default zhash_144_5)\n"
        << "  --wallet ADDR        Wallet-Adresse\n"
        << "  --worker NAME        Worker-Name\n"
        << "  --password PASS      Passwort für Pool (default 'x')\n"
        << "  --pool HOST          Pool-Adresse (default 2miners)\n"
        << "  --port PORT          Port (default 4040)\n";
        return 0;
    }
}

std::cout << "🚀 Starte XBTGPUARC mit Algo: " << algo << "\n";
std::cout << "👤 Worker: " << wallet << "." << worker << "\n";
std::cout << "🎛️  Platform: " << platform_index << " | Device: " << device_index << " | Intensity: " << intensity << "\n";
std::cout << "🌐 Pool: " << pool_host << ":" << pool_port << "\n";

list_opencl_devices();

GpuResources resources;
init_opencl("kernels/zhash.cl", algo, platform_index, device_index, intensity, resources);

run_stratum_listener(pool_host, pool_port, wallet, worker, password, intensity, resources);

cleanup_opencl(resources);
return 0;

}

Makefile:

CXX := g++
CXXFLAGS := -std=c++17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP
LDFLAGS := -lOpenCL -lboost_system -lboost_json -lpthread

Quellcode-Dateien

SRC := main.cpp
miner_loop.cpp
opencl_utils.cpp
stratum_notify_listener.cpp
globals.cpp

OBJ := $(SRC:.cpp=.o)
DEPS := $(OBJ:.o=.d)
OUT := xbtgpuarc

Optional: CPU-Miner

CPU_SRC := cpu_miner.cpp globals.cpp
CPU_OBJ := $(CPU_SRC:.cpp=.o)
CPU_OUT := cpu_miner

Standard-Target

all: $(OUT)

Build des GPU-Miners

$(OUT): $(OBJ)
$(CXX) $(CXXFLAGS) -o $@ $^ $(LDFLAGS)

Optional: CPU-Miner separat bauen

$(CPU_OUT): $(CPU_OBJ)
$(CXX) $(CXXFLAGS) -o $@ $^ $(LDFLAGS)

Generisches Compile-Ziel

%.o: %.cpp
$(CXX) $(CXXFLAGS) -c $< -o $@

Clean

clean:
rm -f $(OUT) $(CPU_OUT) *.o *.d

-include $(DEPS)

miner_loop.cpp:

#include "miner_loop.hpp"
#include "mining_job.hpp"
#include "opencl_utils.hpp"

#include
#include
#include
#include
#include
#include
#include
#include
#include

extern std::atomic abort_mining;
extern std::atomic socket_valid;

// 🧱 Erstellt den Eingabepuffer aus dem MiningJob
void build_input_from_job(const MiningJob& job, std::vector<cl_uchar>& input_buffer) {
input_buffer.clear();

// Alle Felder als Hex interpretieren → zu Bytes umwandeln
auto append_hex = [&](const std::string& hex) {
    for (size_t i = 0; i + 1 < hex.size(); i += 2) {
        std::string byte_str = hex.substr(i, 2);
        input_buffer.push_back(static_cast<cl_uchar>(std::stoul(byte_str, nullptr, 16)));
    }
};

append_hex(job.version);
append_hex(job.prevhash);
append_hex(job.ntime);
append_hex(job.coinb1);
append_hex(job.extranonce1);
append_hex(job.extranonce2);
append_hex(job.coinb2);

for (const auto& hash : job.merkle_branch)
    append_hex(hash);

}

// 🚀 GPU-Miningloop
void miner_loop(const MiningJob& job,
const std::function<void(uint32_t, const std::array<uint8_t, 32>&, const MiningJob&)>& on_valid_share,
const GpuResources& resources,
int intensity)
{
std::mt19937 rng(std::random_device{}());
std::uniform_int_distribution<uint32_t> dist;

constexpr size_t LOCAL_INPUT_SIZE = 140;
constexpr size_t LOCAL_HASH_SIZE  = 32;
const size_t batch_size = intensity;

std::vector<cl_uchar> hashes(HASH_SIZE * batch_size);
std::vector<cl_uint>  nonces(batch_size);

std::vector<cl_uchar> host_input_buffer;
build_input_from_job(job, host_input_buffer);

if (host_input_buffer.size() != INPUT_SIZE) {
    std::cerr << "⚠️ Warnung: Eingabepuffer hat nicht exakt 140 Bytes (" << host_input_buffer.size() << ")\n";
}

// Target aus nbits berechnen
uint32_t bits = std::stoul(job.nbits, nullptr, 16);
std::vector<uint8_t> current_target = bits_to_target(bits);

clEnqueueWriteBuffer(resources.queue, resources.input_buffer, CL_TRUE, 0, INPUT_SIZE, host_input_buffer.data(), 0, nullptr, nullptr);

clSetKernelArg(resources.kernel, 0, sizeof(cl_mem), &resources.input_buffer);
clSetKernelArg(resources.kernel, 1, sizeof(cl_mem), &resources.output_hashes_buffer);
clSetKernelArg(resources.kernel, 2, sizeof(cl_mem), &resources.pool_target_buffer);

uint64_t total_hashes = 0;
auto start_time = std::chrono::steady_clock::now();

std::cout << "🔄 Neuer Job empfangen. Starte Hashing-Loop...\n";

while (!abort_mining.load()) {
    if (!socket_valid.load()) break;

    uint32_t start_nonce = dist(rng);
    clSetKernelArg(resources.kernel, 3, sizeof(uint32_t), &start_nonce);

    size_t global_work_size = batch_size;
    clEnqueueNDRangeKernel(resources.queue, resources.kernel, 1, nullptr, &global_work_size, nullptr, 0, nullptr, nullptr);

    clEnqueueReadBuffer(resources.queue, resources.output_hashes_buffer, CL_TRUE, 0, hashes.size(), hashes.data(), 0, nullptr, nullptr);

    for (size_t i = 0; i < batch_size; ++i) {
        std::array<uint8_t, 32> current_hash;
        std::memcpy(current_hash.data(), &hashes[i * HASH_SIZE], HASH_SIZE);

        if (is_valid_hash(current_hash, current_target)) {
            on_valid_share(start_nonce + i, current_hash, job);
        }
    }

    total_hashes += batch_size;
    auto now = std::chrono::steady_clock::now();
    auto elapsed_ms = std::chrono::duration_cast<std::chrono::milliseconds>(now - start_time).count();

    if (elapsed_ms >= 5000) {
        double khs = total_hashes / (double)elapsed_ms;
        std::cout << "⚙️ Leistung: " << std::fixed << std::setprecision(2) << khs << " kH/s\n";
        total_hashes = 0;
        start_time = now;
    }
}

}

std::atomic abort_mining{false};
std::atomic socket_valid{false};

void stop_mining() {
abort_mining.store(true);
}

miner_loop.hpp:

#pragma once

#include "mining_job.hpp"
#include "opencl_utils.hpp"
#include
#include
#include
#include

extern std::atomic abort_mining;

void stop_mining();

void miner_loop(const MiningJob& job,
const std::function<void(uint32_t, const std::array<uint8_t, 32>&, const MiningJob&)>& on_valid_share,
const GpuResources& resources,
int intensity);

mining_job.hpp:

#pragma once

#include
#include
#include
#include

// Enthält einen vollständigen MiningJob basierend auf einem Notify-Event
struct MiningJob {
std::string job_id;
std::string prevhash;
std::string coinb1;
std::string coinb2;
std::vector<std::string> merkle_branch;
std::string version;
std::string nbits;
std::string ntime;
std::string extranonce1;
std::string extranonce2;
std::string bits;
};

// Vergleich: Hash < Target
inline bool is_valid_hash(const std::array<uint8_t, 32>& hash, const std::vector<uint8_t>& target) {
for (size_t i = 0; i < 32; ++i) {
if (hash[i] < target[i]) return true;
if (hash[i] > target[i]) return false;
}
return true;
}

// Wandelt Compact-Format aus "bits" (z. B. 0x1d00ffff) in 256-bit Target
inline std::vector<uint8_t> bits_to_target(uint32_t bits) {
uint32_t exponent = bits >> 24;
uint32_t mantissa = bits & 0x007fffff;

std::vector<uint8_t> target(32, 0);

if (exponent <= 3) {
    // Mantisse nach rechts schieben
    mantissa >>= 8 * (3 - exponent);
    target[31] = mantissa & 0xFF;
    target[30] = (mantissa >> 8) & 0xFF;
    target[29] = (mantissa >> 16) & 0xFF;
} else if (exponent <= 32) {
    // Platzierung der Mantisse ab dem richtigen Byte
    int idx = 32 - exponent;
    target[idx]     = (mantissa >> 16) & 0xFF;
    target[idx + 1] = (mantissa >> 8) & 0xFF;
    target[idx + 2] = mantissa & 0xFF;
} else {
    // Exponent außerhalb gültiger Range
    // → Return leeres Ziel (niemals gültig)
    return std::vector<uint8_t>(32, 0xFF);
}

return target;

}

notify_parser.hpp:

#pragma once

#include
#include
#include
#include
#include <boost/json.hpp>
#include "mining_job.hpp"

inline std::optional parse_notify(const std::string& json_line) {
auto get_string_safe = [](const boost::json::value& v) -> std::string {
return v.is_string() ? v.as_string().c_str() : "";
};

try {
    auto val = boost::json::parse(json_line);
    if (!val.is_object()) return std::nullopt;
    const auto& obj = val.as_object();

    if (obj.if_contains("method") == nullptr ||
        obj.at("method").as_string() != "mining.notify")
        return std::nullopt;

    const auto* param_ptr = obj.if_contains("params");
    if (!param_ptr || !param_ptr->is_array()) return std::nullopt;
    const auto& params = param_ptr->as_array();

    if (params.size() < 8) return std::nullopt;

    MiningJob job;
    job.job_id    = get_string_safe(params[0]);
    job.version   = get_string_safe(params[1]);
    job.prevhash  = get_string_safe(params[2]);
    job.coinb1    = get_string_safe(params[3]);
    job.coinb2    = get_string_safe(params[4]);

    job.merkle_branch.clear();
    if (params[5].is_array()) {
        for (const auto& el : params[5].as_array())
            job.merkle_branch.push_back(get_string_safe(el));
        job.bits  = get_string_safe(params[6]);
        job.ntime = get_string_safe(params[7]);
    } else {
        std::cerr << "⚠️ Merkle-Branch kein Array, wird übersprungen.\n";
        job.bits  = get_string_safe(params[5]);
        job.ntime = get_string_safe(params[6]);
    }

    job.extranonce1 = "";
    job.extranonce2 = "00000000";

    return job;

} catch (const std::exception& e) {
    std::cerr << "❌ parse_notify Fehler: " << e.what() << "\n";
    return std::nullopt;
}

}

opencl_list_devices.cpp:

#include <CL/cl.h>
#include
#include
#include

void check_error(cl_int err, const std::string& msg) {
if (err != CL_SUCCESS) {
std::cerr << "❌ Fehler: " << msg << " (" << err << ")\n";
exit(1);
}
}

int main() {
cl_uint num_platforms = 0;
cl_int err = clGetPlatformIDs(0, nullptr, &num_platforms);
check_error(err, "clGetPlatformIDs (count)");

std::vector<cl_platform_id> platforms(num_platforms);
err = clGetPlatformIDs(num_platforms, platforms.data(), nullptr);
check_error(err, "clGetPlatformIDs (fetch)");

std::cout << "🌍 Gefundene OpenCL-Plattformen: " << num_platforms << "\n";

for (cl_uint i = 0; i < num_platforms; ++i) {
    char name[128], vendor[128], version[128];
    clGetPlatformInfo(platforms[i], CL_PLATFORM_NAME, sizeof(name), name, nullptr);
    clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, nullptr);
    clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(version), version, nullptr);

    std::cout << "\n[Plattform " << i << "]\n";
    std::cout << "  Name:    " << name << "\n";
    std::cout << "  Vendor:  " << vendor << "\n";
    std::cout << "  Version: " << version << "\n";

    cl_uint num_devices = 0;
    err = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
    if (err != CL_SUCCESS || num_devices == 0) {
        std::cout << "  ⚠️  Keine Geräte gefunden.\n";
        continue;
    }

    std::vector<cl_device_id> devices(num_devices);
    clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_ALL, num_devices, devices.data(), nullptr);

    for (cl_uint j = 0; j < num_devices; ++j) {
        char devname[128];
        clGetDeviceInfo(devices[j], CL_DEVICE_NAME, sizeof(devname), devname, nullptr);
        std::cout << "    [Device " << j << "] " << devname << "\n";
    }
}

return 0;

}

opencl_utils.cpp:

#include "opencl_utils.hpp"
#include "globals.hpp"
#include
#include
#include
#include
#include
#include

// 🧱 Konstante Buffergrößen
constexpr size_t TARGET_SIZE_BYTES = 32;

// 🔁 Hilfsfunktion: Hex-String in Bytearray wandeln
void hex_string_to_bytes(const std::string& hex_str, uint8_t* output, size_t output_len) {
if (hex_str.length() < output_len * 2)
throw std::invalid_argument("Hex-String ist zu kurz.");

for (size_t i = 0; i < output_len; ++i) {
    std::string byte_str = hex_str.substr(i * 2, 2);
    output[i] = static_cast<uint8_t>(std::stoul(byte_str, nullptr, 16));
}

}

// 📤 Target (hex) an GPU übergeben
void update_opencl_target(const GpuResources& resources, const std::string& hex_target) {
try {
std::string reversed_hex = hex_target;
std::reverse(reversed_hex.begin(), reversed_hex.end());

    uint8_t target_bytes[TARGET_SIZE_BYTES] = {};
    hex_string_to_bytes(reversed_hex, target_bytes, TARGET_SIZE_BYTES);

    cl_int err = clEnqueueWriteBuffer(
        resources.queue,
        resources.pool_target_buffer,
        CL_TRUE,
        0,
        TARGET_SIZE_BYTES,
        target_bytes,
        0, nullptr, nullptr
    );

    if (err != CL_SUCCESS) {
        std::cerr << "❌ Fehler beim Schreiben in Target-Puffer: " << err << "\n";
    }
} catch (const std::exception& e) {
    std::cerr << "❌ Target-Konvertierungsfehler: " << e.what() << "\n";
}

}

// 🧼 GPU-Resourcen vollständig freigeben
void cleanup_opencl(GpuResources& resources) {
if (resources.input_buffer) clReleaseMemObject(resources.input_buffer);
if (resources.output_hashes_buffer) clReleaseMemObject(resources.output_hashes_buffer);
if (resources.pool_target_buffer) clReleaseMemObject(resources.pool_target_buffer);
if (resources.kernel) clReleaseKernel(resources.kernel);
if (resources.program) clReleaseProgram(resources.program);
if (resources.queue) clReleaseCommandQueue(resources.queue);
if (resources.context) clReleaseContext(resources.context);

resources.input_buffer = nullptr;
resources.output_hashes_buffer = nullptr;
resources.pool_target_buffer = nullptr;
resources.kernel = nullptr;
resources.program = nullptr;
resources.queue = nullptr;
resources.context = nullptr;

}

// 🛠️ Initialisiert alle GPU-Objekte
void init_opencl(const std::string& kernel_path, const std::string& kernel_func_name,
int platform_index, int device_index, int intensity, GpuResources& resources)
{
cl_int err;

// Plattform & Device wählen
cl_uint num_platforms = 0;
clGetPlatformIDs(0, nullptr, &num_platforms);
std::vector<cl_platform_id> platforms(num_platforms);
clGetPlatformIDs(num_platforms, platforms.data(), nullptr);

if ((cl_uint)platform_index >= num_platforms) {
    std::cerr << "❌ Ungültiger Plattform-Index\n";
    std::exit(1);
}

cl_platform_id platform = platforms[platform_index];

cl_uint num_devices = 0;
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
std::vector<cl_device_id> devices(num_devices);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices.data(), nullptr);

if ((cl_uint)device_index >= num_devices) {
    std::cerr << "❌ Ungültiger Geräte-Index\n";
    std::exit(1);
}

resources.device = devices[device_index];

// 🧠 VRAM-Auslastung anzeigen
cl_ulong mem_total = 0;
clGetDeviceInfo(resources.device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &mem_total, nullptr);
double mem_total_mib = mem_total / 1024.0 / 1024.0;
size_t estimated_usage = intensity * (INPUT_SIZE + HASH_SIZE);
double usage_mib = estimated_usage / 1024.0 / 1024.0;

std::cerr << std::fixed << std::setprecision(1);
std::cerr << "🧠 Gerät: " << mem_total_mib << " MiB VRAM | Nutzung: " << usage_mib << " MiB\n";

if (usage_mib > mem_total_mib * 0.7)
    std::cerr << "⚠️  Warnung: >70% VRAM belegt\n";

// 📦 Kontext & Queue erzeugen
resources.context = clCreateContext(nullptr, 1, &resources.device, nullptr, nullptr, &err);
resources.queue = clCreateCommandQueueWithProperties(resources.context, resources.device, 0, &err);

// 📄 Kernel-Code laden
std::ifstream file(kernel_path);
if (!file) {
    std::cerr << "❌ Kernel-Datei nicht gefunden: " << kernel_path << "\n";
    std::exit(1);
}

std::string source((std::istreambuf_iterator<char>(file)), std::istreambuf_iterator<char>());
const char* src = source.c_str();
size_t size = source.size();

// 🔨 Programm bauen
resources.program = clCreateProgramWithSource(resources.context, 1, &src, &size, &err);
err = clBuildProgram(resources.program, 1, &resources.device, nullptr, nullptr, nullptr);

if (err != CL_SUCCESS) {
    size_t log_size = 0;
    clGetProgramBuildInfo(resources.program, resources.device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
    std::vector<char> log(log_size);
    clGetProgramBuildInfo(resources.program, resources.device, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr);
    std::cerr << "❌ Build-Fehler:\n" << log.data() << "\n";
    std::exit(1);
}

// 🧠 Kernel erzeugen
resources.kernel = clCreateKernel(resources.program, kernel_func_name.c_str(), &err);
if (err != CL_SUCCESS) {
    std::cerr << "❌ Kernel konnte nicht erzeugt werden: " << kernel_func_name << "\n";
    std::exit(1);
}

// 📦 GPU-Puffer anlegen
resources.input_buffer = clCreateBuffer(resources.context, CL_MEM_READ_WRITE, INPUT_SIZE * intensity, nullptr, &err);
resources.output_hashes_buffer = clCreateBuffer(resources.context, CL_MEM_READ_WRITE, HASH_SIZE * intensity, nullptr, &err);
resources.pool_target_buffer = clCreateBuffer(resources.context, CL_MEM_READ_ONLY, TARGET_SIZE_BYTES, nullptr, &err);

}

opencl_utils.hpp:

#pragma once

#include
#include <CL/cl.h>
#include "globals.hpp"

void init_opencl(const std::string& kernel_path,
const std::string& kernel_func_name,
int platform_index,
int device_index,
int intensity,
GpuResources& resources);

void cleanup_opencl(GpuResources& resources);

void update_opencl_target(const GpuResources& resources, const std::string& hex_target);

void set_kernel_args(const GpuResources& resources,
cl_mem solution_indexes_buffer,
uint32_t start_nonce);

run.sh:

#!/bin/bash

export GPU_MAX_HEAP_SIZE=2048
export GPU_MAX_USE_SYNC_OBJECTS=1
export GPU_SINGLE_ALLOC_PERCENT=100
export GPU_MAX_ALLOC_PERCENT=100
export GPU_MAX_SINGLE_ALLOC_PERCENT=100
export GPU_ENABLE_LARGE_ALLOCATION=100
export GPU_MAX_WORKGROUP_SIZE=2048

./xbtgpuarc
--platform 1
--device 0
--algo zhash_144_5
--pool solo-btg.2miners.com
--port 4040
--wallet Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8
--worker A730m
--password x
--intensity 256

stratum_notify_listener.cpp:

#include "stratum_notify_listener.hpp"
#include "globals.hpp"
#include "miner_loop.hpp"
#include "notify_parser.hpp"
#include "opencl_utils.hpp"

#include
#include <boost/asio.hpp>
#include <boost/json.hpp>
#include
#include
#include
#include

using boost::asio::ip::tcp;

extern int next_request_id;
extern std::string current_job_id;
extern std::string worker_name;

void submit_share(tcp::socket& socket, const std::string& nonce_hex, const std::string& ntime_hex) {
using namespace boost::json;

array params;
params.emplace_back(worker_name);
params.emplace_back(current_job_id);
params.emplace_back("00000000");
params.emplace_back(ntime_hex);
params.emplace_back(nonce_hex);

object request;
request["id"] = next_request_id++;
request["method"] = "mining.submit";
request["params"] = params;

std::string message = serialize(request) + "\n";
boost::asio::write(socket, boost::asio::buffer(message));

std::cout << "📤 Share gesendet:\n" << message;

}

void run_stratum_listener(const std::string& pool_host,
int pool_port,
const std::string& wallet,
const std::string& worker,
const std::string& password,
int intensity,
GpuResources& gpu_resources)
{
const std::string port_str = std::to_string(pool_port);
const std::string full_user = wallet + "." + worker;

try {
    boost::asio::io_context io_context;
    tcp::resolver resolver(io_context);
    auto endpoints = resolver.resolve(pool_host, port_str);
    tcp::socket socket(io_context);
    boost::asio::connect(socket, endpoints);

    std::cout << "📡 Verbunden mit " << pool_host << ":" << port_str << "\n";

    std::string subscribe = R"({"id": 1, "method": "mining.subscribe", "params": []})" "\n";
    std::string authorize = R"({"id": 2, "method": "mining.authorize", "params": [")"
        + full_user + R"(", ")" + password + R"("]})" "\n";

    boost::asio::write(socket, boost::asio::buffer(subscribe));
    boost::asio::write(socket, boost::asio::buffer(authorize));

    std::string buffer;
    static std::thread mining_thread;

    for (;;) {
        char reply[4096];
        boost::system::error_code error;
        size_t len = socket.read_some(boost::asio::buffer(reply), error);
        buffer.append(reply, len);

        size_t pos = 0;
        while ((pos = buffer.find('\n')) != std::string::npos) {
            std::string line = buffer.substr(0, pos);
            buffer.erase(0, pos + 1);

            boost::json::value json_value = boost::json::parse(line);
            if (json_value.is_object()) {
                auto& obj = json_value.as_object();
                if (obj.contains("method") && obj["method"].as_string() == "mining.set_target") {
                    auto& params = obj["params"].as_array();
                    if (!params.empty()) {
                        std::string hex_target = boost::json::value_to<std::string>(params[0]);
                        update_opencl_target(gpu_resources, hex_target); // ← Fix hier
                    }
                }
            }

            auto job = parse_notify(line);
            if (job) {
                std::cout << "🎯 Job ID: " << job->job_id << "\n";
                std::cout << "🧱 PrevHash: " << job->prevhash << "\n";

                if (mining_thread.joinable()) {
                    stop_mining();
                    mining_thread.join();
                }

                auto share_submitter = [&](uint32_t nonce, const std::array<uint8_t, 32>& hash, const MiningJob& job) {
                    std::stringstream ss_nonce, ss_ntime;
                    ss_nonce << std::hex << std::setw(8) << std::setfill('0') << nonce;
                    ss_ntime << std::hex << std::setw(8) << std::setfill('0') << std::stoul(job.ntime, nullptr, 16);
                    submit_share(socket, ss_nonce.str(), ss_ntime.str());
                };

                mining_thread = std::thread([&, job]() {
                    miner_loop(*job, share_submitter, gpu_resources, intensity);
                });
            }

            std::cout << "🌐 Nachricht:\n" << line << "\n";
        }

        if (error == boost::asio::error::eof) break;
        else if (error) throw boost::system::system_error(error);
    }
} catch (const std::exception& e) {
    std::cerr << "❌ Fehler: " << e.what() << "\n";
}

}

stratum_notify_listener.hpp:

#pragma once
#include
#include "opencl_utils.hpp"

// Startet Stratum-Connection, empfängt Jobs, startet Miner
void run_stratum_listener(const std::string& pool_host,
int pool_port,
const std::string& wallet,
const std::string& worker,
const std::string& password,
int intensity,
GpuResources& gpu_resources); // ← NICHT const

XBTGPUARC Programmupdate. | Ecency