Teil 3/4 //--This Kernel is Based on the Github Bitcoin Gold Published Network!-- //--All this here, is reworked from Github to ARC and ALL other GPUs for run with the main programm in its base!-- //--All the other files in my Account are not, but this here is completly copied and reworked from me and Ai over and over again in now six month work beginn-- //--I never tryed out compile this until these days. I just learn and change some things the sounds good for me with the help of Vibe Coding and Leaning from AI.-- Here is the only other place you will find changes from me in the future. This is my Angel Version, the Evergreen Fallback to mark a history point in my personal life. Nothing nessecarry for you i just would say copie and work with it all the time you want. Enjoy upcoming updates in this post first and all the other three around. #pragma OPENCL EXTENSION cl_khr_byte_addressable_store : enable //--Ermöglicht den byteweisen Zugriff auf Speicherbereiche, was für die Handhabung von Hashes und Teil-Lösungen nützlich ist.-- #pragma OPENCL EXTENSION cl_intel_subgroups : enable //--Ermöglicht die Nutzung von Subgroups (Untergruppen von Work-Items innerhalb einer Workgroup), //--eine leistungsstarke Funktion für intra-Workgroup-Synchronisation und Datenaustausch auf Intel-GPUs.-- #define N 144 #define K 5 //--Die Kernparameter des Equihash-Algorithmus. N=144 definiert die Bitlänge der Lösung, K=5 bestimmt die Anzahl der Runden und die //--erforderliche Speicherbandbreite (der-- //--Speicherbedarf skaliert mit 2^(N/(K+1))).-- #define INPUT_SIZE (140) //--Beispielwert, anpassen-- #define ENTRY_SIZE (32) //--Beispielwert (Blake2b Blockgröße)-- #define MAX_SOLS (2000) //--200 sind angepeilt bei voller Leistung DG2-- #define MAX_COLLISIONS (16) //--Puffergröße für Kollisionen-- #define WORKGROUP_SIZE (64) //--Optimale Größe für ARC L1-Cache fraglich siehe weitere Anmerkungen-- #define HASH_SIZE 32 #define HT_SIZE 9 //--Konfigurationsparameter für den Kernel.-- //--INPUT_SIZE: Länge der Eingabedaten für den Hash (hier beispielhaft 140 Byte).-- //--ENTRY_SIZE: Größe eines Eintrags in der Hashtabelle (32 Byte für einen 256-Bit Hash-Wert + Metadaten).-- //--MAX_SOLS: Maximale Anzahl an Lösungen, die der Kernel pro Lauf zurückgeben kann.-- //--MAX_COLLISIONS: Maximale Anzahl von Kollisionen, die ein Work-Item lokal zwischenspeichern kann, bevor sie verarbeitet werden.-- //--WORKGROUP_SIZE: Optimale Anzahl von Threads (Work-Items) pro Workgroup. 64 ist eine gängige Größe, die gut auf den L1-Cache vieler //--GPU-Architekturen abgestimmt ist.-- //--Für Intel Arc Grafikkarten wurde ein höherer Wert möglich ermittelt. Da dieses Programm ausschließlich dem Mining mit ARC GPUs //--dienen soll, wird es auch höhere Werte beinhalten-- //--können und die modernen Speichermengen beachten. Hier reden wir von einem L1 Cache im durschnitt der großen DG2-Chips 512/448 //--von 2,5-4-6MiB L1 und mindestens 6-12-16 MiB L2 Cache. Diese-- //--Werte werden in das Programm als Hauptreferenz Einfließen, weil sie dem groben Druchschnitt entsprechen und nur höher werden in //--den bisher Battlemage Generationswechseln.-- //--Entsprechend weiterer Generationen wie Calestial und Druid wird hier nachgearbeitet werden müssen bei Bedarf.-- __constant uint unit IV[8] = { 0x6A09E667, 0xBB67AE85, 0x3C6EF372, 0xA54FF53A, 0x510E527F, 0x9B05688C, 0x1F83D9AB, 0x5BE0CD19 }; //--Konstantendefinitionen für die Hash-Funktionen.-- //--IV: Initialisierungsvektor (Initial Values) für Blake2s (32-Bit Worte).-- //--BLAKE2B_IV: Initialisierungsvektor für Blake2b (64-Bit Worte), wird im Hauptkernel "zhash_144_5" verwendet.-- //--Bei "IV[8]" handelt es sich um Blake2s (32-Bit Wörter), bei "BLAKE2B_IV[8]" um Blake2b (64-Bit).-- //--Die parallele Nutzung von Blake2s und Blake2b ist eine Eigenart des Equihash-Kernels, da Blake2s für kürzere Blöcke effizient ist,-- //--während Blake2b für die Hauptkompression genutzt wird.“-- __constant ulong8 BLAKE2b_IV[8] = { 0x6a09e667f3bcc908, 0xbb67ae8584caa73b, 0x3c6ef372fe94f82b, 0xa54ff53a5f1d36f1, 0x510e527fade682d1, 0x9b05688c2b3e6c1f, 0x1f83d9abfb41bd6b, 0x5be0cd19137e2179 }; //--"sigma": Eine Permutationstabelle, die die Reihenfolge festlegt, in der die Nachrichtenblöcke in jeder Runde der Blake2b-Kompressionsfunktion adressiert werden.-- //--Sie sorgt für Vermischung und Sicherheit.-- //--Konstant definiert und "12×16" groß, weil Blake2b 12 Runden hat.-- __constant uchar sigma[12][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[2000][512]; uint valid[2000]; } sols_t; inline uint rotr32(uint x, uint n) { return (x >> n) | (x << (32 - n)); } //--Inline Rotation Static zu Vector korrigieren.-- //--"ROTR64_8" Inline Rotation von Statisch Static zu Vector(en) korrigieren-- //--Führt eine bitweise Rotation nach rechts (right rotate) für einen Vektor aus acht 64-Bit-Ganzzahlen "ulong8" um "n" Bits durch. Dies ist-- //--eine fundamentale Operation in vielen-- //--Kryptographie-Hash-Funktionen.-- //--Diese Funktion rotiert jedes Element des Vektors "ulong8" unabhängig um "n" Bits. Dadurch laufen 8-- //--Rotationen parallel in einem "SIMD-Befehl“.-- //--"rotr64_8" nimmt einen "ulong8" (8 parallele 64-Bit-Werte im Vektorregister) und rotiert jedes Element um "n" Bits nach rechts.-- //--Das nutzt OpenCLs SIMD-Architektur: statt 8 Rotationen nacheinander zu machen,-- //--passieren alle gleichzeitig in einem einzigen Vektor-Befehl.-- inline ulong8 rotr64_8(ulong8 x, uint n) { return (x >> n) | (x << (64 - n)); } //--Eine Implementierung der Blake2s-Kompressionsfunktion. Sie komprimiert einen Eingabeblock (input)-- //--fester Länge unter Verwendung des internen Zustands (Teile des IV) und erzeugt einen-- //--Ausgabewert (out). Diese Funktion ist für die anfängliche Generierung der Hash-Werte verantwortlich.-- //--G_VEC: Das Makro innerhalb der Funktion definiert die Mischoperation für einen Vektor aus vier Zustandsworten""v_a, v_b, v_c, v_d" unter //--Verwendung zweier Nachrichtenworte "m_i, m_j".-- //--Entwurf: für Blake2s. Einige Variablen wie input_data und out_data sind Platzhalter und müssten im Kernelkontext definiert werden.“-- //--Entwurf: Blake2s (32-Bit-Variante) den "sigma-Permutations-Array" verwendet.-- //--Implementierung "h[] (Hash-State), input_data[] (Block), und out_data[] (Output)" korrekt initialisieren. Der Code Skizze Blake2s-- //--Parallelisiert werden könnte, keine vollständige Funktion.-- void blake2s_core( __global const uchar* input, uint len, __global uchar* 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] << 8) | (input[i*4 + 2] << 16) | (input[i*4 + 3] << 24); } //--Lokales Laden des Input-Blocks-- for (int i = 0; i < 16; ++i) m[i] = *(__private uint*)&input_data[i * 4]; uint v[16]; for (int i = 0; i < 8; ++i) { v[i] = IV[i]; v[i + 8] = IV[i]; } v[12] ^= len; for (int r = 0; r < 10; ++r) { const __constant uchar* s = sigma[r]; //--G_VEC (512-bit, Arbeitet auf ulong8!!!)-- #define 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, 63); \ } 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]]); #undef G } for (int i = 0; i < 8; ++i) { *(__private uint*)&out_data[i * 4] = v[i] ^ v[i + 8]; out[i*4 + 0] = h & 0xFF; out[i*4 + 1] = (h >> 8) & 0xFF; out[i*4 + 2] = (h >> 16) & 0xFF; out[i*4 + 3] = (h >> 24) & 0xFF; } } //--Ein sehr simpler Kernel, der die Hashtabelle initialisiert. Jeder Thread setzt den Zähler für einen bestimmten Bereich der Tabelle auf 0.-- __kernel void kernel_init_ht(__global uchar *ht) { uint gid = get_global_id(0); *(__global uint *)(ht + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32) = 0; } uint ht_store(uint round, __global uchar *ht, uint i, ulong8 xi0, ulong8 xi1, ulong8 xi2, ulong8 xi3) { uint row; __global uchar *p; uint cnt; row = select( (uint)(((xi0 & 0xf0000) >> 0) | ((xi0 & 0xf00) << 4) | ((xi0 & 0xf00000) >> 12) | ((xi0 & 0xf) << 4) | ((xi0 & 0xf000) >> 12)), (uint)((xi0 & 0xffff) | ((xi0 & 0xf00000) >> 4)), !(round % 2) ); xi0 = (xi0 >> 16) | (xi1 << (64 - 16)); xi1 = (xi1 >> 16) | (xi2 << (64 - 16)); xi2 = (xi2 >> 16) | (xi3 << (64 - 16)); p = ht + row * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32; cnt = atomic_inc((__global uint *)p); if (cnt >= ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)) return 1; p += cnt * 32 + (8 + ((round) / 2) * 4); *(__global uint *)(p - 4) = i; if (round == 0 || round == 1) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; *(__global ulong8 *)(p + 16) = xi2; } else if (round == 2) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; *(__global uint *)(p + 16) = xi2; } else if (round == 3 || round == 4) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; } else if (round == 5) { *(__global ulong8 *)(p + 0) = xi0; *(__global uint *)(p + 8) = xi1; } else if (round == 6 || round == 7) { *(__global ulong8 *)(p + 0) = xi0; } else if (round == 8) { *(__global uint *)(p + 0) = xi0; } return 0; } //--ht_store reduziert pro Runde die zu speichernde Datenmenge.-- //--In Runde 0 werden alle Hashes komplett gespeichert.-- //--In späteren Runden werden nur die oberen Bits gespeichert (weil die unteren Bits schon kollidiert sind).-- //--Das ist der Schlüssel zum Memory-Efficiency-Trick von Equihash: weniger Speicherbedarf, aber weiterhin vollständige Kollisionssuche möglich-- //--"Andere Rundenlogik“ ist entscheidend. Das ist der Kern des Speicher-Spar-Tricks von Equihash.-- //--In jeder Runde wird die gespeicherte Datenmenge reduziert, da nur die höheren Bits relevant bleiben.-- //--Dies reduziert Speicherlast und zwingt die GPU, Kollisionen effizienter zu verarbeiten.“-- //--Offset in der Hashtabelle basierend auf dem Work-Item Index-- //--Speichert einen Hash-Wert "xi0, xi1, xi2, xi3" zusammen mit einer ID "xi_id" in-- //--der Hashtabelle ht an einer bestimmten Zeile "row_index" Funktionsweise:-- //--Berechnet den Zeiger "p" auf den Anfang der gewünschten Zeile in der Tabelle.-- //--Erhöht atomar den Zähler für die Anzahl der Einträge in dieser Zeile "atomic_inc". Dies verhindert Race Conditions, wenn mehrere Threads gleichzeitig-- //--in dieselbe Zeile schreiben wollen.-- //--Wenn die Zeile voll ist, wird 1 (Fehler) zurückgegeben.-- //--Andernfalls werden die Daten an der entsprechenden Position "p + cnt * 32" gespeichert.-- //--Wichtig: Die genaue Struktur und Größe der gespeicherten Daten hängt von der aktuellen Equihash-Runde "round" ab. In späteren Runden werden weniger Daten gespeichert nur die-- //--höherwertigen Bits), um Speicher zu sparen.-- __global uchar *p = ht + row_index * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32; uint cnt = atomic_inc((__global uint *)p); if (cnt >= ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)) { return 1; } p = cnt * 32 + (8); //--8 ist die Größe des Blake-Zustands-- *(__global uint *)(p - 4) = xi_id; *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; *(__global ulong8 *)(p + 16) = xi2; //--Andere Rundenlogik-- return 0; } //--Hilfsfunktion zur Suche nach Kollisionen innerhalb einer Workgroup-- //--Diese Funktion scheint ein Entwurf oder eine alternative Implementierung zu sein-- //--und wird nicht von den Hauptkerneln aufgerufen). Die Idee dahinter ist:-- //--Verwendet Shared Memory "__local", um Daten innerhalb einer Workgroup zwischen allen Threads zugänglich zu machen.-- //--Jeder Thread lädt seinen Teil der Hashtabelleneinträge in diesen schnellen, gemeinsamen Speicher.-- //--Nach einer Synchronisationsbarriere "barrier" durchsucht jeder Thread den gemeinsamen Datensatz nach Kollisionen (Werten, die in-- //--bestimmten Bit-Positionen übereinstimmen).-- //--Dies kann die Kollisionssuche innerhalb einer Workgroup erheblich beschleunigen.-- //--Dese Funktion wurde vermutlich als Optimierung für Intel GPUs mit großen Shared-Memory-Blöcken gedacht.-- //--Sie ersetzt die teure Suche über global memory durch eine-- //--kollaborative Suche im schnelleren "local memory".-- //--find_collisions versucht Kollisionen nicht im global memory (langsam), sondern im local memory (schnell, pro Workgroup).-- //--"local_ht[]" ist ein Hash-Table nur für Threads dieser Workgroup.-- //--Jeder Thread trägt seine Werte ein und sucht parallel nach gleichen Präfixen "Kollisionen".-- //--Das ist eine experimentelle Optimierung, die nicht in allen Treibern stabil läuft (Intel GPUs mögen sowas, Nvidia weniger).-- void find_collisions(__global uchar *ht_src, __global uchar *ht_dst, __global sols_t *sols, uint round) { uint gid = get_global_id(0); uint tid = get_local_id(0); uint group_id = get_group_id(0); __global uchar *p = ht_src + gid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32; uint cnt = *(__global uint *)p; cnt = min(cnt, (uint)((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)); //--Geteilter Speicherbereich für Arbeitsgruppenweite Kollisionserkennung-- __local ulong8 shared_data[WORKGROUP_SIZE * MAX_COLLISIONS]; //--Jeder Strang(Thread) lädt seinen Hash-Wert in den lokalen Speicher-- if (tid < cnt) { shared_data[tid] = *(__global ulong8 *)(ht_src + (gid * ENTRY_SIZE * cnt) + tid * ENTRY_SIZE + (8 + ((round-1) / 2) * 4)); } barrier(CLK_LOCAL_MEM_FENCE); //--Kollisionssuche innerhalb der Arbeitsgruppe-- uint coll_count = 0; for (uint i = tid; i < cnt; i += get_sub_group_size()) { ulong8 val_a = shared_data[i]; for (uint j = i + 1; j < cnt; j++) { ulong8 val_b = shared_data[j]; if (val_a.x == val_b.x) { //--Beispiel: Kollision auf dem ersten 64-bit Wort-- //--Kollision gefunden, verarbeite sie-- //--Subgroup/Untergruppen-Funktionen evtl Effizienter-- } } } } row = select( (uint)(((xi0 & 0xf0000) >> 0) | ((xi0 & 0xf00) << 4) | ((xi0 & 0xf00000) >> 12) | ((xi0 & 0xf) << 4) | ((xi0 & 0xf000) >> 12)), (uint)((xi0 & 0xffff) | ((xi0 & 0xf00000) >> 4)), !(round % 2) ); xi0 = (xi0 >> 16) | (xi1 << (64 - 16)); xi1 = (xi1 >> 16) | (xi2 << (64 - 16)); xi2 = (xi2 >> 16) | (xi3 << (64 - 16)); p = ht + row * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32; cnt = atomic_inc((__global uint *)p); //--Hier ATOMIC_INC-- if (cnt >= ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)) return 1; p = cnt * 32 + (8 + ((round) / 2) * 4); *(__global uint *)(p - 4) = i; if (round == 0 || round == 1) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; *(__global ulong8 *)(p + 16) = xi2; } else if (round == 2) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; *(__global uint *)(p + 16) = xi2; } else if (round == 3 || round == 4) { *(__global ulong8 *)(p + 0) = xi0; *(__global ulong8 *)(p + 8) = xi1; } else if (round == 5) { *(__global ulong8 *)(p + 0) = xi0; *(__global uint *)(p + 8) = xi1; } else if (round == 6 || round == 7) { *(__global ulong8 *)(p + 0) = xi0; } else if (round == 8) { *(__global uint *)(p + 0) = xi0; } return 0; } SCHLUSS TEIL: 3/4 Zhash.cl intel arc + all

full

Experts

Show more authors