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--
//--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 6MiB L1 und mindestens 16MiB 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 BLAKE2b_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} };
//--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 BLAKE2b_IV_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; }
//--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 9--
//--Hier laufen die G-Funktionen von BLAKE2b, mehrfach mit Rotationen der Vektoren--
//--Das entspricht den 12 Standardrunden, nur hier auf 9 reduziert, evtl 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