21.08.25
XBTGPUARC Zhash_144_5 Open Source Miner von Alucian Intel ARC GPU Spezialisiert Blake2b Vector.
Inhalt:
kernels:
zhash.cl (Exclude)
Dateien A-Z:
globals.cpp
globals.hpp
main.cpp
Makefile
miner_loop.cpp
miner_loop.hpp
mining_job.hpp
notify_parser.hpp
opencl_utils_devices.cpp
opencl_utils.cpp
opencl_utils.hpp
runs.sh
stratum_notify_listener.cpp
stratum_notify_listener.hpp
Inhalt Dateien A-Z
globals.cpp
#include "globals.hpp"
#include "miner_loop.hpp"
// Globale Variablen definieren
std::atomic abort_mining{false};
std::atomic socket_valid{false};
int next_request_id = 1;
std::string current_job_id = "";
std::string worker_name = "";
std::array<uint8_t, 32> current_target = {};
// Funktion implementieren
void stop_mining() { abort_mining.store(true); }
globals.hpp
#pragma once
#include <CL/cl.h>
#include
#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 "globals.hpp"
#include "miner_loop.hpp"
#include "mining_job.hpp"
#include "notify_parser.hpp"
#include "opencl_utils.hpp"
#include "stratum_notify_listener.hpp"
#include <CL/cl.h>
#include
#include
#include
#include
#include
#include
// OpenCL-Geräte auflisten (zur Orientierung)
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();
// Initialisiere OpenCL
GpuResources resources;
init_opencl("kernels/zhash.cl", algo, platform_index, device_index, intensity,
resources);
// Starte Stratum-Listener + Mining-Thread
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
CXXFLAGS := -std=c++17 -Wall -O3 -march=native -DCL_TARGET_OPENCL_VERSION=300
LDFLAGS := -lOpenCL -lboost_system -lboost_json -lpthread
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 <CL/cl.h>
#include
#include
#include
#include // für std::isxdigit
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include
#include
extern std::atomic abort_mining;
extern std::atomic socket_valid;
extern std::atomic job_wurde_übernommen;
// Globale OpenCL-Objekte
cl_context context = nullptr;
cl_command_queue queue = nullptr;
cl_kernel kernel = nullptr;
cl_program program = nullptr;
cl_device_id device = nullptr;
// 🧱 Erstellt den Eingabepuffer aus dem MiningJob
void build_input_from_job(const MiningJob &job,
std::vector<cl_uchar> &input_buffer) {
input_buffer.clear();
auto append_hex = [&](const std::string &hex) {
if (hex.empty())
return;
try {
for (size_t i = 0; i + 1 < hex.size(); i += 2) {
std::string byte_str = hex.substr(i, 2);
// Prüfe auf gültige Hex-Ziffern
if (byte_str.find_first_not_of("0123456789abcdefABCDEF") !=
std::string::npos) {
throw std::invalid_argument("Ungültiges Hex-Zeichen");
}
input_buffer.push_back(
static_cast<cl_uchar>(std::stoul(byte_str, nullptr, 16)));
}
} catch (const std::exception &e) {
std::cerr << "❌ Hex-Parsing-Fehler: " << e.what() << "\n";
input_buffer.clear(); // Bei Fehlern Puffer leeren
}
};
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);
}
// Lade OPENCL Zeugs für die Grafikkarte
void init_opencl() {
if (context)
return;
cl_int err;
cl_platform_id platform;
clGetPlatformIDs(1, &platform, nullptr);
clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, nullptr);
cl_uint num_platforms = 0;
clGetPlatformIDs(0, nullptr, &num_platforms);
std::vector<cl_platform_id> platforms(num_platforms);
clGetPlatformIDs(num_platforms, platforms.data(), nullptr);
for (auto p : platforms) {
cl_uint num_devices = 0;
clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, 0, nullptr, &num_devices);
std::vector<cl_device_id> devices(num_devices);
clGetDeviceIDs(p, CL_DEVICE_TYPE_GPU, num_devices, devices.data(), nullptr);
for (auto d : devices) {
char vendor[256];
clGetDeviceInfo(d, CL_DEVICE_VENDOR, sizeof(vendor), vendor, nullptr);
if (strstr(vendor, "Intel(R) Corporation")) { // ARC GPU-Filter
device = d;
platform = p;
break;
}
}
}
context = clCreateContext(nullptr, 1, &device, nullptr, nullptr, &err);
queue = clCreateCommandQueueWithProperties(context, device, 0, &err);
std::ifstream file("kernels/zhash.cl");
if (!file) {
std::cerr << "❌ Konnte Kernel-Datei nicht öffnen.\n";
std::exit(1);
}
std::string source((std::istreambuf_iterator(file)),
std::istreambuf_iterator());
const char *src = source.c_str();
size_t size = source.size();
program = clCreateProgramWithSource(context, 1, &src, &size, &err);
err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
// nach clBuildProgram(program, ...)
char build_log[65536];
size_t log_size = 0;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
sizeof(build_log), build_log, &log_size);
std::cerr << "--- CL BUILD LOG (" << log_size << " bytes) ---\n"
<< build_log << "\n-----------------\n";
if (err != CL_SUCCESS) {
size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, nullptr,
&log_size);
std::vector log(log_size);
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size,
log.data(), nullptr);
std::cerr << "❌ Build-Fehler:\n" << log.data() << "\n";
std::exit(1);
}
if (err != CL_SUCCESS) {
std::cerr << "❌ Build fehlgeschlagen.\n";
}
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
sizeof(build_log), build_log, &log_size);
std::cerr << "--- CL BUILD LOG ---\n" << build_log << "\n";
kernel = clCreateKernel(program, "zhash_144_5", &err);
if (kernel == nullptr) {
std::cerr << "ERROR: Kernel creation failed\n";
}
}
// Wenn/Dann Funktion für gefundenes Päärchen
bool is_valid_candidate_pair(const std::vector<uint8_t> &hashA,
const std::vector<uint8_t> &hashB) {
if (hashA.empty() || hashB.empty())
return false;
return (hashA[0] ^ hashB[0]) == 0;
}
std::string sanitize_hex_string(const std::string &input) {
std::string output;
output.reserve(input.size());
for (char c : input) {
if (std::isxdigit(static_cast(c))) {
output.push_back(c);
}
}
return output;
}
// --- Hilfsfunktion: Sicheres Umwandeln von Hex-String in uint32_t
std::optional<uint32_t> safe_stoul_hex(const std::string &hex_str) {
try {
size_t idx = 0;
uint32_t value = std::stoul(hex_str, &idx, 16);
if (idx != hex_str.size()) {
// Es wurden nicht alle Zeichen gelesen → ungültig
return std::nullopt;
}
return value;
} catch (...) {
return std::nullopt;
}
}
// 🚀 GPU-Miningloop (auf ulong8, work_size immer 64 umbauen)
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
) {
// 1. Initialisierung
std::random_device rd;
std::mt19937 rng(rd());
std::uniform_int_distribution<uint32_t> dist;
// SYCL-Beispiel für ulong8-Vergleiche
sycl::vec<ulong, 8> nonce = ...; // 512-bit Nonce-Berechnung
sycl::vec<ulong, 8> hash = blake2b(nonce); // SIMD-optimierte Hashfunktion
const size_t local_work_size = 64;
const size_t batch_size = intensity * 4096;
size_t global_work_size =
((batch_size + local_work_size - 1) / local_work_size) * local_work_size;
// 2. Target berechnen
auto clean_bits = sanitize_hex_string(job.nbits);
auto maybe_bits = safe_stoul_hex(clean_bits);
if (!maybe_bits)
return;
std::vector<uint8_t> target = bits_to_target(*maybe_bits);
// 3. Puffer vorbereiten
std::vector<cl_uchar> host_input_buffer;
build_input_from_job(job, host_input_buffer);
host_input_buffer.resize(512, 0);
std::vector<cl_uchar> output_buffer(32 * batch_size);
std::vector<cl_uint> index_buffer(2 * batch_size, 0);
// 4. OpenCL-Buffer
cl_int err;
cl_mem cl_input =
clCreateBuffer(resources.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
512, host_input_buffer.data(), &err);
cl_mem cl_output = clCreateBuffer(resources.context, CL_MEM_WRITE_ONLY,
output_buffer.size(), nullptr, &err);
cl_mem cl_indexes =
clCreateBuffer(resources.context, CL_MEM_WRITE_ONLY,
index_buffer.size() * sizeof(cl_uint), nullptr, &err);
cl_mem cl_solution_count = clCreateBuffer(
resources.context, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &err);
if (err != CL_SUCCESS) { /* Fehlerbehandlung */
}
// 5. Hauptschleife
while (!abort_mining.load() && socket_valid.load()) {
uint32_t start_nonce = batch_size;
;
cl_uint zero = 0;
clEnqueueWriteBuffer(resources.queue, cl_solution_count, CL_TRUE, 0,
sizeof(cl_uint), &zero, 0, nullptr, nullptr);
// Vor der Kernel-Ausführung
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);
clSetKernelArg(resources.kernel, 3, sizeof(uint32_t), &intensity);
std::cout << "Starting mining loop with:\n";
std::cout << " Job ID: " << job.job_id << "\n";
std::cout << " PrevHash: " << job.prevhash << "\n";
std::cout << " Target: " << job.nbits << "\n";
std::cout << " Intensity: " << intensity << "\n";
cl_event event;
err = clEnqueueNDRangeKernel(resources.queue, resources.kernel, 1, nullptr,
&global_work_size, &local_work_size, 0,
nullptr, &event);
clWaitForEvents(1, &event);
// Lösungen verarbeiten
clEnqueueReadBuffer(resources.queue, cl_output, CL_TRUE, 0,
output_buffer.size(), output_buffer.data(), 0, nullptr,
nullptr);
clEnqueueReadBuffer(resources.queue, cl_indexes, CL_TRUE, 0,
index_buffer.size() * sizeof(cl_uint),
index_buffer.data(), 0, nullptr, nullptr);
for (size_t i = 0; i < batch_size; ++i) {
uint32_t idxA = index_buffer[i * 2];
uint32_t idxB = index_buffer[i * 2 + 1];
if (idxA == 0 && idxB == 0)
continue;
if (idxA >= batch_size || idxB >= batch_size)
continue;
std::array<uint8_t, 32> final_hash;
for (int j = 0; j < 32; ++j) {
final_hash[j] =
output_buffer[idxA * 32 + j] ^ output_buffer[idxB * 32 + j];
}
if (is_valid_hash(final_hash, target)) {
on_valid_share(start_nonce + idxA, final_hash, job);
}
}
}
// Aufräumen
clReleaseMemObject(cl_input);
clReleaseMemObject(cl_output);
clReleaseMemObject(cl_indexes);
clReleaseMemObject(cl_solution_count);
clReleaseKernel(resources.kernel);
clReleaseProgram(resources.program);
clReleaseCommandQueue(resources.queue);
clReleaseContext(resources.context);
}
miner_loop.hpp
#pragma once
#include "mining_job.hpp"
#include "opencl_utils.hpp"
#include "optional"
#include
#include
#include
#include
// Externe Globals
extern std::atomic abort_mining;
extern std::atomic socket_valid;
// Mining-Steuerung
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);
// Hex-Helper
std::string sanitize_hex_string(const std::string &input);
std::optional<uint32_t> safe_stoul_hex(const std::string &hex_str);
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;
bool clean_job; // <-- DIESES FELD WURDE HINZUGEFÜGT
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 "mining_job.hpp"
#include <boost/json.hpp>
#include
#include
#include
#include
// Diese Funktion parst eine JSON-Zeile vom Pool und wandelt sie in ein
// MiningJob-Objekt um.
inline std::optional parse_notify(const std::string &line) {
using namespace boost::json;
boost::system::error_code ec;
value json_value = parse(line, ec);
if (ec || !json_value.is_object()) {
return std::nullopt; // Keine gültige JSON-Nachricht
}
const object &obj = json_value.as_object();
// Sicherstellen, dass es eine "mining.notify"-Nachricht ist
if (!obj.contains("method") ||
value_to<std::string>(obj.at("method")) != "mining.notify") {
return std::nullopt;
}
if (!obj.contains("params") || !obj.at("params").is_array()) {
std::cerr << "❌ Fehler: 'mining.notify' hat keine gültigen Parameter.\n";
return std::nullopt;
}
const array ¶ms = obj.at("params").as_array();
// Die Parameter-Anzahl für BTG auf 2miners ist typischerweise 8 oder mehr
if (params.size() < 8) {
std::cerr << "❌ Fehler: 'mining.notify' hat zu wenige Parameter ("
<< params.size() << "). Erwartet >= 8.\n";
return std::nullopt;
}
MiningJob job;
// ==================================================================
// KORREKTE PARAMETER-ZUWEISUNG BASIEREND AUF DEINEM LOG
// ==================================================================
// Stratum-Protokoll für solo-btg.2miners.com:
// params[0]: job_id
// params[1]: version
// params[2]: prevhash
// params[3]: coinb1
// params[4]: coinb2
// params[5]: nbits
// params[6]: ntime
// params[7]: clean_job (boolean)
//
// WICHTIG: Dieser Pool sendet KEINEN separaten 'merkle_branch'.
// Der Merkle-Root muss vom Miner selbst berechnet werden, indem
// die Coinbase-Transaktion gehasht wird. Die 'merkle_branch'
// Liste bleibt also absichtlich leer.
job.job_id = value_to<std::string>(params.at(0));
job.version = value_to<std::string>(params.at(1));
job.prevhash = value_to<std::string>(params.at(2));
job.coinb1 = value_to<std::string>(params.at(3));
job.coinb2 = value_to<std::string>(params.at(4));
job.nbits = value_to<std::string>(params.at(5));
job.ntime = value_to<std::string>(params.at(6));
job.clean_job = params.at(7).as_bool();
// Die Merkle Branch Liste wird explizit geleert, da sie nicht vom Pool kommt.
job.merkle_branch.clear();
std::cout << "🌿 Job korrekt geparst. Merkle Branch ist leer, wie vom Pool "
"erwartet.\n";
// Alte Felder für Kompatibilität füllen
job.bits = job.nbits;
job.extranonce1 = ""; // Wird später vom Subscriber gesetzt
job.extranonce2 = "00000000"; // Platzhalter
// Debug-Ausgabe
std::cout << "🔍 Debug Notify: bits = '" << job.nbits << "', ntime = '"
<< job.ntime << "'\n";
return job;
}
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
#include
// 🎯 Umwandlung von "bits" (kompakte Schwierigkeit) in Ziel-Hash
inline std::vector<uint8_t> bits_to_target(uint32_t bits) {
std::vector<uint8_t> target(32, 0);
uint8_t exponent = bits >> 24;
uint32_t mantissa = bits & 0x007fffff;
std::cout << std::hex << std::setfill('0') << "🧮 bits_to_target: bits = 0x"
<< std::setw(8) << bits << " (exp=" << std::dec << +exponent
<< ", mant=0x" << std::hex << mantissa << ")\n";
if (exponent <= 3) {
mantissa >>= 8 * (3 - exponent);
target[31] = (mantissa >> 16) & 0xff;
target[30] = (mantissa >> 8) & 0xff;
target[29] = mantissa & 0xff;
} else {
int idx = 32 - exponent;
if (idx >= 0 && idx + 3 <= 32) {
target[idx] = (mantissa >> 16) & 0xff;
target[idx + 1] = (mantissa >> 8) & 0xff;
target[idx + 2] = mantissa & 0xff;
} else {
std::cerr << "❌ bits_to_target: Exponent " << +exponent
<< " liegt außerhalb gültiger Range\n";
}
}
return target;
}
// 🧱 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(file)),
std::istreambuf_iterator());
// const char* src = source.c_str();
size_t size = source.size();
// 🔨 Programm bauen
const char *src = "-cl-std=CL2.0 -cl-fast-relaxed-math";
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 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);
}
// 🔍 Max Work-Group-Size prüfen
size_t device_max_wg_size = 0;
clGetDeviceInfo(resources.device, CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(device_max_wg_size), &device_max_wg_size, nullptr);
std::cout << "Device max work group size: " << device_max_wg_size << "\n";
size_t kernel_max_wg_size = 0;
clGetKernelWorkGroupInfo(
resources.kernel, resources.device, CL_KERNEL_WORK_GROUP_SIZE,
sizeof(kernel_max_wg_size), &kernel_max_wg_size, nullptr);
std::cout << "Kernel max work group size: " << kernel_max_wg_size << "\n";
// 📏 Effektives Limit bestimmen
size_t effective_max_wg = std::min(device_max_wg_size, kernel_max_wg_size);
if ((size_t)intensity > effective_max_wg) {
std::cerr << "⚠️ Intensity reduziert von " << intensity << " auf "
<< effective_max_wg << " (Work-group limit)\n";
intensity = static_cast(effective_max_wg);
}
// 📦 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=100
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=262144
./xbtgpuarc
--platform 1
--device 0
--algo zhash_144_5
--pool solo-btg.2miners.com
--port 4040
--wallet Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8
--worker A730m
--password x
--intensity 262144
stratum_notify_listener.cpp
#include "stratum_notify_listener.hpp"
#include "globals.hpp"
#include "miner_loop.hpp"
#include "notify_parser.hpp" // Wir nutzen jetzt den externen Parser
#include "opencl_utils.hpp"
#include <boost/asio.hpp>
#include <boost/json.hpp>
#include
#include
#include
#include
#include
using boost::asio::ip::tcp;
// Globale Variablen, die von anderen Teilen des Programms gesetzt werden
extern int next_request_id;
extern std::string current_job_id;
extern std::string worker_name;
// Funktion zum Senden eines gefundenen Shares an den Pool
void submit_share(tcp::socket &socket, const std::string &nonce_hex,
const std::string &ntime_hex, const std::string &job_id) {
using namespace boost::json;
// Erstellt die Parameter für die "mining.submit" Methode
array params;
params.emplace_back(worker_name);
params.emplace_back(job_id);
params.emplace_back(
"00000000"); // extranonce2 Platzhalter, oft nicht benötigt
params.emplace_back(ntime_hex);
params.emplace_back(nonce_hex);
// Baut die komplette JSON-RPC-Anfrage zusammen
object request;
request["id"] = next_request_id++;
request["method"] = "mining.submit";
request["params"] = params;
// Sendet die Nachricht an den Pool
std::string message = serialize(request) + "\n";
boost::asio::write(socket, boost::asio::buffer(message));
std::cout << "📤 Share für Job " << job_id << " gesendet:\n" << message;
}
// Hauptfunktion, die die Verbindung zum Stratum-Pool hält und auf Nachrichten
// lauscht
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);
worker_name = wallet + "." + worker; // Setzt den globalen Worker-Namen
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";
// Standard-Nachrichten zur Anmeldung am Pool
std::string subscribe =
R"({"id": 1, "method": "mining.subscribe", "params": []})"
"\n";
std::string authorize =
R"({"id": 2, "method": "mining.authorize", "params": [")" +
worker_name + R"(", ")" + password +
R"("]})"
"\n";
boost::asio::write(socket, boost::asio::buffer(subscribe));
boost::asio::write(socket, boost::asio::buffer(authorize));
std::string buffer; // Puffer für eingehende Daten vom Socket
static std::thread mining_thread;
// Endlosschleife zum Lesen von Nachrichten
for (;;) {
char reply[4096];
boost::system::error_code error;
size_t len = socket.read_some(boost::asio::buffer(reply), error);
if (len == 0 && error)
break; // Verbindung geschlossen oder Fehler
buffer.append(reply, len);
size_t pos = 0;
// Verarbeite jede vollständige Zeile (getrennt durch '\n') im Puffer
while ((pos = buffer.find('\n')) != std::string::npos) {
std::string line = buffer.substr(0, pos);
buffer.erase(0, pos + 1);
std::cout << "🌐 Nachricht:\n" << line << "\n";
// Versuch, die Nachricht als Mining-Job zu parsen
auto job_opt = parse_notify(line);
if (job_opt) { // Wenn ein gültiger Job empfangen wurde
auto &job = *job_opt;
current_job_id = job.job_id; // Update der globalen Job-ID
std::cout << "🎯 Job ID: " << job.job_id << "\n";
std::cout << "🧱 PrevHash: " << job.prevhash << "\n";
// Stoppe den alten Mining-Prozess, falls er noch läuft
if (mining_thread.joinable()) {
stop_mining();
mining_thread.join();
}
// Definiere eine Lambda-Funktion, die aufgerufen wird, wenn eine
// Lösung gefunden wird
auto share_submitter = [&](uint32_t nonce,
const std::array<uint8_t, 32> &hash,
const MiningJob &job) {
std::stringstream ss_nonce;
ss_nonce << std::hex << std::setw(8) << std::setfill('0') << nonce;
// Keine Umwandlung von job.ntime nötig, ist schon hex-String
submit_share(socket, ss_nonce.str(), job.ntime, job.job_id);
};
// Starte den neuen Mining-Prozess in einem eigenen Thread
mining_thread = std::thread([&, job]() {
miner_loop(job, share_submitter, gpu_resources, intensity);
});
}
}
if (error == boost::asio::error::eof)
break;
else if (error)
throw boost::system::system_error(error);
}
if (mining_thread.joinable()) {
stop_mining();
mining_thread.join();
}
} catch (const std::exception &e) {
std::cerr << "❌ Stratum-Fehler: " << e.what() << "\n";
}
}
stratum_notify_listener.hpp
#pragma once
#include "opencl_utils.hpp"
#include
// 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