#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 //--Erklärung: 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 //--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.-- __constant uint 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.-- __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.-- __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} }; //--Inline Rotation Static zu Vector korrigieren- //--Erklärung: 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.-- 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). void blake2s_core( __global const uchar* input, uint len, __global uchar* out) { //--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; } } //--Erklärung: 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 char *ht) { uint gid = get_global_id(0); *(__global uint *)(ht + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32) = 0; } uint ht_store(__global char *ht, uint row_index, uint xi_id, ulong8 xi0, ulong8 xi1, ulong8 xi2, ulong8 xi3) { //--Offset in der Hashtabelle basierend auf dem Work-Item Index-- //--Erklärung: 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 char *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-- //-- Erklärung: (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.-- void find_collisions(__global char *ht_src, __global char *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 char *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)); //--Shared Memory Puffer für Workgroup-weite Kollisionserkennung-- __local ulong8 shared_data[WORKGROUP_SIZE * MAX_COLLISIONS]; //--Jeder 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 Workgroup-- 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; } //--Erklärung: Der Hauptkernel für die erste Runde (Round 0). Seine Aufgabe ist es, die initiale Menge von Hash-Werten zu generieren. //--Eingabe: Ein initialer Blake2b-Zustand (blake_state), der wahrscheinlich aus einem Block-Header abgeleitet ist. //--Verarbeitung: //--Jeder Thread berechnet eine Reihe von unterschiedlichen Hash-Werten, indem er einen Zähler (input) mit dem Blake2b-Zustand mischt (hier durch Addition von word1 = (ulong8)input << 32 simuliert). //--Für jeden Wert wird eine vereinfachte Version der Blake2b-Runden durchlaufen (dargestellt durch die G_VEC-Makros und die v1 = v1.yzwx-Permutationen, die die diagonale Mischung in Blake2b nachahmen). //--Der finale Hash-Wert h[0..7] wird durch XOR des ursprünglichen Zustands mit dem aktuellen Zustand und dem IV berechnet. //--Ausgabe: Die generierten Hash-Werte werden sofort mittels ht_store in der Hashtabelle ht abgelegt. Pro input-Wert werden oft zwei Einträge gespeichert, wobei der zweite ein um 8 Bit geshifteter Wert des ersten ist, um die Kollisionssuche zu starten. __kernel __attribute__((reqd_work_group_size(64,1,1))) void zhash_144_5(__global ulong8 *blake_state, __global char *ht, __global uint *debug) { uint gid = get_global_id(0); uint inputs_per_thread = (1 << (200 / 10)) / get_global_size(0); //--10 = 9+1-- uint input = tid * inputs_per_thread; uint input_end = (tid + 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]); v3_init.x ^= 140 + 4; v3_init.z ^= -1; 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; //--Runden 1 bis 9-- 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; //--Final XOR--(In Hauptprogramm eintragen)-- ulong8 h[8]; 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; 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; } } } //--Dies ist eine ulttimativ oft Bearbeitet Version für Intel Grafikkarten, bitte Lösche sie niemals einfach so. uint xor_and_store(uint round, __global char *ht_dst, uint row, uint slot_a, uint slot_b, __global ulong8 *a, __global ulong8 *b) { ulong8 xi0 = 0UL, xi1 = 0UL, xi2 = 0UL; ulong8 subgroup_reduce(ulong8 val) { for (int i = 1; i < get_sub_group_size(); i <<= 1) val = sub_group_shuffle_xor(val, i); return val; } //--Vektorisierte Verarbeitung mit ulong8 für maximale GPU-Effizienz-- if (round == 1 || round == 2) { xi0 = a[0] ^ b[0]; //--Volles ulong8 XOR-- xi1 = a[1] ^ b[1]; //--Volles ulong8 XOR-- xi2 = a[2] ^ b[2]; //--Volles ulong8 XOR-- if (round == 2) { //--Vektorisierte Shifts für alle 8 Elemente-- xi0 = (xi0 >> 8) | (xi1 << (64 - 8)); xi1 = (xi1 >> 8) | (xi2 << (64 - 8)); xi2 = (xi2 >> 8); } } else if (round == 3) { xi0 = a[0] ^ b[0]; xi1 = a[1] ^ b[1]; //--Für 32-bit Zugriffe: konvertiere zu uint8 für vektorisierten Zugriff-- xi2 = convert_ulong8(convert_uint8(a[2]) ^ convert_uint8(b[2])); } else if (round == 4 || round == 5) { xi0 = *a++ ^ *b++; xi1 = *a ^ *b; xi2 = 0; if (round == 4) { xi0 = (xi0 >> 8) | (xi1 << (64 - 8)); xi1 = (xi1 >> 8); } } else if (round == 6) { xi0 = *a++ ^ *b++; //--Für 32-bit Zugriffe: konvertiere zu uint8 für vektorisierten Zugriff-- xi1 = xi2 = convert_ulong8(convert_uint8(a[2]) ^ convert_uint8(b[2])); xi2 = 0; if (round == 6) { xi0 = (xi0 >> 8) | (xi1 << (64 - 8)); xi1 = (xi1 >> 8); } } else if (round == 7 || round == 8) { xi0 = a[0] ^ b[0]; xi1 = (ulong8)0; xi2 = (ulong8)0; if (round == 8) { xi0 = (xi0 >> 8); } } // Überprüfe ob alle Elemente in xi0 und xi1 Null sind ulong8 zero = (ulong8)0; bool all_zero = all(xi0 == zero) && all(xi1 == zero); if (all_zero) return 0; // ID aus row und slots codieren uint id = (row << 12) | ((slot_b & 0x3f) << 6) | (slot_a & 0x3f); // ulong8 Werte an ht_store übergeben return ht_store(round, ht_dst, id, xi0, xi1, xi2, (ulong8)0); } //--Erklärung: Diese Funktion ist das Herzstück der Equihash-Runden 1 bis 8. Sie sucht nach Paaren von Einträgen, die in den relevanten Bit-Positionen übereinstimmen (Kollisionen). Für jedes gefundene Paar ruft sie //--xor_and_store auf, um die Teillösung zu kombinieren und in der Zieltabelle für die nächste Runde zu speichern void equihash_round(uint round, __global char *ht_src, __global char *ht_dst, __global uint *debug) { uint tid = get_global_id(0); __global char *p; uint cnt; uchar first_words[((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)]; uchar mask; uint i, j; ushort collisions[((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 3]; uint nr_coll = 0; uint n; uint dropped_coll, dropped_stor; __global ulong8 *a, *b; uint xi_offset; xi_offset = (8 + ((round - 1) / 2) * 4); mask = 0; p = (ht_src + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32); cnt = *(__global uint *)p; cnt = min(cnt, (uint)((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)); p += xi_offset; for (i = 0; i < cnt; i++, p += 32) first_words[i] = *(__global uchar *)p; nr_coll = 0; dropped_coll = 0; for (i = 0; i < cnt; i++) for (j = i + 1; j < cnt; j++) if ((first_words[i] & mask) == (first_words[j] & mask)) { if (nr_coll >= sizeof (collisions) / sizeof (*collisions)) dropped_coll++; else collisions[nr_coll++] = ((ushort)j << 8) | ((ushort)i & 0xff); } dropped_stor = 0; for (n = 0; n < nr_coll; n++) { i = collisions[n] & 0xff; j = collisions[n] >> 8; a = (__global ulong8 *)(ht_src + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32 + i * 32 + xi_offset); b = (__global ulong8 *)(ht_src + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32 + j * 32 + xi_offset); dropped_stor += xor_and_store(round, ht_dst, tid, i, j, a, b); } if (round < 8) *(__global uint *)(ht_src + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32) = 0; } //--Equihash Round Kernels-- //--Eine Reihe von Kerneln für die Equihash-Runden 1 bis 8. Jeder Kernel führt equihash_round() für seine jeweilige Runde aus. //--equihash_round(round, ht_src, ht_dst, debug): //--Lade Kandidaten: Jeder Thread liest die Einträge "seiner" Zeile aus der Quelltabelle ht_src der vorherigen Runde. //--Suche Kollisionen: Innerhalb dieser Zeile sucht er nach Paaren von Einträgen, deren Hash-Werte in den für die aktuelle Runde relevanten Bit-Positionen übereinstimmen (die mask wird /verwendet, um die nicht relevanten Bits auszublenden). //--Kombiniere und speichere: Für jedes gefundene Paar wird xor_and_store aufgerufen, um die Teillösung zu kombinieren und in der Zieltabelle ht_dst für die nächste Runde zu speichern. //--In Runde 8 (kernel_round8) wird auch der Lösungszähler sols->nr zurückgesetzt. __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round1(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(1, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round2(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(2, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round3(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(3, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round4(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(4, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round5(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(5, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round6(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(6, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round7(__global char *ht_src, __global char *ht_dst, __global uint *debug) { equihash_round(7, ht_src, ht_dst, debug); } __kernel __attribute__((reqd_work_group_size(64, 1, 1))) void kernel_round8(__global char *ht_src, __global char *ht_dst, __global uint *debug, __global sols_t *sols) { uint tid = get_global_id(0); equihash_round(8, ht_src, ht_dst, debug); if (!tid) sols->nr = sols->likely_invalids = 0; } //--Rekonstruiert die Referenz (die xi_id) eines gespeicherten Eintrags. Diese ID kodiert die Herkunft des Eintrags (welche zwei Einträge der vorherigen Runde kombiniert wurden).-- uint expand_ref(__global char *ht, uint xi_offset, uint row, uint slot) { return *(__global uint *)(ht + row * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32 + slot * 32 + xi_offset - 4); } //--Rekonstruiert einen gesamten Lösungsbaum. Diese Funktion wird rekursiv für jede Runde aufgerufen. //--Sie nimmt ein Array von Referenzen (ins) aus der aktuellen Runde. //--Für jede Referenz schaut sie in der Hashtabelle der vorherigen Runde nach und ersetzt die Referenz durch die beiden Referenzen, aus denen sie in dieser vorherigen Runde erzeugt wurde. //--Dies verdoppelt die Anzahl der Referenzen in jedem Schritt, bis die ursprünglichen 200 Indizes aus der ersten Runde rekonstruiert sind.-- void expand_refs(__global uint *ins, uint nr_inputs, __global char **htabs, uint round) { __global 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] >> 12), ((ins[i] >> 6) & 0x3f)); ins[j - 1] = expand_ref(ht, xi_offset, (ins[i] >> 12), (ins[i] & 0x3f)); if (!i) break ; i--; j -= 2; } while (1); } //--Wird aufgerufen, wenn in der finalen Runde eine Kollision gefunden wurde. //--Reserviert einen Slot in der Lösungsliste (sols). //--Speichert die beiden finalen Referenzen. //--Ruft expand_refs für alle vorherigen Runden auf, um den gesamten Lösungsvektor (die 200 Indizes) zu rekonstruieren. //--Markiert die Lösung als gültig (valid = 1). void potential_sol(__global 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 >= 2000) 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 = 9 - 1; do { round--; expand_refs(&(sols->values[sol_i][0]), nr_values, htabs, round); nr_values *= 2; } while (round > 0); sols->valid[sol_i] = 1; } //--Hauptkernel für die finale Kollisionssuche Er kombiniert mehrere Konzepte: //--Laden in Shared Memory: Die relevanten Teile der globalen Hashtabelle werden in den schnellen Shared Memory der Workgroup geladen. //--Lokale Kollisionssuche: Alle Threads der Workgroup durchsuchen gemeinsam die im Shared Memory liegenden Daten nach Kollisionen. //--Dies ist effizienter, als wenn jeder Thread nur auf seinen eigenen globalen Speicherbereich zugreift. //--Lösungsausgabe: Gefundene Kollisionspaare werden zur endgültigen Verarbeitung an potential_sol übergeben. //--Grundkern!!! Wichtig Beibehaltung-- __kernel __attribute__((reqd_work_group_size(WORKGROUP_SIZE, 1, 1))) void zhash_144_5( __global ulong8 *blake2s_state, __global uchar *ht0, __global uchar *ht1, __global sols_t *sols, __local uint64x8 *shared_cache) //--Lokaler Speicher für Vektordaten-- { //--Jeder Thread bearbeitet seinen eigenen Datenblock-- uint gid = get_global_id(0); uint lid = get_local_id(0); uint group_id = get_group_id(0); //--Globale Tabellen anlegen-- __global uchar *htabs[2] = { ht0, ht1 }; uint ht_i = (9 - 1) % 2; //--WERT vom Tisch nehmen als Gültig-- //--Lokale Puffer Vektorisiert-- uint64x8 local_block[INPUT_SIZE / sizeof(uint64x8)]; uint64x8 local_collisions[MAX_COLLISIONS]; uint coll = 0; const uint mask = 0x00ffffffU; //--Schritt 1: Daten aus globalem Speicher in lokale Kopie laden-- __global uint64x8* global_block = ( __global uint64x8*)( htabs[ht_i] + gid * (MAX_ENTRIES * ENTRY_SIZE)); uint cnt = vload8(0, (__global uint*)global_block).s0; //--Erstes Element des Vektors-- cnt = min(cnt, (uint)((1 << (((200 / (9 + 1)) + 1) - 20)) * 9)); uint xi_offset = (8 + ((9 - 1) / 2) * 4); //--Daten in lokale Arbeitgruppe laden-- for (uint i = lid; i < cnt * (ENTRY_SIZE / sizeof(uint64x8)); i += WORKGROUP_SIZE) { shared_cache[i] = global_block[i + (8 + ((9 - 1) / 2) * 4) / sizeof(uint64x8)]; } barrier(CLK_LOCAL_MEM_FENCE); //--Synchronisiere vor Zugriff-- //--Schritt 2: Kollisionen im lokalen Puffer suchen-- for (uint i = 0; i < cnt; i++) { uint64x8 val_a = shared_cache[i] & mask; //--Maskierung vektorisiert-- //--Werte aus dem lokalen Block holen und maskieren-- for (uint j = i + 1; j < cnt; j++) { uint64x8 val_b = shared_cache[j] & mask; if (any(val_a == val_b)) { //--SIMD-Vergleich-- if (coll < MAX_COLLISIONS) { local_collisions[coll++] = (uint64x8)(i, j, 0, 0, 0, 0, 0, 0); //--PaketMarkierung-- } else { //--Puffer voll -> globales Zählen als "wahrscheinlich ungültig"-- atomic_inc(&sols->likely_invalids); } } } } //--Wenn keine Kollision gefunden wurde, fertig-- if (!coll) return; //--Schritt 3: Gefundene Kollisionen verarbeiten-- barrier(CLK_LOCAL_MEM_FENCE); //--Synchronisiere vor globalem Zugriff-- for (uint k = 0; k < coll; k++) { potential_sol(htabs, sols, local_collisions[k].s0, local_collisions[k].s1); } } //--Zusammenfassung ABLAUF //--1--Initialisierung: Die Hashtabelle wird geleert (kernel_init_ht). //--2--Runde 0: Der Kernel zhash_144_5 (der erste) füllt die Tabelle mit 2^(200/10) = 2^20 (≈1 Million) initialen Blake2b-Hashwerten. //--3--Runden 1-8: Die Kernel kernel_round1 bis kernel_round8 werden nacheinander ausgeführt. In jeder Runde i: //--3.1--Werden Kollisionen in den unteren Bits der Einträge aus Runde i-1 gesucht. //--3.2--Die gefundenen Paare werden geXORt und die Ergebnisse (nach einem Rechts-Shift) in die Tabelle für Runde i geschrieben. //--3.3--Die Größe der Problemstellung halbiert sich mit jeder Runde effektiv. //--4--Finale Prüfung: In Runde 8 sucht der Kernel nach einer Kollision über die verbleibenden Bits. Wird eine gefunden, ist eine potenzielle Gesamtlösung gefunden. //--5--Rekonstruktion: Die Funktion potential_sol rekonstruiert aus den beiden finalen Referenzen den gesamten Pfad durch alle vorherigen Runden zurück zu den ursprünglichen 200 Indizes. //--Diese Indizesequenz ist der gültige Proof-of-Work. //--Diese Anmerkungen bleiben im Kern bis alles Fertig ist und entsprechend Nachgearbeitet werden kann bei Bedarf

Topics