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