arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-29 18:08:13 +00:00
Code Issues Releases Wiki Activity
17fdf8754d
hashcat/OpenCL/m16600_a3-optimized.cl

882 lines
32 KiB
Common Lisp
Raw Normal View History

Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
//#define NEW_SIMD_CODE
#include "inc_vendor.cl"
#include "inc_hash_constants.h"
#include "inc_hash_functions.cl"
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_cipher_aes.cl"
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
DECLSPEC void m16600 (SHM_TYPE u32a *s_te0, SHM_TYPE u32a *s_te1, SHM_TYPE u32a *s_te2, SHM_TYPE u32a *s_te3, SHM_TYPE u32a *s_te4, SHM_TYPE u32a *s_td0, SHM_TYPE u32a *s_td1, SHM_TYPE u32a *s_td2, SHM_TYPE u32a *s_td3, SHM_TYPE u32a *s_td4, u32 *w, const u32 pw_len, KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
/**
* loop
*/
u32 w0l = w[0];
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
{
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
const u32x w0r = ix_create_bft (bfs_buf, il_pos);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
const u32x w0 = w0l | w0r;
u32x w0_t = w0;
u32x w1_t = w[ 1];
u32x w2_t = w[ 2];
u32x w3_t = w[ 3];
u32x w4_t = w[ 4];
u32x w5_t = w[ 5];
u32x w6_t = w[ 6];
u32x w7_t = w[ 7];
u32x w8_t = w[ 8];
u32x w9_t = w[ 9];
u32x wa_t = w[10];
u32x wb_t = w[11];
u32x wc_t = w[12];
u32x wd_t = w[13];
u32x we_t = w[14];
u32x wf_t = w[15];
u32x a = SHA256M_A;
u32x b = SHA256M_B;
u32x c = SHA256M_C;
u32x d = SHA256M_D;
u32x e = SHA256M_E;
u32x f = SHA256M_F;
u32x g = SHA256M_G;
u32x h = SHA256M_H;
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C10);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C11);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C12);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C13);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C14);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C15);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C16);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C17);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C18);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C19);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C1a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C1b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C1c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C1d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C1e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C1f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C20);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C21);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C22);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C23);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C24);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C25);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C26);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C27);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C28);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C29);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C2a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C2b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C2c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C2d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C2e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C2f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C30);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C31);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C32);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C33);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C34);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C35);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C36);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C37);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C38);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C39);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C3a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C3b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C3c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C3d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C3e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C3f);
a += SHA256M_A;
b += SHA256M_B;
c += SHA256M_C;
d += SHA256M_D;
e += SHA256M_E;
f += SHA256M_F;
g += SHA256M_G;
h += SHA256M_H;
w0_t = a;
w1_t = b;
w2_t = c;
w3_t = d;
w4_t = e;
w5_t = f;
w6_t = g;
w7_t = h;
w8_t = 0x80000000;
w9_t = 0;
wa_t = 0;
wb_t = 0;
wc_t = 0;
wd_t = 0;
we_t = 0;
wf_t = 32 * 8;
a = SHA256M_A;
b = SHA256M_B;
c = SHA256M_C;
d = SHA256M_D;
e = SHA256M_E;
f = SHA256M_F;
g = SHA256M_G;
h = SHA256M_H;
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C10);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C11);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C12);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C13);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C14);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C15);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C16);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C17);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C18);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C19);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C1a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C1b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C1c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C1d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C1e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C1f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C20);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C21);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C22);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C23);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C24);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C25);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C26);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C27);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C28);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C29);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C2a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C2b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C2c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C2d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C2e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C2f);
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C30);
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C31);
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C32);
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C33);
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C34);
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C35);
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C36);
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C37);
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C38);
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C39);
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C3a);
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C3b);
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C3c);
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C3d);
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C3e);
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C3f);
a += SHA256M_A;
b += SHA256M_B;
c += SHA256M_C;
d += SHA256M_D;
e += SHA256M_E;
f += SHA256M_F;
g += SHA256M_G;
h += SHA256M_H;
u32 ukey[8];
ukey[0] = swap32_S (a);
ukey[1] = swap32_S (b);
ukey[2] = swap32_S (c);
ukey[3] = swap32_S (d);
ukey[4] = swap32_S (e);
ukey[5] = swap32_S (f);
ukey[6] = swap32_S (g);
ukey[7] = swap32_S (h);
#define KEYLEN 60
u32 ks[KEYLEN];
aes256_set_decrypt_key (ks, ukey, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);
u32 encrypted[4];
encrypted[0] = esalt_bufs[digests_offset].encrypted[0];
encrypted[1] = esalt_bufs[digests_offset].encrypted[1];
encrypted[2] = esalt_bufs[digests_offset].encrypted[2];
encrypted[3] = esalt_bufs[digests_offset].encrypted[3];
u32 out[4];
aes256_decrypt (ks, encrypted, out, s_td0, s_td1, s_td2, s_td3, s_td4);
u32 iv[4];
iv[0] = esalt_bufs[digests_offset].iv[0];
iv[1] = esalt_bufs[digests_offset].iv[1];
iv[2] = esalt_bufs[digests_offset].iv[2];
iv[3] = esalt_bufs[digests_offset].iv[3];
out[0] ^= iv[0];
out[1] ^= iv[1];
out[2] ^= iv[2];
out[3] ^= iv[3];
if (esalt_bufs[digests_offset].salt_type == 1)
{
if (is_valid_hex_32 (out[0]) == 0) continue;
if (is_valid_hex_32 (out[1]) == 0) continue;
if (is_valid_hex_32 (out[2]) == 0) continue;
if (is_valid_hex_32 (out[3]) == 0) continue;
Add salt type 2 support for optimized functions
2018-12-05 18:42:23 +00:00
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
if (esalt_bufs[digests_offset].salt_type == 2)
{
if ((u8) (out[0] >> 0) != 'x') continue;
if ((u8) (out[0] >> 8) != 'p') continue;
if ((u8) (out[0] >> 16) != 'r') continue;
if ((u8) (out[0] >> 24) != 'v') continue;
if (is_valid_base58_32 (out[1]) == 0) continue;
if (is_valid_base58_32 (out[2]) == 0) continue;
if (is_valid_base58_32 (out[3]) == 0) continue;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
if (atomic_inc (&hashes_shown[digests_offset]) == 0)
{
mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, digests_offset + 0, gid, il_pos);
}
}
}
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = 0;
w[ 5] = 0;
w[ 6] = 0;
w[ 7] = 0;
w[ 8] = 0;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_m08 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = 0;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_m16 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = pws[gid].i[ 8];
w[ 9] = pws[gid].i[ 9];
w[10] = pws[gid].i[10];
w[11] = pws[gid].i[11];
w[12] = pws[gid].i[12];
w[13] = pws[gid].i[13];
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_s04 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = 0;
w[ 5] = 0;
w[ 6] = 0;
w[ 7] = 0;
w[ 8] = 0;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_s08 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = 0;
w[ 9] = 0;
w[10] = 0;
w[11] = 0;
w[12] = 0;
w[13] = 0;
w[14] = 0;
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 22:35:58 +00:00
__kernel void m16600_s16 (KERN_ATTR_ESALT (electrum_wallet_t))
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
/**
* base
*/
const u64 gid = get_global_id (0);
const u64 lid = get_local_id (0);
const u64 lsz = get_local_size (0);
/**
* aes shared
*/
#ifdef REAL_SHM
__local u32 s_td0[256];
__local u32 s_td1[256];
__local u32 s_td2[256];
__local u32 s_td3[256];
__local u32 s_td4[256];
__local u32 s_te0[256];
__local u32 s_te1[256];
__local u32 s_te2[256];
__local u32 s_te3[256];
__local u32 s_te4[256];
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 16:04:33 +00:00
for (MAYBE_VOLATILE u32 i = lid; i < 256; i += lsz)
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
{
s_td0[i] = td0[i];
s_td1[i] = td1[i];
s_td2[i] = td2[i];
s_td3[i] = td3[i];
s_td4[i] = td4[i];
s_te0[i] = te0[i];
s_te1[i] = te1[i];
s_te2[i] = te2[i];
s_te3[i] = te3[i];
s_te4[i] = te4[i];
}
barrier (CLK_LOCAL_MEM_FENCE);
#else
__constant u32a *s_td0 = td0;
__constant u32a *s_td1 = td1;
__constant u32a *s_td2 = td2;
__constant u32a *s_td3 = td3;
__constant u32a *s_td4 = td4;
__constant u32a *s_te0 = te0;
__constant u32a *s_te1 = te1;
__constant u32a *s_te2 = te2;
__constant u32a *s_te3 = te3;
__constant u32a *s_te4 = te4;
#endif
if (gid >= gid_max) return;
/**
* base
*/
u32 w[16];
w[ 0] = pws[gid].i[ 0];
w[ 1] = pws[gid].i[ 1];
w[ 2] = pws[gid].i[ 2];
w[ 3] = pws[gid].i[ 3];
w[ 4] = pws[gid].i[ 4];
w[ 5] = pws[gid].i[ 5];
w[ 6] = pws[gid].i[ 6];
w[ 7] = pws[gid].i[ 7];
w[ 8] = pws[gid].i[ 8];
w[ 9] = pws[gid].i[ 9];
w[10] = pws[gid].i[10];
w[11] = pws[gid].i[11];
w[12] = pws[gid].i[12];
w[13] = pws[gid].i[13];
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 14:26:46 +00:00
const u32 pw_len = pws[gid].pw_len & 63;
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
/**
* main
*/
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 10:21:42 +00:00
m16600 (s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4, w, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3)
2018-01-25 14:28:21 +00:00
}
Reference in New Issue Copy Permalink