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[i4 + 1] << 8) |
(input[i4 + 2] << 16) |
(input[i4 + 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
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
CPU_SRC := cpu_miner.cpp globals.cpp
CPU_OBJ := $(CPU_SRC:.cpp=.o)
CPU_OUT := cpu_miner
all: $(OUT)
$(OUT): $(OBJ)
$(CXX) $(CXXFLAGS) -o $@ $^ $(LDFLAGS)
$(CPU_OUT): $(CPU_OBJ)
$(CXX) $(CXXFLAGS) -o $@ $^ $(LDFLAGS)
%.o: %.cpp
$(CXX) $(CXXFLAGS) -c $< -o $@
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