XBTGPUARC — Hive

//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 <CL/cl.h>
#include <array>
#include <cstdint>
#include <string>
#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 <CL/cl.h>
#include <array>
#include <cstdlib>
#include <iostream>
#include <random>
#include <string>
#include <vector>
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 <CL/cl.h>
#include <array>
#include <atomic>
#include <cctype>
#include <chrono>
#include <cstring>
#include <fstream>
#include <functional>
#include <iomanip>
#include <iostream>
#include <limits>
#include <optional>
#include <random>
#include <string>
#include <thread>
#include <vector>
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 <array>
#include <functional>
#include <string>
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 <array>
#include <cstdint>
#include <string>
#include <vector>
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 <boost/json.hpp>
#include <iostream>
#include <optional>
#include <string>
#include <vector>
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 <CL/cl.h>
#include <iostream>
#include <string>
#include <vector>
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 <CL/cl.h>
#include <fstream>
#include <iostream>
#include <vector>
#include <algorithm>
#include <iomanip>
#include <stdexcept>
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 <string>
#include <CL/cl.h>
#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 <boost/asio.hpp>
#include <boost/json.hpp>
#include <iomanip>
#include <iostream>
#include <sstream>
#include <string>
#include <thread>
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 <string>
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 <CL/cl2.hpp>
#include <vector>
#include <cstdint>
#include <atomic>
#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.
Sort:
Trending
[-]
bitandi (75) 14 days ago
Ich hoffe es geht Dir gut und alles lief ok mit deiner Unterkunft usw.
$0.02
2 votes
alucian (72) 13 days ago
danke dir andi, ich binineiner wwunderschoenen single bude in pankow und fuehle mich dort sehr wohl, leider nur im oeffentlichem internet alle paar tage. hoffe auch bei dir alles gut? mach dir keine sorgen liebe gruesse
1 vote
bitandi (75) 13 days ago
Freut mich zu hören.
Bei mir auch alles klar soweit.
LG
!BBH
$0.00
alucian (72) 9 days ago
Des freut mich ebenso wie deine Nachfrage. :-) Schoen dabei belassen! 😃