//--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 zhash_144_5((reqd_work_group_size(64,1,1)))
//--erzwingt Workgroups von 64 Threads (in x-Richtung), wichtig für die Parallelisierung.--
void zhash_144_5(__global ulong8 *blake_state, __global uchar *ht, __global uint *debug) {
uint gid = get_global_id(0); //--die globale ID des Threads.--
uint inputs_per_thread = (1 << (200 / 10)) / get_global_size(0); //--"10 = 9+1" wie viele Inputs ein einzelner Thread bearbeiten soll.--
//--Wird durch die Gesamtzahl der Threads geteilt.--
uint input = tid * inputs_per_thread;
uint input_end = (tid + 1) * inputs_per_thread;
uint dropped = 0;
//--Vier Vektoren "ulong8" bilden den Startzustand
//--"v2_init" und "v3_init" kommen aus den IV-Konstanten von BLAKE2b.
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]);
//--Ergibt 144. Der aktuelle Wert von "v3_init.x" wird bitweise mit "144"--
//--(in Binär geschrieben) XOR-verknüpft. Ergebnis ersetzt "v3_init.x".--
v3_init.x ^= 140 + 4;
//--Ergibt 144. Der aktuelle Wert von "v3_init.x" wird bitweise mit "144" (in Binär geschrieben) XOR-verknüpft. Ergebnis ersetzt "v3_init.x".--
v3_init.z ^= -1;
//-- in Binärdarstellung ist (im Zweierkomplement) eine Folge von lauter 1-Bits. Ein XOR mit lauter Einsen kehrt alle Bits um
//--Bitweise Invertierung. Ergebnis: "v3_init.z" wird umgedreht (Bitwise NOT).--
//--x bekommt durch das XOR eine Art „geheimes“ Verrechnen mit 144.
//--z wird komplett invertiert.
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 12--
//--Hier laufen die G-Funktionen von BLAKE2b, mehrfach mit Rotationen der Vektoren--
//--Das entspricht den 12 Standardrunden, nur hier auf 9 reduziert, wieder auf 12 aufbauen--
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)--
//--Damit wird ein Hashblock gebildet.--
//--Gemischte Initialwerten (v2_init, v3_init) und dem Originalstate (blake_state) bereiningen. Original Blake2b Einhalten.--
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;
//--Ergebnisse werden in ht geschrieben (HashTable?). 2-Slot Speicherung pro Input. Ändern.--
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 ultimativ oft Bearbeitet Version, für Intel ARC Grafikkarten, bitte Löschen Sie diese niemals einfach so!.--
//--Die Funktion enthält außerdem einen integrierten subgroup_reduce,--
//--der eigentlich nach außen müsste. Dieser Teil deutet darauf hin, dass verschiedene Varianten getestet wurden--
//--(Workgroup-basiert vs. Subgroup-basiert).“--
//--Diese Funktion kombiniert zwei Hashwerte (val1 und val2) durch XOR.--
//--Zusätzlich ist ein Subgroup-Reduce-Mechanismus eingebaut:--
//--Damit kann die GPU Threads innerhalb einer Subgroup (z. B. 8 oder 16 Threads)--
//--direkt synchronisieren, ohne globalen Speicher zu nutzen.--
//--Das deutet darauf hin, dass Equihash-Kollisionen in manchen Varianten--
//--direkt innerhalb einer Subgroup verarbeitet werden, ohne zurück ins große Hash-Table zu schreiben.--
uint xor_and_store(uint round, __global uchar *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]));
//--xi2 = *(__global uint *)a ^ *(__global uint *)b; /* FIXME: potential unaligned 32-bit read */--
} 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--
if (!xi0 && !xi1)
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, ((row << 12)) | ((slot_b & 0x3f) << 6) | (slot_a & 0x3f)), id, xi0, xi1, xi2, (ulong8)0);
}
//--In jeder Runde werden Hashes kombiniert und auf Kollisionen geprüft.--
//--Input: Hashes aus vorheriger Runde (hash_table + ht_prefix).--
//--Verarbeitung: find_collisions oder direkte xor_and_store--.
//--Output: neue Hashes für die nächste Runde.--
//--Wichtig: Nach der letzten Runde (n/k Runden) entstehen noch keine fertigen Lösungen.--
//--Erst die Kombination aller Pfade (in solve_equihash) erzeugt die 200-Bit gültigen Proofs.--
//--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--
//--Dieser TEil produziert noch keine Lösung. Die tatsächliche Lösung (200-Bit Wert)
//--entsteht erst nach Round 8 durch Rekombination aller Pfade.--
void equihash_round(uint round, __global uchar *ht_src, __global uchar *ht_dst, __global uint *debug)
{
uint tid = get_global_id(0);
__global uchar *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);
//--Verarbeitung von Kollisionslisten in einer Hash-Tabelle, wobei jeder Thread einen eigenen Bucket bearbeitet.--
mask = 0;
p = (ht_src + tid * ((1 << (((200 / (9 + 1)) + 1) - 20)) * 9) * 32);
//--ht_src: Basisadresse der Hash-Tabelle im Speicher--
//--tid: Thread-ID (welcher Thread gerade arbeitet) 9 + 1 = 10--
//--200 / 10 = 20--
//--20 + 1 = 21--
//--21 - 20 = 1--
//--1 << 1 = 2 (Bitshift nach links um 1 Position = Verdopplung)--
//--2 * 9 = 18--
//--18 * 32 = 576--
//--Vereinfacht: p = ht_src + tid * 576--
//--Berechne die Speicheradresse für diesen Thread. Jeder Thread bekommt einen eigenen
//--576-Byte-Bereich in der Hash-Tabelle zugewiesen.--
cnt = *(__global uint *)p;
//--*(...)p: Dereferenzierung - liest den Wert an Adresse p- Liest den Zählerwert aus dem Speicher an der berechneten Position "p".--
cnt = min(cnt, (uint)((1 << (((200 / (9 + 1)) + 1) - 20)) * 9));
//--Berechnung des Maximums--
//--Gleiche Berechnung wie oben: 2 * 9 = 18--
//--Zweck: Begrenzt den gelesenen Wert auf maximal 1--
p += xi_offset;
//--Zweck: Verschiebt den Pointer um eine bestimmte Offset-Distanz weiter--
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--
//--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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 uchar *ht_src, __global uchar *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 "i_id" eines gespeicherten Eintrags.--
//--Diese ID kodiert die Herkunft des Eintrags (welche zwei Einträge der vorherigen Runde kombiniert wurden).--
uint expand_ref(__global uchar *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 uchar **htabs, uint round)
{
__global uchar *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 uchar **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 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--
//--5--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--
//--1--Initialisierung
//--Es gibt feste Startwerte (IVs).
//--Daraus baust du deinen internen Zustand (meist 8 oder 16 Wörter).
//--2--Daten in Blöcke packen
//--Blake2b nimmt 128-Byte Blöcke.
//--Blake2s nimmt 64-Byte Blöcke.
//--3--Mix-Funktion (das Herzstück)
//--Jeder Block wird durch eine G-Funktion „durchgemischt“.
//--Dabei wird rotiert, addiert, XORt.
//--Genau das sorgt für die „Kryptomagie“.
//--4--Finalisierung
//--Am Ende werden die Zustände zusammengefügt.
//--Ergebnis ist dein Hash (z. B. 32 oder 64 Byte).
SCHLUSS TEIL 4/4