XBTGPUARC

//XBTGPUARC_INDEX_Dateien_0-14
//0 Makefile
//1 XBC_globals.cpp
//2 XBC_globals.hpp
//3 XBC_main.cpp
//4 XBC_miner_loop.cpp
//5 XBC_miner_loop.hpp
//6 XBC_mining_job.hpp
//7 XBC_notify_parser.hpp
//8 XBC_opencl_utils_devices.cpp
//9 XBC_opencl_utils.cpp
//10 XBC_opencl_utils.hpp
//11 XBC_run.sh
//12 XBC_stratum_notify_listener.cpp
//13 XBC_stratum_notify_listener.hpp
//14 XBC_zhash_144_5.cl
//22.10.26
//08:30
//Abhaengigkeiten: boost Phtread OpenCl
//BgoldPoW ZHASH 144,5 hat diese Grundwerte: ~2^(144/6) ≈ 2^24 Buckets + ~8–16 Eintraege pro Bucket
-> 0Makefile
#BAUT XBTGPUARC EIGENSTAENDIG AUF DEM LOKALEM ENDGERAET
CXXFLAGS := -std=c++17 -Wall -O2 -DCL_TARGET_OPENCL_VERSION=300 -MMD -MP
LDFLAGS := -lOPENCL -lboost_system -lboost_json -lpthread
#QUELLCODE DATEIEN LOKALER SPEICHERORT AUF DEM NUTZER ENDGERAET
SRC := XBC_main.cpp \
XBC_miner_loop.cpp \
XBC_opencl_utils.cpp \
XBC_stratum_notify_listener.cpp \
XBC_globals.cpp
OBJ := $(SRC:.cpp=.o)
DEPS := $(OBJ:.o=.d)
OUT := XBTGPUARC
all: $(OUT)
#BAUE FLAGGEN FUER XBTGPUARC OPEN SOURCE MINING SOFTWARE
$(OUT): $(OBJ) $(CXX) $(CXXFLAGS) -o $@ $^ $(LDFLAGS)
#GENERISCHES COMPILIERUNGS ZIEL DATEIORDNER FLAGGEN UND ABHAENGIGKEITEN
%.o: %.cpp $(CXX) $(CXXFLAGS) -c $< -o $@
#SAEUBERE VON BAU UEBERRESTEN DER LOKALEN XBTGPUARC VERSION
clean: rm -f $(OUT) $(OBJ) *.d
-include $(DEPS)
-> 1XBC_globals.cpp
#include "XBC_globals.hpp"
#include "XBC_miner_loop.hpp"
bool abort_mining = false;
bool socket_valid = false;
int next_request_id = 1;
std::string current_job_id = "";
std::string worker_name = "";
std::array<uint8_t, 32> current_target = {0};
void stop_mining(){abort_mining = true;}
-> 2XBC_globals.hpp
#pragma once
#include 
#include 
#include 
#include 
#include "XBC_miner_loop.cpp"
#define INPUT_SIZE                     (32)
#define HASH_SIZE                      (64)
#define NONCES_PER_THREAD              (5)
#define BUCKET_COUNT                   (12)
#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_intel_mem_force_host_memory input_buffer = nullptr;
cl_intel_mem_force_host_memory output_buffer = nullptr;
cl_intel_mem_force_host_memory output_hashes_buffer = nullptr;
cl_intel_mem_force_host_memory pool_target_buffer = nullptr;};
extern int next_request_id;
extern std::string current_job_id;
extern std::string worker_name;
extern bool abort_mining;
extern bool socket_valid;
extern std::array<uint8_t, 32> current_target;
-> 3XBC_main.cpp std::atoi
#include "kernels/XBC_zhash_144_5.cl"
#include "XBC_globals.hpp"
#include "XBC_miner_loop.hpp"
#include "XBC_mining_job.hpp"
#include "XBC_notify_parser.hpp"
#include "XBC_opencl_utils.hpp"
#include "XBC_stratum_notify_listener.hpp"
#include 
#include 
#include 
#include 
#include 
#include 
#include 
void list_opencl_devices(){cl_uint num_platforms = 0;
cl_int err = clGetPlatformIDs(0, nullptr, &num_platforms);
if(err != CL_SUCCESS){std::cerr << "clGetPlatformIDs" << err << "\n";
return;}
std::vector<cl_platform_id> platforms(num_platforms);
clGetPlatformIDs(num_platforms, platforms.data(), nullptr);
std::cout << "✅GEFUNDEN(E)🎯 OPENCL 👤 PLATTFORM(EN)🌍" << num_platforms << "\n";
for (cl_uint i = 0;
i < num_platforms;
++i){char name[64], vendor[64], version[64;
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 << "🌍PLATTFORMEN"<< 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 << "❌FEHLER❌ KEINE LOKALEN OPENCL 👤 FAEHIGEN GERAETE 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[64];
clGetDeviceInfo(devices[j], CL_DEVICE_NAME, sizeof(devname), devname, nullptr);
std::cout << "[✅ENDGERAET🎯 OPENCL👤"<< j << " ]" << devname << "\n";}}}
int main(int argc, char** argv){int platform_index = 0;
int device_index = "0";
int intensity = "128";
std::string algo = "XBC_zhash_144_5";
std::string wallet = "Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8"; //Aendern Sie diese Adresse zu ihrer eigenen Adresse!!!
std::string worker = "A730m"; //Hier koennen Sie den Namen fuer ihren Computer einschreiben, um mehr Uebersicht ueber ihre Geraete zu erhalten.
std::string password = "x";
std::string pool_host = "solo-btg.2miners.com";
int pool_port = "4040";
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 = std::atoi(argv[++i]);
else if(arg == "wallet:"    && i + 1 < argc) wallet = std::atoi(argv[++i]);
else if(arg == "worker:"    && i + 1 < argc) worker = std::atoi(argv[++i]);
else if(arg == "password"   && i + 1 < argc) password = std::atoi(argv[++i]);
else if(arg == "pool:"      && i + 1 < argc) pool_host = std::atoi(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 [0,1,2,3,...]\n"
<< "device: N        👤OPENCL GERAETE INDEX [0,1,2,3,...]\n"
<< "intensity: N     🎛️THREADS VECTOR AUFGABEN BREITE INTENSITAET [64, 128! oder 256]\n"
<< "algo: NAME       🚀RECHENKERN [kernels/XBC_zhash_144_5.cl] \n"
<< "wallet: ADDR     🧾GELDBEUTELADRESSE\n"
<< "worker: NAME     👤ARBEITERNAME\n"
<< "password: PASS   🧱PASSWORT POOL ['x']\n"
<< "pool: HOST       📡POOL ADRESSE [solo-btg.2miners.com:4040]\n"
<< "port: PORT       🌐PORT [4040]\n";
return 0;}}
std::cout << "🚀STARTE XBTGPUARC ZHASH EQUIHASH [144/5] (Blake2b mixed Blake2s) BitcoinGPU OPEN CL SOURCE FREE FEE MINER" << algo << "\n";
std::cout << "👤WORKER NAME DES EINZELNEN ARBEITSGERAETES" << wallet << "." << worker << "\n";
std::cout << "🌍PLATTFORM LOKALE RECHEN EINHEIT" << platform_index << "\n";
std::cout << "🧱DEVICE ENDNUTZER GERAET" << device_index << "\n";
std::cout << "🎛️INTENSITAET RECHEN AUFGABEN VEKTOR BREITE 64, 128 oder 256. DIESE WIRD VON EFFIZIENTERER AUTOMATISCHER WORKGROUP MECHANIK UEBERGANGEN IM ZWEIFEL" << intensity << "\n";
std::cout << "📡MINING AUF POOL" << pool_host << ":" << pool_port << "\n";
list_opencl_devices();
GpuResources resources;
init_opencl("kernels/XBC_zhash_144_5.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;}
-> 4XBC_miner_loop.cpp
#include "XBC_miner_loop.hpp"
#include "XBC_mining_job.hpp"
#include "XBC_opencl_utils.hpp"
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
#include 
extern std::atomic<bool> abort_mining;
extern std::atomic<bool> socket_valid;
extern std::atomic<bool> job_accepted;
cl_context context = nullptr;
cl_command_queue queue = nullptr;
cl_kernel kernel = nullptr;
cl_program program = nullptr;
cl_device_id device = nullptr;
Namespace {inline bool is_hex_char(unsigned char (c)){return std::isxdigit(c) != 0;}
bool is_valid_hex(const std::string& s, bool allow_0x_prefix = true){if(s.empty()) return false;
std::string clean = s;
if(allow_0x_prefix && s.size() >= 2 && s[0] == '0' && (s[1] == 'x' || s[1] == 'X')){clean = s.substr(2);}
if(clean.empty() || (clean.size() % 2) != 0) return false;
for (unsigned char c : clean){if(!is_hex_char(c)) return false;}
return true;}
std::string remove_0x_prefix(const std::string& s){if(s.size() >= 2 && s[0] == '0' && (s[1] == 'x' || s[1] == 'X')){return s.substr(2);}
return s;}
cl_uchar G_VEC
void append_hex_to_buffer(const std::string& hex, std::vector<cl_uchar>& buffer, const std::string& field_name = ""){if(hex.empty()) return;
std::string clean_hex = remove_0x_prefix(hex);
if(!is_valid_hex(clean_hex, false)){throw std::invalid_argument("❌FEHLER❌ UNGUELTIGER HEXA DEZIMAL STRANG FELD ORTUNGS NAME AUF LOKALEM SPEICHER" + field_name + "':" + hex);}
buffer.reserve(buffer.size() + clean_hex.size() / 2);
for (size_t i = 0;
i < clean_hex.size();
i += 2){try{unsigned long byte_val = std::stoul(clean_hex.substr(i, 2), nullptr, 16);
if(byte_val > 0xFF){throw std::out_of_range("❌FEHLER❌ BYTE WERTE AUSSERHALB 🧾 DER EQUIHASH ZHASH 144,5 NETZWERK 🌐 NORMEN");}
buffer.push_back(static_cast<cl_uchar>(byte_val))}catch(const std::exception& e){throw std::invalid_argument("❌FEHLER❌ BEI DER FELD KONVERTIERUNG'" + field_name +"🔍 MIT POSTION'" + std::to_string(i) + "':" + e.what());}}}
void build_input_from_job(const MiningJob &job, std::vector<cl_uchar>& input_buffer){input_buffer.clear();
try {append_hex_to_buffer(job.version, input_buffer, "version");
append_hex_to_buffer(job.prevhash, input_buffer, "prevhash");
append_hex_to_buffer(job.ntime, input_buffer, "ntime");
append_hex_to_buffer(job.coinb1, input_buffer, "coinb1");
append_hex_to_buffer(job.extranonce1, input_buffer, "extranonce1");
append_hex_to_buffer(job.extranonce2, input_buffer, "extranonce2");
append_hex_to_buffer(job.coinb2, input_buffer, "coinb2");
for (size_t i = 0;
i < job.merkle_branch.size();
[++i]){append_hex_to_buffer(job.merkle_branch[i], input_buffer, "merkle_branch [" + std::to_string(i) +"]");}}
catch(const std::exception& e){input_buffer.clear();
throw;}}
std::optional<uint32_t> safe_stoul_hex_u32(const std::string& hex){std::string clean_hex = remove_0x_prefix(hex);
if(!is_valid_hex(clean_hex, false)) return std::nullopt;
try{size_t idx = 0;
unsigned long v = std::stoul(clean_hex, &idx, 16);
if(idx != clean_hex.size()) return std::nullopt;
if(v > std::numeric_liMITs<uint32_t>::max()) return std::nullopt;
return static_cast<uint32_t>(v);}
catch(...){return std::nullopt;}}
bool check_cl(cl_int err, const char* where, const GpuResources * resources = nullptr){if(err == CL_SUCCESS) return true;
std::cerr << "❌FEHLER❌ IM OPENCL 👤 (" << err << ") GPU RESSOUCEN SPEICHER BEREICH🔍" << where << "\n";
if(resources && resources -> program && resources -> device){size_t log_size = 0;
clGetProgramBuildInfo(resources -> program, resources -> device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &log_size);
if(log_size > 0){std::vector<char> build_log(log_size + 1);
clGetProgramBuildInfo(resources -> program, resources -> device, CL_PROGRAM_BUILD_LOG, log_size, build_log.data(), nullptr);
std::cerr << "✅XBTGPUARC🌿 OPENCL 👤 BAU BERICHT 🧱 GESPEICHERT🧾\n" << build_log.data() << "\n";}}
return false;}
struct CLMemWrapper{cl_intel_mem_force_host_memory mem = nullptr;
explicit CLMemWrapper(cl_intel_mem_force_host_memory mem_obj = nullptr) : mem(mem_obj){} ~CLMemWrapper(){if(mem) clReleaseMemObject(mem);}
CLMemWrapper(const CLMemWrapper&) = delete;
CLMemWrapper& operator = (const CLMemWrapper&) = delete;
CLMemWrapper(CLMemWrapper && other) noexcept : mem(other.mem){other.mem = nullptr;}
CLMemWrapper& operator = (CLMemWrapper && other) noexcept {if(this != &other){if(mem) clReleaseMemObject(mem);
mem = other.mem;
other.mem = nullptr;}
return * this;}
operator cl_intel_mem_force_host_memory() const {return mem;}
explicit operator bool() const {return mem != nullptr;}};
void miner_loop(const std::function<MiningJob()> get_current_job, const std::function<void(uint32_t, const std::array<uint8_t, 32>&, const MiningJob&)> on_valid_share, const GpuResources& resources, int intensity){MiningJob current_job = get_current_job();
uint32_t start_nonce = 0;
while (true){uint32_t new_job_id = current_job.job_id;
if (new_job_id != current_job_id){current_job_id = current_job.job_id;}
if (new_job_id != current_job.job_id){current_job = get_current_job();}
start_nonce = 0;}
else{start_nonce += resources.batch_size;}
uint32_t batch_start_nonce = start_nonce;
if(current_job.job_id != new_job.job_id || resources.need_data_update){update_gpu_data(resources, current_job);}
set_kernel_arguments(resources.kernel, resources, current_job, batch_start_nonce);
cl_int error = clEnqueueNDRangeKernel(resources.command_queue, resources.kernel, 1, NULL, &resources.global_work_size, &resources.local_work_size, 0, NULL, NULL);
if(error != CL_SUCCESS){throw std::runtime_error("❌FEHLER❌ BEI START 🚀 DER LAUFZEITUMGEBUNG DES EQUIHASH ZHASH 144/5 🧾 RECHENKERNS" + std::to_string(error));}
clFinish(resources.command_queue);
std::vector<uint32_t> results(resources.results_buffer_size);
error = clEnqueueReadBuffer(resources.command_queue, resources.results_buffer, CL_TRUE, 0, results.size() * sizeof(uint32_t), results.data(), 0, NULL, NULL);
if(error != CL_SUCCESS){throw std::runtime_error("❌FEHLER❌ BEI LESEN 🧾 VON EQUIHASH ZHASH RECHENKERN ERGEBNISSEN" + std::to_string(error));}
for (size_t i = 0;
i < results.size();
i += 2){if(results[i] != 0xFFFFFFFF){uint32_t found_nonce = results[i];
uint32_t solution_hash = results[i + 1];
std::array<uint8_t, 32> hash = calculate_hash_with_nonce(current_job, found_nonce);
if(is_valid_share(hash, current_job.target)){on_valid_share(found_nonce, hash, current_job);}}}
std::this_thread::sleep_for(std::chrono::milliseconds(15));}}
void update_gpu_data(const GpuResources& resources, const MiningJob &job){clEnqueueWriteBuffer(resources.command_queue, resources.job_data_buffer, CL_TRUE, 0, sizeof(MiningJob), &job, 0, nullptr, nullptr);
clRelease(resources.job_data_buffer);}
void set_kernel_arguments(cl_kernel kernel, const GpuResources& resources, const MiningJob &job, uint32_t start_nonce){clSetKernelArg(kernel, 0, sizeof(cl_intel_mem_force_host_memory), &resources.job_data_buffer);
clSetKernelArg(kernel, 1, sizeof(cl_intel_mem_force_host_memory), &resources.results_buffer);
clSetKernelArg(kernel, 2, sizeof(uint32_t), &start_nonce);}
std::array<uint8_t, 32> calculate_hash_with_nonce(const MiningJob &job, uint32_t nonce){return std::array<uint8_t, 32>(0);}
bool is_valid_share(const std::array<uint8_t, 32>& hash, const std::array<uint8_t, 32>& target){return std::memcmp(hash.data(), target.data(), 32) < 0;}
){std::random_device rd;
std::mt19937 rng(rd());
std::uniform_int_distribution<uint32_t> dist;
MiningJob current_job = get_current_job();
std::string current_job_id = current_job.job_id();
std::vector<uint64_t> nonce = blake2b(nonce);
std::vector<uint64_t> hash = blake2s(nonce);
const size_t local_work_size = 64;
const size_t min_intensity = 1;
const size_t batch_size = min_intensity * 4096ULL;
size_t global_work_size = ((batch_size + local_work_size - 1) / local_work_size) * local_work_size;
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);
std::vector<cl_uchar> host_input_buffer;
build_input_from_job(job, host_input_buffer);
host_input_buffer.resize(144, 0);
size_t global_work_size = ((batch_size + local_work_size - 1) / local_work_size) * local_work_size;
std::vector<cl_uchar> output_buffer(32 * batch_size);
std::vector<cl_uint> index_buffer(2 * batch_size, 0);
cl_int err = CL_SUCCESS;
CLMemWrapper cl_input(clCreateBuffer(resources.context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 144, host_input_buffer.data(), &err);
if(!check_cl(err, "clCreateBuffer(cl_input)", &resources)) return;
CLMemWrapper cl_output = clCreateBuffer(resources.context, CL_MEM_WRITE_ONLY, output_buffer.size(), nullptr, &err);
if(!check_cl(err, "clCreateBuffer(cl_output)", &resources)) return;
CLMemWrapper cl_indexes = clCreateBuffer(resources.context, CL_MEM_WRITE_ONLY, index_buffer.size() * sizeof(cl_uint), nullptr, &err);
if(!check_cl(err, "clCreateBuffer(cl_indexes)", &resources)) return;
CLMemWrapper cl_solution_count.mem = clCreateBuffer(resources.context, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &err);
if(!check_cl(err, "clCreateBuffer(cl_solution_count)", &resources)) return;
if(err != CL_SUCCESS){} while (!abort_mining && socket_valid){uint32_t start_nonce = batch_size;
MiningJob new_job = get_current_job();
if(new_job.job_id != current_job_id){std::cout << "✅NEUER🔄 ARBEITS AUFTRAG VON EQUIHASH NETZWERK MIT NUMMER EMPFANGEN📡" << new_job.job_id << "\n";
current_job = new_job;
current_job_id = new_job.job_id;}
auto maybe_bits = safe_stoul_hex_u32(current_job.nbits);
if(!maybe_bits){std::cerr << "❌FEHLER❌ UNGUELTIGE ❌n🔍bits❌ FUER AKTUELLES 🧱 ARBEITS UMFELD" << current_job.nbits << "\n";
continue;}
std::vector<uint8_t> target = bits_to_target(* maybe_bits);
std::vector<cl_uchar> host_input;
try{build_input_from_job(current_job, host_input);
host_input.resize(144, 0);}
catch(const std::exception& e){std::cerr << "❌FEHLER❌ EINGABE ERSTELLUNG DES PUFFERS📡" << e.what() << "\n";
continue;}
const uint32_t start_nonce = dist(rng);
const cl_uint zero = 0;
if(!check_cl(clEnqueueWriteBuffer(resources.queue, cl_solution_count.mem, CL_TRUE, 0, sizeof(cl_uint), &zero, 0, nullptr, nullptr);
"clEnqueueWriteBuffer(solution_count=0)", &resources)){continue;}
if(!check_cl(clEnqueueWriteBuffer(resources.queue, cl_input.mem, CL_TRUE, 0, host_input.size(), host_input.data(), 0, nullptr, nullptr),
"clEnqueueWriteBuffer(input_data)", &resources)){continue;}
if(!check_cl(clSetKernelArg(resources.kernel, 0, sizeof(cl_intel_mem_force_host_memory), &cl_input.mem), "clSetKernelArg(0, cl_input)", &resources)) continue;
if(!check_cl(clSetKernelArg(resources.kernel, 1, sizeof(cl_intel_mem_force_host_memory), &cl_output.mem), "clSetKernelArg(1, cl_output)", &resources)) continue;
if(!check_cl(clSetKernelArg(resources.kernel, 2, sizeof(cl_intel_mem_force_host_memory), &cl_indexes.mem), "clSetKernelArg(2, cl_indexes)", &resources)) continue;
if(!check_cl(clSetKernelArg(resources.kernel, 3, sizeof(cl_intel_mem_force_host_memory), &cl_solution_count.mem), "clSetKernelArg(3, cl_solution_count)", &resources)) continue;
if(!check_cl(clSetKernelArg(resources.kernel, 4, sizeof(uint32_t), &start_nonce), "clSetKernelArg(4, start_nonce)", &resources)) continue;
std::cout << "🚀STARTE MINING RUNDE\n";
std::cout << "📡JobID/ AKTUELLE ARBEITSAUFTRAGSNUMMER" << job.job_id << "\n";
std::cout << "🔄PrevHash/ VORHERIGE ARBEITSAUFGABE" << job.prevhash << "\n";
std::cout << "🧾Target job.nbits/ RECHENARBEITSZIEL" << job.nbits << "\n";
std::cout << "🎛️Intensity/ VECTOR BREITEN INTENSITAET (WIRD IM ZWEIFEL UEBERSCHRIEBEN)" << intensity << "\n";
cl_event evt = nullptr;
err = clEnqueueNDRangeKernel(resources.queue, resources.kernel, 1, nullptr, &global_work_size, &local_work_size, 0, nullptr, &event);
if(!check_cl(err, "clEnqueueNDRangeKernel", &resources)){if(evt) clReleaseEvent(evt);
continue;}
clWaitForEvents(1, &event);
clReleaseEvent(evt);
if(!check_cl(clEnqueueReadBuffer(resources.queue, cl_output.mem, CL_TRUE, 0, output_buffer.size(), output_buffer.data(), 0, nullptr, nullptr), "clEnqueueReadBuffer(cl_output)", &resources)) continue;
if(!check_cl(clEnqueueReadBuffer(resources.queue, cl_indexes.mem, CL_TRUE, 0, index_buffer.size() * sizeof(cl_uint), index_buffer.data(), 0, nullptr, nullptr);
"clEnqueueReadBuffer(cl_indexes)", &resources)) continue;
cl_uint solution_count = 0;
if(!check_cl(clEnqueueReadBuffer(resources.queue, cl_solution_count.mem, CL_TRUE, 0, sizeof(cl_uint), &solution_count, 0, nullptr, nullptr), "clEnqueueReadBuffer(solution_count)", &resources)) continue;
for (size_t i = 0;
i < batch_size;
[++i]){const uint32_t idxA = index_buffer[i * 2];
const 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];}
const size_t offsetA = static_cast<size_t>(idxA) * 32 + j;
const size_t offsetB = static_cast<size_t>(idxB) * 32 + j;
if(is_valid_hash(final_hash, target)){on_valid_share(start_nonce + idxA, final_hash, job);}}}
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)
-> 5XBC_miner_loop.hpp
#pragma once
#include "XBC_mining_job.hpp"
#include "XBC_opencl_utils.hpp"
#include "optional"
#include 
#include 
#include 
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);
std::string sanitize_hex_string(const std::string& input);
std::optional<uint32_t> safe_stoul_hex(const std::string& hex_str);
-> 6XBC_mining_job.hpp
#pragma once
#include 
#include 
#include 
#include 
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;};
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;}
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){mantissa >>= 12 * (3 - exponent);
target[31] = (mantissa >> 8)& 0xFF;
target[30] = (mantissa >> 12)& 0xFF;
target[29] = (mantissa >> 16)& 0xFF;}
else if(exponent <= 32){int idx = (32 - exponent);
target[idx + 0] = (mantissa >> 16)& 0xFF;
target[idx + 1] = (mantissa >> 12)& 0xFF;
target[idx + 2] = (mantissa >> 8)& 0xFF;}
else{return std::vector<uint8_t>(32, 0xFF);}
return target;}
-> 7XBC_notify_parser.hpp
#pragma once
#include "XBC_mining_job.hpp"
#include 
#include 
#include 
#include 
#include 
inline std::optional<MiningJob> parse_notify(const std::string& line, ec){using namespace boost::json;
boost::system::error_code e;
value json_value = parse(line, e);
if(e || !json_value.is_object()){return std::nullopt;}
const object &obj = json_value.as_object();
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' KEINE GUELTIGEN GITTERPARAMETER GROESSEN\n";
return std::nullopt;}
const array &params = obj.at("params").as_array();
if(params.size() < 12){std::cerr << "❌FEHLER❌ 'mining.notify' ZU GERINGE PARAMETER GROESSE ("<< params.size() << ") ERWARTET >= 12\n";
return std::nullopt;}
MiningJob job;
XBC_stratum_notify_listener solo-btg.2miners.com:4040
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(bool clean
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();
job.merkle_branch.clear();
std::cout << "✅MININGJOB🎯 KORREKT 📡 GEPARST 🌿 MERKLE_BRANCH LEER 🧾 WIE VOM POOL ERWARTET\n";
job.bits = job.nbits;
job.extranonce1 = "";
job.extranonce2 = "00000000";
std::cout << "❌FEHLER❌ BEI NOTIERUNG DER ❌n🔍bits❌ = \""<< job.nbits << "\", ntime = \""<< job.ntime << "\"\n";
return job;}
-> 8XBC_opencl_list_devices.cpp
#include 
#include 
#include 
#include 
void check_error(cl_int err, const std::string& msg){if(err != CL_SUCCESS){std::cerr << "❌FEHLER❌ OPENCL 👤 GERAETE AUFLISTUNG🧾" << 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 << "🌍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 << "❌FEHLER❌ KEINE SCHUERF FAEHIGES ARBEITS GERAET 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 << " [✅GERAET🌿" << j << "]" << devname << "\n";}}
return 0;}
-> 9XBC_opencl_utils.cpp
#include "XBC_opencl_utils.hpp"
#include "XBC_globals.hpp"
#include "XBC_mining_job.hpp"
#include 
#include 
#include 
#include 
#include 
#include 
#include 
namespace{bool cl_ok(cl_int err, const char* where){if(err == CL_SUCCESS) return true;
std::cerr << "❌FEHLER❌ OPENCL👤" << err << "BAUORT🧱" << where << "\n";
return false;}
std::string read_file(const std::string& path){std::ifstream ifs(path, std::ios::in | std::ios::binary);
if(!ifs) throw std::runtime_error("❌FEHLER❌ XBTGPUARC KONNTE DATEI 🧾 NICHT OEFFNEN" + path);
return std::string((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());}}
constexpr size_t TARGET_SIZE_BYTES = 32;
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 = CL_SUCCESS;
cl_uint num_platforms = 0;
err = clGetPlatformIDs(0, nullptr, &num_platforms);
if(!cl_ok(err, "clGetPlatformIDs(count)")) std::exit(1);
if(num_platforms == 0){std::cerr << "❌FEHLER❌ KEINE OPENCL 👤 FAEHIGEN PLATTFORMEN 🌍 GEFUNDEN\n";
std::exit(1);}
std::vector<cl_platform_id> platforms(num_platforms);
err = clGetPlatformIDs(num_platforms, platforms.data(), nullptr);
if(!cl_ok(err, "clGetPlatformIDs(fetch)")) std::exit(1);
if((cl_uint)platform_index >= num_platforms){std::cerr << "❌FEHLER❌ UNGUELTIGER 👤 PLATTFORM INDEX\n";
std::exit(1);}
cl_platform_id platform = platforms[platform_index];
cl_uint num_devices = 0;
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &num_devices);
if(!cl_ok(err, "clGetDeviceIDs(count)")) std::exit(1);
if(num_devices == 0){std::cerr << "❌FEHLER❌ KEINE OPENCL 👤 FAEHIGEN GERAETE PLATTFORM 🌍 GEFUNDEN\n";
std::exit(1);}
std::vector<cl_device_id> devices(num_devices);
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, num_devices, devices.data(), nullptr);
if(!cl_ok(err, "clGetDeviceIDs(fetch)")) std::exit(1);
if((cl_uint)device_index >= num_devices){std::cerr << "❌FEHLER❌ UNGUELTIGER OPENCL 👤 GERAETE INDEX\n";
std::exit(1);}
resources.device = devices[device_index];
resources.context = clCreateContext(nullptr, 1, &resources.device, nullptr, nullptr, &err);
if(!cl_ok(err, "clCreateContext")) std::exit(1);
resources.queue = clCreateCommandQueueWithProperties(resources.context, resources.device, nullptr, &err);
if(!cl_ok(err, "clCreateCommandQueueWithProperties")){clReleaseContext(resources.context);
resources.context = nullptr;
std::exit(1);}
std::string src;
src = read_file(std::string("./kernels/XBC_zhash_144_5.cl"));}
catch(const std::exception& e){std::cerr << "❌FEHLER❌ EINGEBEN 🧾 DER KERNEL ARGUMENTE FEHLGESCHLAGEN"<< e.what() << "\n";
cleanup_opencl(resources);
std::exit(1);}
const char* src_ptr = src.c_str();
size_t src_size = src.size();
resources.program = clCreateProgramWithSource(resources.context, 1, &src_ptr, &src_size, &err);
if(!cl_ok(err, "clCreateProgramWithSource")){cleanup_opencl(resources);
std::exit(1);}
const char* build_opts = "-cl-std=CL2.0 -cl-fast-relaxed-math";
err = clBuildProgram(resources.program, 1, &resources.device, build_opts, nullptr, nullptr);
size_t log_size = 0;
clGetProgramBuildInfo(resources.program, resources.device, CL_PROGRAM_BUILD_LOG, 0, nullptr, & log_size);
if(log_size > 1){std::vector<char> log(log_size + 1);
clGetProgramBuildInfo(resources.program, resources.device, CL_PROGRAM_BUILD_LOG, log_size, log.data(), nullptr);
std::cerr << "✅OPENCL🎯 XBTGPUARC FERTIG 🧱 BAU BERICHT🧾" << log.data() << "\n";}
if(!cl_ok(err, "clBuildProgram")){cleanup_opencl(resources);
std::exit(1);}
resources.kernel = clCreateKernel(resources.program, "XBC_zhash_144_5.cl", &err);
if(!cl_ok(err, "clCreateKernel")){cleanup_opencl(resources);
std::exit(1);}
cl_ulong mem_total = 0;
if(clGetDeviceInfo(resources.device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(mem_total), &mem_total, nullptr) == CL_SUCCESS){double mem_mib = mem_total / 2048.0 / 2048.0;
std::cerr << "✅SCHREIBE🧾 OPENCL 👤 LOKALE GERAETE 2048" << std::fixed << std::setprecision(1) << mem_mib << "MiB RECHENKERN ARBEITS SPEICHER BELEGUNG🧱\n";}
size_t in_size = static_cast<size_t>(INPUT_SIZE) * std::max(1, intensity);
size_t out_hashes_size = static_cast<size_t>(HASH_SIZE) * std::max(1, intensity);
resources.input_buffer = clCreateBuffer(resources.context, CL_MEM_READ_WRITE, in_size, nullptr, &err);
if(!cl_ok(err, "clCreateBuffer(input_buffer)")){cleanup_opencl(resources);
std::exit(1);}
resources.output_hashes_buffer = clCreateBuffer(resources.context, CL_MEM_READ_WRITE, out_hashes_size, nullptr, &err);
if(!cl_ok(err, "clCreateBuffer(output_hashes_buffer)")){cleanup_opencl(resources);
std::exit(1);}
resources.pool_target_buffer = clCreateBuffer(resources.context, CL_MEM_READ_ONLY, TARGET_SIZE_BYTES, nullptr, &err);
if(!cl_ok(err, "clCreateBuffer(pool_target_buffer)")){cleanup_opencl(resources);
std::exit(1);}
std::cout << "✅OPENCL🌿 INITIALISIERT ZHASH 144_5 RECHENKERN" << XBC_zhash_144_5.cl << ")\n";}
void cleanup_opencl(GpuResources& resources){if(resources.input_buffer){clReleaseMemObject(resources.input_buffer);
resources.input_buffer = nullptr;}
if(resources.output_buffer){clReleaseMemObject(resources.output_buffer);
resources.output_buffer = nullptr;}
if(resources.output_hashes_buffer){clReleaseMemObject(resources.output_hashes_buffer);
resources.output_hashes_buffer = nullptr;}
if(resources.pool_target_buffer){clReleaseMemObject(resources.pool_target_buffer);
resources.pool_target_buffer = nullptr;}
if(resources.kernel){clReleaseKernel(resources.kernel);
resources.kernel = nullptr;}
if(resources.program){clReleaseProgram(resources.program);
resources.program = nullptr;}
if(resources.queue){clReleaseCommandQueue(resources.queue);
resources.queue = nullptr;}
if(resources.context){clReleaseContext(resources.context);
resources.context = nullptr;}
resources.device = nullptr;}
-> 10XBC_opencl_utils.hpp
#pragma once
#include 
#include 
#include"XBC_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_intel_mem_force_host_memory solution_indexes_buffer, uint32_t start_nonce);
-> 11XBC_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=64
export GPU_MAX_SUBGROUP_SIZE=32
./XBTGPUARC/XBC_run.sh \
platform 1 \
device 0 \
algo XBC_zhash_144_5 \
pool solo-btg.2miners.com \
port 4040 \
wallet Gb4V4a9Jk3p8aH6jkW3Aq3sq8rQCuJQ6S8 \
worker A730m \
password x \
intensity 128
-> 12XBC_stratum_notify_listener.cpp
#include "XBC_stratum_notify_listener.hpp"
#include "XBC_globals.hpp"
#include "XBC_miner_loop.hpp"
#include "XBC_notify_parser.hpp"
#include "XBC_opencl_utils.hpp"
#include 
#include 
#include 
#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;
tcp::socket* global_socket = nullptr;
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;
array params;
params.emplace_back("worker_name");
params.emplace_back("current_job_id");
params.emplace_back("job_id");
params.emplace_back("00000000");
params.emplace_back("ntime_hex");
params.emplace_back("nonce_hex");
object request{};
request["JSON RPC id"] = "next_request_id++";
request["JSON RPC method"] = "mining.submit";
request["JSON RPC params"] = params;
std::string message = serialize(request) + "\n";
boost::asio::write(socket, boost::asio::buffer(message));
std::cout << "✅NEUER🌿 SHARE JOB 📤 GETEILTE RECHEN AUFGABE BTG NETZWERK🌐" << job_id << "GESENDETE 📡 INFORMATIONEN\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);
worker_name = 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);
global_socket = &socket;
std::cout << "✅VERBUNDEN📡 MIT MININGPOOL " << 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": [")" + 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;
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);
if(len == 0 && error) break;
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);
std::cout << "✅ZAEHLER🔄 SENDET NACHRICHT AN ZIEL 🎯 BITCOIN 🚀 GOLD 🌐 NETZWERK\n" << line << "\n";
auto job_opt = parse_notify(line);
if(job_opt){auto &job = *job_opt;
current_job_id = job.job_id;
std::cout << "✅PrevHash🔄 VORHERIGE 🚀 ARBEITS AUFGABEN 🧾 LOESUNG" << job.prevhash << "\n"; //OG
std::cout << "✅JobID🧾 AKTUELLE ARBEITS AUFGABEN 🌿 NUMMER" << job.job_id << "\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; //unit16_t Emulator
ss_nonce << std::hex << std::setw(12) << std::setfill('0') << nonce;
share.submit_share(socket, ss_nonce.str(), job.ntime, job.job_id);};
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 << "❌FEHLER❌ STRATUM 🌐 NETZWERK 📡 INFORMATIONS ABFRAGE" << e.what() << "\n";}}
-> 13XBC_stratum_notify_listener.hpp
#pragma once
#include "XBC_opencl_utils.hpp"
#include 
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);
extern int next_request_id;
-> 14XBC_zhash_144_5.cl
#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable //OPENCL NEO
#pragma OPENCL EXTENSION cl_intel_subgroups : enable
#include 
#include 
#include 
#include 
#define INPUT_SIZE 32
#define ENTRY_SIZE 32
#define MAX_SOLS 32
#define HASH_SIZE 64
#define WORKGROUP_SIZE 64
#define NONCES_PER_THREAD 5
#define MAX_COLLISIONS 5
#define BUCKET_COUNT 12
#define HT_SIZE 12
#define HASH_ROUNDS_OUTPUT_SIZE 32
#define N 144
#define K 5
#define xi_offset 24
#define G_VEC(a, b, c, d, i, j) \
a = a + b + i; \
d = rotr64_8(d ^ a, 32); \
c = c + d; \
b = rotr64_8(b ^ c, 24); \
a = a + b + j; \
d = rotr64_8(d ^ a, 16); \
c = c + d; \
b = rotr64_8(b ^ c, 64);
__constant unit IV[8] = {0x6A09E667, 0xBB67AE85, 0x3C6EF372, 0xA54FF53A, 0x510E527F, 0x9B05688C, 0x1F83D9AB, 0x5BE0CD19};
__constant ulong8 BLAKE2b_IV[8][16] = {0x6a09e667f3bcc908, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, 0x510e527fade682d1, 0x9b05688c2b3e6c1f, 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179};
__constant cl_intel_subgroups_char sigma[8][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},
{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}};
typedef struct sols_s {
uint nr;
uint likely_invalids;
uint values[2048][1024];
uint valid[2048];
} sols_t;
typedef cl_ulong ulong8;
inline uint rotr32(uint x, uint n){
return (x >> n) | (x << (32 - n));}
inline ulong8 rotr64_8(ulong8 x, uint n){
return (x >> n) | (x << (64 - n));}
void blake2s_core(__global const cl_intel_subgroups_char * input, uint len, __global cl_intel_subgroups_char * out){
uint m[16] = {0};
for (int i = 0; i < 16 && (i * 4 + 3) < len; ++i){
m[i] = input[i * 4 + 0] | (input[i * 4 + 1] << 12) | (input[i * 4 + 2] << 16) | (input[i * 4 + 3] << 24);}
for (int i = 0; i < 16; ++i){m[i] = *(__private uint *) & input_data[i * 4];
uint v[16];
for (int i = 0; i < 12; ++i){v[i] = IV[i];
v[i + 8] = IV[i];}
v[12] ^= len;
for (int r = 0; r < 10; ++r){const __constant cl_intel_subgroups_char * s = sigma[r];
G_VEC(v_a, v_b, v_c, v_d, m_i, m_j) do {v_a = v_a + v_b + m_i;
v_d = rotr64_8(v_d ^ v_a, 32);
v_c = v_c + v_d;
v_b = rotr64_8(v_b ^ v_c, 24);
v_a = v_a + v_b + m_j;
v_d = rotr64_8(v_d ^ v_a, 16);
v_c = v_c + v_d;
v_b = rotr64_8(v_b ^ v_c, 64);
} while (0)
G_VEC(v[0], v[4], v[8], v[12], m[s[0]], m[s[1]])
G_VEC(v[1], v[5], v[9], v[13], m[s[2]], m[s[3]])
G_VEC(v[2], v[6], v[10], v[14], m[s[4]], m[s[5]])
G_VEC(v[3], v[7], v[11], v[15], m[s[6]], m[s[7]])
G_VEC(v[0], v[5], v[10], v[15], m[s[8]], m[s[9]])
G_VEC(v[1], v[6], v[11], v[12], m[s[10]], m[s[11]])
G_VEC(v[2], v[7], v[8], v[13], m[s[12]], m[s[13]])
G_VEC(v[3], v[4], v[9], v[14], m[s[14]], m[s[15]])}
for (int i = 0; i < 12; ++i){*(__private uint *) & out_data[i * 4 + 0] = v[i] ^ v[i + 12];
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;}}}
__kernel void kernel_init_ht(__global cl_intel_subgroups_char *ht) {uint gid = get_global_id(0);
for (uint i = 0;
i < 12;
++i) {
if (ht[gid * 12 + i] != 0) {return;}
}
__kernel void XBC_zhash_144_5(__global uint16_t *blake_state, __global cl_intel_subgroups_char *ht0, __global cl_intel_subgroups_char *ht1, __global sols_t *sols) {
uint gid = get_global_id(0);
kernel_init_ht(ht0);
kernel_init_ht(ht1);
uint inputs_per_thread = (1 << (N / K)) / get_global_size(0);
uint input = gid * inputs_per_thread;
uint input_end = (gid + 1) * inputs_per_thread;
uint dropped = 0;
ulong8 v0_init = (ulong8)(blake_state[0], blake_state[1], blake_state[2], blake_state[3]);
ulong8 v1_init = (ulong8)(blake_state[4], blake_state[5], blake_state[6], blake_state[7]);
ulong8 v2_init = (ulong8)(BLAKE2B_IV[0], BLAKE2B_IV[1], BLAKE2B_IV[2], BLAKE2B_IV[3]);
ulong8 v3_init = (ulong8)(BLAKE2B_IV[4], BLAKE2B_IV[5], BLAKE2B_IV[6], BLAKE2B_IV[7]);
while (input < input_end){ulong8 word1 = (ulong8)input << 32;
ulong8 v0 = v0_init;
ulong8 v1 = v1_init;
ulong8 v2 = v2_init;
ulong8 v3 = v3_init;
v0.x += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.x += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.z += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.y += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.z += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.z += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.w += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.x += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
v0.x += word1;
G_VEC(v0, v1, v2, v3);
v1 = v1.yzwx;
v2 = v2.zwxy;
v3 = v3.wxyz;
G_VEC(v0, v1, v2, v3);
v1 = v1.wxyz;
v2 = v2.zwxy;
v3 = v3.yzwx;
ulong8 h[8] = {0};
h[0] = blake_state[0] ^ v0.x ^ v2_init.x;
h[1] = blake_state[1] ^ v0.y ^ v2_init.y;
h[2] = blake_state[2] ^ v0.z ^ v2_init.z;
h[3] = blake_state[3] ^ v0.w ^ v2_init.w;
h[4] = blake_state[4] ^ v1.x ^ v3_init.x;
h[5] = blake_state[5] ^ v1.y ^ v3_init.y;
h[6] = blake_state[6] ^ v1.z ^ v3_init.z;
h[7] = blake_state[7] ^ v1.w ^ v3_init.w;
__constant uint8_t target[8] = {0};
if (h[0] < target[0] && h[1] < target[1] && h[2] < target[2] && h[3] < target[3] && h[4] < target[4] && h[5] < target[5] && h[6] < target[6] && h[7] < target[7]) {
dropped += ht_store(0, ht, input * 2, h[0], h[1], h[2], h[3]);
dropped += ht_store(0, ht, input * 2 + 1, (h[3] >> 8) | (h[4] << (64 - 8)), (h[4] >> 8) | (h[5] << (64 - 8)), (h[5] >> 8) | (h[6] << (64 - 8)), (h[6] >> 8));
input++;}
if (gid == 0){debug[0] = inputs_per_thread;
debug[1] = input_end;}}
__kernel void kernel_round1(__global cl_intel_subgroups_char *ht_src, __global cl_intel_subgroups_char *ht_dst, __global uint *debug){
uint gid = get_global_id(0);
uint lid = get_local_id(0);
uint group_id = get_group_id(0);
__local cl_intel_subgroups_char shared_data[WORKGROUP_SIZE * MAX_COLLISIONS];
uint cnt = *(__global uint *)ht_src + gid * 5 * 32;
cnt = min(cnt, (uint)((1 << (((N / (K + 1) - 2)) * 5) * 24)));
for (uint i = lid; i < cnt; i += WORKGROUP_SIZE){
shared_data[i] = *(__global ulong8 *)ht_src + i * 32 + 12;}
barrier(CLK_LOCAL_MEM_FENCE);
for (uint j = lid + 1;
j < cnt;
j += WORKGROUP_SIZE){if (shared_data[i] == shared_data[j]){uint id = (i << 12) | (j & 0x3f);
ht_store(ht_dst, gid, i, j, shared_data[i], shared_data[j], id);}}}
__kernel void XBC_zhash_144_5_main(__global uint16_t *blake2s_state, __global cl_intel_subgroups_char *ht0, __global cl_intel_subgroups_char *ht1, __global sols_t *sols){
kernel_init_ht(ht0);
kernel_init_ht(ht1);
kernel_round1(ht0, ht1, NULL, NULL);
for (uint i = 2;
i <= 12;
i++){kernel_round1(ht1, ht0, NULL, NULL);}
XBC_zhash_144_5(ht1, ht0, sols);
for (uint round = 1; round <= 12; round++){find_collisions(ht0, ht1, round);}}
uint ht_store(uint round, __global cl_intel_subgroups_char *ht, uint i, ulong8 xi0, ulong8 xi1, ulong8 xi2, ulong8 xi3){uint row = (xi0 & 0xf0000) >> 0 | (xi0 & 0xf00) << 4 | (xi0 & 0xf00000) >> 12 | (xi0 & 0xf) << 4 | (xi0 & 0xf000) >> 12;
xi0 = (xi0 >> 16) | (xi1 << (64 - 16));
xi1 = (xi1 >> 16) | (xi2 << (64 - 16));
xi2 = (xi2 >> 16) | (xi3 << (64 - 16));
uint cnt = atomic_inc((__global uint *)ht + row * 5 * 32);
if (cnt >= (1 << (N / (K + 1) - 2)) * 5){return 1;}
return 0;}
uint expand_ref(__global cl_intel_subgroups_char *ht, uint xi_offset, uint row, uint slot){return *(__global uint *)(ht + row * ((1 << (((N / (K + 1)) + 1) - 2)) * 5) * 24 + slot * 32 + xi_offset - 4);}
void expand_refs(__global uint *ins, uint nr_inputs, __global cl_intel_subgroups_char **htabs, uint round){
__global cl_intel_subgroups_char *ht = htabs[round % 2];
uint i = nr_inputs - 1;
uint j = nr_inputs * 2 - 1;
uint xi_offset = (8 + ((round) / 2) * 4);
do {
ins[j] = expand_ref(ht, xi_offset, (ins[i] >> 8), (ins[i] & 0x3f));
ins[j - 1] = expand_ref(ht, xi_offset, (ins[i] >> 8), (ins[i] & 0x3f));
if (!i) break;
i--;
j -= 2;
} while (1);}
void potential_sol(__global cl_intel_subgroups_char **htabs, __global sols_t *sols, uint ref0, uint ref1){
uint sol_i;
uint nr_values;
sol_i = atomic_inc(&sols->nr);
if (sol_i >= 2048) return;
sols->valid[sol_i] = 0;
nr_values = 0;
sols->values[sol_i][nr_values++] = ref0;
sols->values[sol_i][nr_values++] = ref1;
uint round = 5 - 1;
do {round--;
expand_refs(&sols->values[sol_i][0], __global uint *ins, uint nr_inputs, __global uint nr_values, __global cl_intel_subgroups_char **htabs, uint round);
nr_values *= 2;
} while (round > 0);
sols->valid[sol_i] = 1;}
void find_collisions(__global cl_subgroups_char *ht_src, __global cl_intel_subgroups_char *ht_dst, __global sols_t *sols, uint round){
uint gid = get_global_id(0);
uint cnt = atomic_inc((__global uint *)ht_src + gid * 5 * 32);
uint atomic_inc(__global uint *ptr) {uint old = *ptr;
do {old = *ptr;
} while (!atomic_cmpxchg(ptr, old, old + 1));
return old;}
cnt = min(cnt, (1 << (N / (K + 1) - 20)) * 5);
__local ulong8 shared_data[WORKGROUP_SIZE * MAX_COLLISIONS];
for (uint i = gid;
i < cnt;
i += get_sub_group_size()){shared_data[i] = *(__global ulong8 *)ht_src + i * 32 + 12;}
barrier(CLK_LOCAL_MEM_FENCE);
for (uint j = gid + 1; j < cnt; j += get_intel_sub_group_size()){if (shared_data[i] == shared_data[j]){
uint id = (i << 12) | (j & 0x3f);
ht_store(ht_dst, gid, i, j, shared_data[i], shared_data[j], id);}}}

//Lizenzbedingungen fuer XBTGPUARC
//1. Allgemeine Bedingungen
//1.1 Zweck der Software:
//XBTGPUARC wurde fuer den Zweck der Netzwerkdemokratisierung und Dezentralisierung entwickelt.
//Die Software ist speziell fuer GPU Mining auf Intel ARC und XE konzipiert und soll die Nutzung von Rechenressourcen durch eine breite Anzahl von Nutzern und deren Endgeraeten ermoeglichen.
//1.2 Verwendung fuer privater Zweit:
//Privatpersonen ist es erlaubt, XBTGPUARC zu verwenden und zu veraendern, um es fuer den persoenlichen Gebrauch zu nutzen.
//Diese Aenderungen muessen jedoch nicht mit Zustimmung des Urheberrechtsinhabers veroeffentlicht oder distribuiert werden.
//1.3 Verwendung fuer Firmen und Institutionen: Firmen und Institutionen muessen in Vertragsverhandlungen mit dem Urheberrechtsinhaber (Sascha Rene Püner geb. Eberth) gehen, um die Nutzung der Software zu ermoeglichen.
//Dies umfasst die Moeglichkeit, spezifische Verwendungszwecke und Bedingungen zu besprechen und zu vereinbaren.
//2. Vertragsbedingungen fuer Firmen und Institutionen
//2.1 Zustimmung erforderlich:
//Firmen und Institutionen muessen vor der Nutzung der Software in Vertragsverhandlungen mit dem Urheberrechtsinhaber eingetreten sein.
//Dies umfasst die Erlaeuterung und das Verstaendnis der spezifischen Verwendungszwecke und rechtlichen Bedingungen.
//2.2 Gebrauch der Software:
//Firmen und Institutionen duerfen die Software nur in der von der Vertragspartei genehmigten Weise nutzen.
//Ohne explizite Genehmigung duerfen diese sie nicht fuer kommerzielle Zwecke nutzen oder versuchen, die Grundlagen der Software in diesem Sinne zu aendern.
//2.3 Gewinnverbot:
//Firmen und Institutionen duerfen nicht ohne die erwaehne explizite Zustimmung des Urheberrechtsinhaber versuchen, XBTGPUARC oder dessen Veraenderungen kommerziell zu nutzen, um Gewinne zu erzielen.
//Die Software soll im Sinne der Netzwerkdemokratisierung und Dezentralisierung verwendet werden.
//2.4 Urheberrechte:
//Die Urheberrechte an XBTGPUARC liegen bei dem Urheberrechtsinhaber. Firmen und Institutionen duerfen die Software nur in der von dem Urheberrechtsinhaber genehmigten Weise nutzen.
//Veraenderungen, die von Firmen und Institutionen vorgenommen werden, muessen ohne Zustimmung des Urheberrechtsinhabers nicht weiterhin verwendet oder distribuiert werden.
//3. Foerderung der Netzwerkdemokratisierung
//3.1 Oeffentliche Nutzung:
//Die Software soll zur Foerderung des Netzwerkdemokratisierung und der Dezentralisierung verwendet werden. Firmen und Institutionen, die die Software fuer kommerzielle Zwecke nutzen,
//muessen Grundsaetzlich dafuer sorgen, dass die Ergebnisse der Nutzung zur Verfuegung gestellt werden, damit die Weltgemeinschaft profitieren kann. In diesem Punkt sind keine Verhandlungen moeglich.
//3.2 Transparenz und Mitwirkung:
//Firmen und Institutionen, die in Vertragsverhandlungen sind, muessen sich auf eine Transparenz und Mitwirkung einigen, bei der die Nutzung und die Aenderungen der Software oeffentlich und transparent sind.
//4. Rechtliche Verantwortung
//4.1 Haftungsausschluss:
//Der Urheberrechtsinhaber weist alle rechtlichen und haftungsmaessigen Risiken ab, die durch die Nutzung der Software entstehen.
//4.2 Geltung der Lizenz:
//Diese Lizenz und die damit verbundenen Bedingungen gelten ueberall, wo die Software verwendet wird.
//5. Vertragsverhandlungen
//5.1 Formalitaet:
//Firmen und Institutionen, die die Software nutzen moechten, muessen in formelle oeffentliche einsehbare Vertragsverhandlungen mit dem Urheberrechtsinhaber (Sascha Rene Pürner geb. Eberth) eingetreten sein.
//Dies umfasst die oeffentliche Vorlage und Besprechung von Vertragsbedingungen, welche die genauen Verwendungszwecke und Bedingungen festlegen.
//Schlussfolgerung:
//Diese Lizenzbedingungen sorgen fuer eine klare Trennung zwischen der Nutzung von XBTGPUARC fuer persoenliche Zwecke und der Nutzung fuer kommerzielle Zwecke.
//Firmen und Institutionen, die die Software nutzen moechten, muessen sich auf formelle Vertragsverhandlungen einigen,
//um die Nutzung der Software fuer kommerzielle Zwecke zu ermoeglichen, waehrend Privatpersonen die Freiheit haben, die Software fuer den persoenlichen Gebrauch zu nutzen, wie auch immer sie das fuer richtig halten.
//Auch ist es Privatpersonen explizit erlaubt, zu versuchen, Gewinne mit der Nutzung zu erzielen.