```cpp //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 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 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 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 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 abort_mining; extern std::atomic socket_valid; extern std::atomic 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& 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(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& 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 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::max()) return std::nullopt; return static_cast(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 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 get_current_job, const std::function&, 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 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 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 calculate_hash_with_nonce(const MiningJob &job, uint32_t nonce){return std::array(0);} bool is_valid_share(const std::array& hash, const std::array& target){return std::memcmp(hash.data(), target.data(), 32) < 0;} ){std::random_device rd; std::mt19937 rng(rd()); std::uniform_int_distribution dist; MiningJob current_job = get_current_job(); std::string current_job_id = current_job.job_id(); std::vector nonce = blake2b(nonce); std::vector 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 target = bits_to_target(* maybe_bits); std::vector 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 output_buffer(32 * batch_size); std::vector 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 target = bits_to_target(* maybe_bits); std::vector 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 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(idxA) * 32 + j; const size_t offsetB = static_cast(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&, const MiningJob&)> &on_valid_share, const GpuResources& resources, int intensity); std::string sanitize_hex_string(const std::string& input); std::optional 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 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 &hash, const std::vector &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 bits_to_target(uint32_t bits){uint32_t exponent = bits >> 24; uint32_t mantissa = bits& 0x007fffff; std::vector 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(32, 0xFF);} return target;} -> 7XBC_notify_parser.hpp #pragma once #include "XBC_mining_job.hpp" #include #include #include #include #include inline std::optional 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(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 ¶ms = 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(params.at(0)); job.version = value_to(params.at(1)); job.prevhash = value_to(params.at(2)); job.coinb1 = value_to(params.at(3)); job.coinb2 = value_to(params.at(4)); job.nbits = value_to(params.at(5)); job.ntime = value_to(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 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 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(ifs)), std::istreambuf_iterator());}} 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 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 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 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(INPUT_SIZE) * std::max(1, intensity); size_t out_hashes_size = static_cast(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& 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. ```

#include

Trending Posts