mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-14 11:48:57 +00:00
Merge branch 'master' into fix_15300_bof
This commit is contained in:
Jens Steube
GitHub
commit
3c1542b0b1
@ -2066,6 +2066,12 @@ DECLSPEC int hc_enc_next (hc_enc_t *hc_enc, const u32 *src_buf, const int src_le
|
||||
|
||||
int src_pos = hc_enc->pos;
|
||||
|
||||
#if VENDOR_ID == 8
|
||||
// Work around segmentation fault in Intel JiT
|
||||
// Tested with 2021.12.6.0.19_160000
|
||||
volatile
|
||||
#endif
|
||||
|
||||
int dst_pos = hc_enc->clen;
|
||||
|
||||
dst_buf[0] = hc_enc->cbuf;
|
||||
@ -2197,6 +2203,12 @@ DECLSPEC int hc_enc_next_global (hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf,
|
||||
|
||||
int src_pos = hc_enc->pos;
|
||||
|
||||
#if VENDOR_ID == 8
|
||||
// Work around segmentation fault in Intel JiT
|
||||
// Tested with 2021.12.6.0.19_160000
|
||||
volatile
|
||||
#endif
|
||||
|
||||
int dst_pos = hc_enc->clen;
|
||||
|
||||
dst_buf[0] = hc_enc->cbuf;
|
||||
@ -2300,6 +2312,8 @@ DECLSPEC int hc_enc_next_global (hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf,
|
||||
|
||||
if ((dst_pos + 2) == dst_sz)
|
||||
{
|
||||
// this section seems to break intel opencl runtime but is unknown why
|
||||
|
||||
dst_ptr[dst_pos++] = (a >> 0) & 0xff;
|
||||
dst_ptr[dst_pos++] = (a >> 8) & 0xff;
|
||||
|
||||
|
||||
@ -113,7 +113,7 @@ DECLSPEC u64 blake2b_rot24_S (const u64 a)
|
||||
|
||||
#else
|
||||
|
||||
return hc_rotr64_S (a, 16);
|
||||
return hc_rotr64_S (a, 24);
|
||||
|
||||
#endif
|
||||
}
|
||||
|
||||
@ -269,6 +269,35 @@ DECLSPEC int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p
|
||||
return (len);
|
||||
}
|
||||
|
||||
DECLSPEC int mangle_toggle_at_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
|
||||
{
|
||||
if (len >= RP_PASSWORD_SIZE) return (len);
|
||||
|
||||
u8 occurence = 0;
|
||||
|
||||
u32 rem = 0;
|
||||
|
||||
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
|
||||
{
|
||||
const u32 t = buf[idx];
|
||||
|
||||
buf[idx] = t | generate_cmask (t);
|
||||
|
||||
u32 out = rem;
|
||||
|
||||
rem = 0;
|
||||
|
||||
if (((t >> 0) & 0xff) == p1) { if (occurence == p0) out = 0x0000ff00; occurence++; }
|
||||
if (((t >> 8) & 0xff) == p1) { if (occurence == p0) out = 0x00ff0000; occurence++; }
|
||||
if (((t >> 16) & 0xff) == p1) { if (occurence == p0) out = 0xff000000; occurence++; }
|
||||
if (((t >> 24) & 0xff) == p1) { if (occurence == p0) rem = 0x000000ff; occurence++; }
|
||||
|
||||
buf[idx] = t ^ (generate_cmask (t) & out);
|
||||
}
|
||||
|
||||
return (len);
|
||||
}
|
||||
|
||||
DECLSPEC int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
|
||||
{
|
||||
for (int l = 0; l < len / 2; l++)
|
||||
@ -725,6 +754,7 @@ DECLSPEC int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED
|
||||
case RULE_OP_MANGLE_UREST_LFIRST: out_len = mangle_urest_lfirst (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_TREST: out_len = mangle_trest (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_TOGGLE_AT: out_len = mangle_toggle_at (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_TOGGLE_AT_SEP: out_len = mangle_toggle_at_sep (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_REVERSE: out_len = mangle_reverse (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_DUPEWORD: out_len = mangle_dupeword (p0, p1, buf, out_len); break;
|
||||
case RULE_OP_MANGLE_DUPEWORD_TIMES: out_len = mangle_dupeword_times (p0, p1, (u8 *) buf, out_len); break;
|
||||
|
||||
@ -21,6 +21,7 @@
|
||||
#define RULE_OP_MANGLE_UREST_LFIRST 'C'
|
||||
#define RULE_OP_MANGLE_TREST 't'
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT 'T'
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT_SEP '3'
|
||||
#define RULE_OP_MANGLE_REVERSE 'r'
|
||||
#define RULE_OP_MANGLE_DUPEWORD 'd'
|
||||
#define RULE_OP_MANGLE_DUPEWORD_TIMES 'p'
|
||||
@ -83,6 +84,7 @@ DECLSPEC int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u
|
||||
DECLSPEC int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_toggle_at_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len);
|
||||
DECLSPEC int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len);
|
||||
|
||||
@ -1163,6 +1163,82 @@ DECLSPEC u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED c
|
||||
return (in_len);
|
||||
}
|
||||
|
||||
DECLSPEC u32 rule_op_mangle_toggle_at_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
|
||||
{
|
||||
if (in_len == 0) return in_len;
|
||||
|
||||
u32 r0 = search_on_register (buf0[0], p1);
|
||||
u32 r1 = search_on_register (buf0[1], p1);
|
||||
u32 r2 = search_on_register (buf0[2], p1);
|
||||
u32 r3 = search_on_register (buf0[3], p1);
|
||||
u32 r4 = search_on_register (buf1[0], p1);
|
||||
u32 r5 = search_on_register (buf1[1], p1);
|
||||
u32 r6 = search_on_register (buf1[2], p1);
|
||||
u32 r7 = search_on_register (buf1[3], p1);
|
||||
|
||||
const u32 rn = (r0 << 0)
|
||||
| (r1 << 4)
|
||||
| (r2 << 8)
|
||||
| (r3 << 12)
|
||||
| (r4 << 16)
|
||||
| (r5 << 20)
|
||||
| (r6 << 24)
|
||||
| (r7 << 28);
|
||||
|
||||
if (rn == 0) return in_len;
|
||||
|
||||
u32 occurence = 0;
|
||||
|
||||
u32 ro = 0;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 32; i++)
|
||||
{
|
||||
if ((rn >> i) & 1)
|
||||
{
|
||||
if (occurence == p0)
|
||||
{
|
||||
ro = 1 << i;
|
||||
|
||||
break;
|
||||
}
|
||||
|
||||
occurence++;
|
||||
}
|
||||
}
|
||||
|
||||
r0 = (ro >> 0) & 15;
|
||||
r1 = (ro >> 4) & 15;
|
||||
r2 = (ro >> 8) & 15;
|
||||
r3 = (ro >> 12) & 15;
|
||||
r4 = (ro >> 16) & 15;
|
||||
r5 = (ro >> 20) & 15;
|
||||
r6 = (ro >> 24) & 15;
|
||||
r7 = (ro >> 28) & 15;
|
||||
|
||||
r0 <<= 1;
|
||||
r1 <<= 1; r1 |= r0 >> 4;
|
||||
r2 <<= 1; r2 |= r1 >> 4;
|
||||
r3 <<= 1; r3 |= r2 >> 4;
|
||||
r4 <<= 1; r4 |= r3 >> 4;
|
||||
r5 <<= 1; r5 |= r4 >> 4;
|
||||
r6 <<= 1; r6 |= r5 >> 4;
|
||||
r7 <<= 1; r7 |= r6 >> 4;
|
||||
|
||||
buf0[0] = toggle_on_register (buf0[0], r0);
|
||||
buf0[1] = toggle_on_register (buf0[1], r1);
|
||||
buf0[2] = toggle_on_register (buf0[2], r2);
|
||||
buf0[3] = toggle_on_register (buf0[3], r3);
|
||||
buf1[0] = toggle_on_register (buf1[0], r4);
|
||||
buf1[1] = toggle_on_register (buf1[1], r5);
|
||||
buf1[2] = toggle_on_register (buf1[2], r6);
|
||||
buf1[3] = toggle_on_register (buf1[3], r7);
|
||||
|
||||
return in_len;
|
||||
}
|
||||
|
||||
DECLSPEC u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len)
|
||||
{
|
||||
reverse_block_optimized (buf0, buf1, buf0, buf1, in_len);
|
||||
@ -2285,6 +2361,7 @@ DECLSPEC u32 apply_rule_optimized (const u32 name, const u32 p0, const u32 p1, u
|
||||
case RULE_OP_MANGLE_UREST_LFIRST: out_len = rule_op_mangle_urest_lfirst (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_TREST: out_len = rule_op_mangle_trest (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_TOGGLE_AT: out_len = rule_op_mangle_toggle_at (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_TOGGLE_AT_SEP: out_len = rule_op_mangle_toggle_at_sep (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_REVERSE: out_len = rule_op_mangle_reverse (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_DUPEWORD: out_len = rule_op_mangle_dupeword (p0, p1, buf0, buf1, out_len); break;
|
||||
case RULE_OP_MANGLE_DUPEWORD_TIMES: out_len = rule_op_mangle_dupeword_times (p0, p1, buf0, buf1, out_len); break;
|
||||
|
||||
@ -21,6 +21,7 @@
|
||||
#define RULE_OP_MANGLE_UREST_LFIRST 'C'
|
||||
#define RULE_OP_MANGLE_TREST 't'
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT 'T'
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT_SEP '3'
|
||||
#define RULE_OP_MANGLE_REVERSE 'r'
|
||||
#define RULE_OP_MANGLE_DUPEWORD 'd'
|
||||
#define RULE_OP_MANGLE_DUPEWORD_TIMES 'p'
|
||||
@ -85,6 +86,7 @@ DECLSPEC u32 rule_op_mangle_lrest_ufirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSE
|
||||
DECLSPEC u32 rule_op_mangle_urest_lfirst (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_trest (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_toggle_at (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_toggle_at_sep (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_reverse (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_dupeword (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
DECLSPEC u32 rule_op_mangle_dupeword_times (MAYBE_UNUSED const u32 p0, MAYBE_UNUSED const u32 p1, MAYBE_UNUSED u32 *buf0, MAYBE_UNUSED u32 *buf1, const u32 in_len);
|
||||
|
||||
@ -16,7 +16,6 @@
|
||||
#define IS_OPENCL
|
||||
#endif
|
||||
|
||||
|
||||
#if defined IS_NATIVE
|
||||
#define CONSTANT_VK
|
||||
#define CONSTANT_AS
|
||||
|
||||
@ -804,7 +804,7 @@ KERNEL_FQ void m12500_loop (KERN_ATTR_TMPS (rar3_tmp_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
@ -970,7 +970,7 @@ KERNEL_FQ void m12500_comp (KERN_ATTR_TMPS (rar3_tmp_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
|
||||
@ -37,7 +37,7 @@ typedef struct zip2
|
||||
u32 verify_bytes;
|
||||
u32 compress_length;
|
||||
u32 data_len;
|
||||
u32 data_buf[0x4000000];
|
||||
u32 data_buf[0x200000];
|
||||
u32 auth_len;
|
||||
u32 auth_buf[4];
|
||||
|
||||
|
||||
@ -912,7 +912,7 @@ KERNEL_FQ void m23700_loop (KERN_ATTR_TMPS_ESALT (rar3_tmp_t, rar3_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
@ -1086,7 +1086,7 @@ KERNEL_FQ void m23700_comp (KERN_ATTR_TMPS_ESALT (rar3_tmp_t, rar3_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
|
||||
@ -823,7 +823,7 @@ KERNEL_FQ void m23800_loop (KERN_ATTR_TMPS_HOOKS_ESALT (rar3_tmp_t, rar3_hook_t,
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
@ -990,7 +990,7 @@ KERNEL_FQ void m23800_hook23 (KERN_ATTR_TMPS_HOOKS_ESALT (rar3_tmp_t, rar3_hook_
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len & 255;
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 salt_len = 8;
|
||||
|
||||
|
||||
@ -14,7 +14,7 @@
|
||||
#include "inc_hash_md5.cl"
|
||||
#endif
|
||||
|
||||
DECLSPEC u32x MurmurHash_w0 (const u32 seed, const u32x w0, const u32 *w, const int pw_len)
|
||||
DECLSPEC u32x MurmurHash_w0 (const u32 seed, const u32x w0, const u32x *w, const int pw_len)
|
||||
{
|
||||
u32x hash = seed;
|
||||
|
||||
|
||||
697
OpenCL/m27000-pure.cl
Normal file
697
OpenCL/m27000-pure.cl
Normal file
@ -0,0 +1,697 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
//#define NEW_SIMD_CODE
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md4.cl"
|
||||
#endif
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
#define PERM_OP(a,b,tt,n,m) \
|
||||
{ \
|
||||
tt = a >> n; \
|
||||
tt = tt ^ b; \
|
||||
tt = tt & m; \
|
||||
b = b ^ tt; \
|
||||
tt = tt << n; \
|
||||
a = a ^ tt; \
|
||||
}
|
||||
|
||||
#define HPERM_OP(a,tt,n,m) \
|
||||
{ \
|
||||
tt = a << (16 + n); \
|
||||
tt = tt ^ a; \
|
||||
tt = tt & m; \
|
||||
a = a ^ tt; \
|
||||
tt = tt >> (16 + n); \
|
||||
a = a ^ tt; \
|
||||
}
|
||||
|
||||
CONSTANT_VK u32a c_SPtrans[8][64] =
|
||||
{
|
||||
{
|
||||
0x02080800, 0x00080000, 0x02000002, 0x02080802,
|
||||
0x02000000, 0x00080802, 0x00080002, 0x02000002,
|
||||
0x00080802, 0x02080800, 0x02080000, 0x00000802,
|
||||
0x02000802, 0x02000000, 0x00000000, 0x00080002,
|
||||
0x00080000, 0x00000002, 0x02000800, 0x00080800,
|
||||
0x02080802, 0x02080000, 0x00000802, 0x02000800,
|
||||
0x00000002, 0x00000800, 0x00080800, 0x02080002,
|
||||
0x00000800, 0x02000802, 0x02080002, 0x00000000,
|
||||
0x00000000, 0x02080802, 0x02000800, 0x00080002,
|
||||
0x02080800, 0x00080000, 0x00000802, 0x02000800,
|
||||
0x02080002, 0x00000800, 0x00080800, 0x02000002,
|
||||
0x00080802, 0x00000002, 0x02000002, 0x02080000,
|
||||
0x02080802, 0x00080800, 0x02080000, 0x02000802,
|
||||
0x02000000, 0x00000802, 0x00080002, 0x00000000,
|
||||
0x00080000, 0x02000000, 0x02000802, 0x02080800,
|
||||
0x00000002, 0x02080002, 0x00000800, 0x00080802,
|
||||
},
|
||||
{
|
||||
0x40108010, 0x00000000, 0x00108000, 0x40100000,
|
||||
0x40000010, 0x00008010, 0x40008000, 0x00108000,
|
||||
0x00008000, 0x40100010, 0x00000010, 0x40008000,
|
||||
0x00100010, 0x40108000, 0x40100000, 0x00000010,
|
||||
0x00100000, 0x40008010, 0x40100010, 0x00008000,
|
||||
0x00108010, 0x40000000, 0x00000000, 0x00100010,
|
||||
0x40008010, 0x00108010, 0x40108000, 0x40000010,
|
||||
0x40000000, 0x00100000, 0x00008010, 0x40108010,
|
||||
0x00100010, 0x40108000, 0x40008000, 0x00108010,
|
||||
0x40108010, 0x00100010, 0x40000010, 0x00000000,
|
||||
0x40000000, 0x00008010, 0x00100000, 0x40100010,
|
||||
0x00008000, 0x40000000, 0x00108010, 0x40008010,
|
||||
0x40108000, 0x00008000, 0x00000000, 0x40000010,
|
||||
0x00000010, 0x40108010, 0x00108000, 0x40100000,
|
||||
0x40100010, 0x00100000, 0x00008010, 0x40008000,
|
||||
0x40008010, 0x00000010, 0x40100000, 0x00108000,
|
||||
},
|
||||
{
|
||||
0x04000001, 0x04040100, 0x00000100, 0x04000101,
|
||||
0x00040001, 0x04000000, 0x04000101, 0x00040100,
|
||||
0x04000100, 0x00040000, 0x04040000, 0x00000001,
|
||||
0x04040101, 0x00000101, 0x00000001, 0x04040001,
|
||||
0x00000000, 0x00040001, 0x04040100, 0x00000100,
|
||||
0x00000101, 0x04040101, 0x00040000, 0x04000001,
|
||||
0x04040001, 0x04000100, 0x00040101, 0x04040000,
|
||||
0x00040100, 0x00000000, 0x04000000, 0x00040101,
|
||||
0x04040100, 0x00000100, 0x00000001, 0x00040000,
|
||||
0x00000101, 0x00040001, 0x04040000, 0x04000101,
|
||||
0x00000000, 0x04040100, 0x00040100, 0x04040001,
|
||||
0x00040001, 0x04000000, 0x04040101, 0x00000001,
|
||||
0x00040101, 0x04000001, 0x04000000, 0x04040101,
|
||||
0x00040000, 0x04000100, 0x04000101, 0x00040100,
|
||||
0x04000100, 0x00000000, 0x04040001, 0x00000101,
|
||||
0x04000001, 0x00040101, 0x00000100, 0x04040000,
|
||||
},
|
||||
{
|
||||
0x00401008, 0x10001000, 0x00000008, 0x10401008,
|
||||
0x00000000, 0x10400000, 0x10001008, 0x00400008,
|
||||
0x10401000, 0x10000008, 0x10000000, 0x00001008,
|
||||
0x10000008, 0x00401008, 0x00400000, 0x10000000,
|
||||
0x10400008, 0x00401000, 0x00001000, 0x00000008,
|
||||
0x00401000, 0x10001008, 0x10400000, 0x00001000,
|
||||
0x00001008, 0x00000000, 0x00400008, 0x10401000,
|
||||
0x10001000, 0x10400008, 0x10401008, 0x00400000,
|
||||
0x10400008, 0x00001008, 0x00400000, 0x10000008,
|
||||
0x00401000, 0x10001000, 0x00000008, 0x10400000,
|
||||
0x10001008, 0x00000000, 0x00001000, 0x00400008,
|
||||
0x00000000, 0x10400008, 0x10401000, 0x00001000,
|
||||
0x10000000, 0x10401008, 0x00401008, 0x00400000,
|
||||
0x10401008, 0x00000008, 0x10001000, 0x00401008,
|
||||
0x00400008, 0x00401000, 0x10400000, 0x10001008,
|
||||
0x00001008, 0x10000000, 0x10000008, 0x10401000,
|
||||
},
|
||||
{
|
||||
0x08000000, 0x00010000, 0x00000400, 0x08010420,
|
||||
0x08010020, 0x08000400, 0x00010420, 0x08010000,
|
||||
0x00010000, 0x00000020, 0x08000020, 0x00010400,
|
||||
0x08000420, 0x08010020, 0x08010400, 0x00000000,
|
||||
0x00010400, 0x08000000, 0x00010020, 0x00000420,
|
||||
0x08000400, 0x00010420, 0x00000000, 0x08000020,
|
||||
0x00000020, 0x08000420, 0x08010420, 0x00010020,
|
||||
0x08010000, 0x00000400, 0x00000420, 0x08010400,
|
||||
0x08010400, 0x08000420, 0x00010020, 0x08010000,
|
||||
0x00010000, 0x00000020, 0x08000020, 0x08000400,
|
||||
0x08000000, 0x00010400, 0x08010420, 0x00000000,
|
||||
0x00010420, 0x08000000, 0x00000400, 0x00010020,
|
||||
0x08000420, 0x00000400, 0x00000000, 0x08010420,
|
||||
0x08010020, 0x08010400, 0x00000420, 0x00010000,
|
||||
0x00010400, 0x08010020, 0x08000400, 0x00000420,
|
||||
0x00000020, 0x00010420, 0x08010000, 0x08000020,
|
||||
},
|
||||
{
|
||||
0x80000040, 0x00200040, 0x00000000, 0x80202000,
|
||||
0x00200040, 0x00002000, 0x80002040, 0x00200000,
|
||||
0x00002040, 0x80202040, 0x00202000, 0x80000000,
|
||||
0x80002000, 0x80000040, 0x80200000, 0x00202040,
|
||||
0x00200000, 0x80002040, 0x80200040, 0x00000000,
|
||||
0x00002000, 0x00000040, 0x80202000, 0x80200040,
|
||||
0x80202040, 0x80200000, 0x80000000, 0x00002040,
|
||||
0x00000040, 0x00202000, 0x00202040, 0x80002000,
|
||||
0x00002040, 0x80000000, 0x80002000, 0x00202040,
|
||||
0x80202000, 0x00200040, 0x00000000, 0x80002000,
|
||||
0x80000000, 0x00002000, 0x80200040, 0x00200000,
|
||||
0x00200040, 0x80202040, 0x00202000, 0x00000040,
|
||||
0x80202040, 0x00202000, 0x00200000, 0x80002040,
|
||||
0x80000040, 0x80200000, 0x00202040, 0x00000000,
|
||||
0x00002000, 0x80000040, 0x80002040, 0x80202000,
|
||||
0x80200000, 0x00002040, 0x00000040, 0x80200040,
|
||||
},
|
||||
{
|
||||
0x00004000, 0x00000200, 0x01000200, 0x01000004,
|
||||
0x01004204, 0x00004004, 0x00004200, 0x00000000,
|
||||
0x01000000, 0x01000204, 0x00000204, 0x01004000,
|
||||
0x00000004, 0x01004200, 0x01004000, 0x00000204,
|
||||
0x01000204, 0x00004000, 0x00004004, 0x01004204,
|
||||
0x00000000, 0x01000200, 0x01000004, 0x00004200,
|
||||
0x01004004, 0x00004204, 0x01004200, 0x00000004,
|
||||
0x00004204, 0x01004004, 0x00000200, 0x01000000,
|
||||
0x00004204, 0x01004000, 0x01004004, 0x00000204,
|
||||
0x00004000, 0x00000200, 0x01000000, 0x01004004,
|
||||
0x01000204, 0x00004204, 0x00004200, 0x00000000,
|
||||
0x00000200, 0x01000004, 0x00000004, 0x01000200,
|
||||
0x00000000, 0x01000204, 0x01000200, 0x00004200,
|
||||
0x00000204, 0x00004000, 0x01004204, 0x01000000,
|
||||
0x01004200, 0x00000004, 0x00004004, 0x01004204,
|
||||
0x01000004, 0x01004200, 0x01004000, 0x00004004,
|
||||
},
|
||||
{
|
||||
0x20800080, 0x20820000, 0x00020080, 0x00000000,
|
||||
0x20020000, 0x00800080, 0x20800000, 0x20820080,
|
||||
0x00000080, 0x20000000, 0x00820000, 0x00020080,
|
||||
0x00820080, 0x20020080, 0x20000080, 0x20800000,
|
||||
0x00020000, 0x00820080, 0x00800080, 0x20020000,
|
||||
0x20820080, 0x20000080, 0x00000000, 0x00820000,
|
||||
0x20000000, 0x00800000, 0x20020080, 0x20800080,
|
||||
0x00800000, 0x00020000, 0x20820000, 0x00000080,
|
||||
0x00800000, 0x00020000, 0x20000080, 0x20820080,
|
||||
0x00020080, 0x20000000, 0x00000000, 0x00820000,
|
||||
0x20800080, 0x20020080, 0x20020000, 0x00800080,
|
||||
0x20820000, 0x00000080, 0x00800080, 0x20020000,
|
||||
0x20820080, 0x00800000, 0x20800000, 0x20000080,
|
||||
0x00820000, 0x00020080, 0x20020080, 0x20800000,
|
||||
0x00000080, 0x20820000, 0x00820080, 0x00000000,
|
||||
0x20000000, 0x20800080, 0x00020000, 0x00820080,
|
||||
}
|
||||
};
|
||||
|
||||
CONSTANT_VK u32a c_skb[8][64] =
|
||||
{
|
||||
{
|
||||
0x00000000, 0x00000010, 0x20000000, 0x20000010,
|
||||
0x00010000, 0x00010010, 0x20010000, 0x20010010,
|
||||
0x00000800, 0x00000810, 0x20000800, 0x20000810,
|
||||
0x00010800, 0x00010810, 0x20010800, 0x20010810,
|
||||
0x00000020, 0x00000030, 0x20000020, 0x20000030,
|
||||
0x00010020, 0x00010030, 0x20010020, 0x20010030,
|
||||
0x00000820, 0x00000830, 0x20000820, 0x20000830,
|
||||
0x00010820, 0x00010830, 0x20010820, 0x20010830,
|
||||
0x00080000, 0x00080010, 0x20080000, 0x20080010,
|
||||
0x00090000, 0x00090010, 0x20090000, 0x20090010,
|
||||
0x00080800, 0x00080810, 0x20080800, 0x20080810,
|
||||
0x00090800, 0x00090810, 0x20090800, 0x20090810,
|
||||
0x00080020, 0x00080030, 0x20080020, 0x20080030,
|
||||
0x00090020, 0x00090030, 0x20090020, 0x20090030,
|
||||
0x00080820, 0x00080830, 0x20080820, 0x20080830,
|
||||
0x00090820, 0x00090830, 0x20090820, 0x20090830,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x02000000, 0x00002000, 0x02002000,
|
||||
0x00200000, 0x02200000, 0x00202000, 0x02202000,
|
||||
0x00000004, 0x02000004, 0x00002004, 0x02002004,
|
||||
0x00200004, 0x02200004, 0x00202004, 0x02202004,
|
||||
0x00000400, 0x02000400, 0x00002400, 0x02002400,
|
||||
0x00200400, 0x02200400, 0x00202400, 0x02202400,
|
||||
0x00000404, 0x02000404, 0x00002404, 0x02002404,
|
||||
0x00200404, 0x02200404, 0x00202404, 0x02202404,
|
||||
0x10000000, 0x12000000, 0x10002000, 0x12002000,
|
||||
0x10200000, 0x12200000, 0x10202000, 0x12202000,
|
||||
0x10000004, 0x12000004, 0x10002004, 0x12002004,
|
||||
0x10200004, 0x12200004, 0x10202004, 0x12202004,
|
||||
0x10000400, 0x12000400, 0x10002400, 0x12002400,
|
||||
0x10200400, 0x12200400, 0x10202400, 0x12202400,
|
||||
0x10000404, 0x12000404, 0x10002404, 0x12002404,
|
||||
0x10200404, 0x12200404, 0x10202404, 0x12202404,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x00000001, 0x00040000, 0x00040001,
|
||||
0x01000000, 0x01000001, 0x01040000, 0x01040001,
|
||||
0x00000002, 0x00000003, 0x00040002, 0x00040003,
|
||||
0x01000002, 0x01000003, 0x01040002, 0x01040003,
|
||||
0x00000200, 0x00000201, 0x00040200, 0x00040201,
|
||||
0x01000200, 0x01000201, 0x01040200, 0x01040201,
|
||||
0x00000202, 0x00000203, 0x00040202, 0x00040203,
|
||||
0x01000202, 0x01000203, 0x01040202, 0x01040203,
|
||||
0x08000000, 0x08000001, 0x08040000, 0x08040001,
|
||||
0x09000000, 0x09000001, 0x09040000, 0x09040001,
|
||||
0x08000002, 0x08000003, 0x08040002, 0x08040003,
|
||||
0x09000002, 0x09000003, 0x09040002, 0x09040003,
|
||||
0x08000200, 0x08000201, 0x08040200, 0x08040201,
|
||||
0x09000200, 0x09000201, 0x09040200, 0x09040201,
|
||||
0x08000202, 0x08000203, 0x08040202, 0x08040203,
|
||||
0x09000202, 0x09000203, 0x09040202, 0x09040203,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x00100000, 0x00000100, 0x00100100,
|
||||
0x00000008, 0x00100008, 0x00000108, 0x00100108,
|
||||
0x00001000, 0x00101000, 0x00001100, 0x00101100,
|
||||
0x00001008, 0x00101008, 0x00001108, 0x00101108,
|
||||
0x04000000, 0x04100000, 0x04000100, 0x04100100,
|
||||
0x04000008, 0x04100008, 0x04000108, 0x04100108,
|
||||
0x04001000, 0x04101000, 0x04001100, 0x04101100,
|
||||
0x04001008, 0x04101008, 0x04001108, 0x04101108,
|
||||
0x00020000, 0x00120000, 0x00020100, 0x00120100,
|
||||
0x00020008, 0x00120008, 0x00020108, 0x00120108,
|
||||
0x00021000, 0x00121000, 0x00021100, 0x00121100,
|
||||
0x00021008, 0x00121008, 0x00021108, 0x00121108,
|
||||
0x04020000, 0x04120000, 0x04020100, 0x04120100,
|
||||
0x04020008, 0x04120008, 0x04020108, 0x04120108,
|
||||
0x04021000, 0x04121000, 0x04021100, 0x04121100,
|
||||
0x04021008, 0x04121008, 0x04021108, 0x04121108,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x10000000, 0x00010000, 0x10010000,
|
||||
0x00000004, 0x10000004, 0x00010004, 0x10010004,
|
||||
0x20000000, 0x30000000, 0x20010000, 0x30010000,
|
||||
0x20000004, 0x30000004, 0x20010004, 0x30010004,
|
||||
0x00100000, 0x10100000, 0x00110000, 0x10110000,
|
||||
0x00100004, 0x10100004, 0x00110004, 0x10110004,
|
||||
0x20100000, 0x30100000, 0x20110000, 0x30110000,
|
||||
0x20100004, 0x30100004, 0x20110004, 0x30110004,
|
||||
0x00001000, 0x10001000, 0x00011000, 0x10011000,
|
||||
0x00001004, 0x10001004, 0x00011004, 0x10011004,
|
||||
0x20001000, 0x30001000, 0x20011000, 0x30011000,
|
||||
0x20001004, 0x30001004, 0x20011004, 0x30011004,
|
||||
0x00101000, 0x10101000, 0x00111000, 0x10111000,
|
||||
0x00101004, 0x10101004, 0x00111004, 0x10111004,
|
||||
0x20101000, 0x30101000, 0x20111000, 0x30111000,
|
||||
0x20101004, 0x30101004, 0x20111004, 0x30111004,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x08000000, 0x00000008, 0x08000008,
|
||||
0x00000400, 0x08000400, 0x00000408, 0x08000408,
|
||||
0x00020000, 0x08020000, 0x00020008, 0x08020008,
|
||||
0x00020400, 0x08020400, 0x00020408, 0x08020408,
|
||||
0x00000001, 0x08000001, 0x00000009, 0x08000009,
|
||||
0x00000401, 0x08000401, 0x00000409, 0x08000409,
|
||||
0x00020001, 0x08020001, 0x00020009, 0x08020009,
|
||||
0x00020401, 0x08020401, 0x00020409, 0x08020409,
|
||||
0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
|
||||
0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
|
||||
0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
|
||||
0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
|
||||
0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
|
||||
0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
|
||||
0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
|
||||
0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x00000100, 0x00080000, 0x00080100,
|
||||
0x01000000, 0x01000100, 0x01080000, 0x01080100,
|
||||
0x00000010, 0x00000110, 0x00080010, 0x00080110,
|
||||
0x01000010, 0x01000110, 0x01080010, 0x01080110,
|
||||
0x00200000, 0x00200100, 0x00280000, 0x00280100,
|
||||
0x01200000, 0x01200100, 0x01280000, 0x01280100,
|
||||
0x00200010, 0x00200110, 0x00280010, 0x00280110,
|
||||
0x01200010, 0x01200110, 0x01280010, 0x01280110,
|
||||
0x00000200, 0x00000300, 0x00080200, 0x00080300,
|
||||
0x01000200, 0x01000300, 0x01080200, 0x01080300,
|
||||
0x00000210, 0x00000310, 0x00080210, 0x00080310,
|
||||
0x01000210, 0x01000310, 0x01080210, 0x01080310,
|
||||
0x00200200, 0x00200300, 0x00280200, 0x00280300,
|
||||
0x01200200, 0x01200300, 0x01280200, 0x01280300,
|
||||
0x00200210, 0x00200310, 0x00280210, 0x00280310,
|
||||
0x01200210, 0x01200310, 0x01280210, 0x01280310,
|
||||
},
|
||||
{
|
||||
0x00000000, 0x04000000, 0x00040000, 0x04040000,
|
||||
0x00000002, 0x04000002, 0x00040002, 0x04040002,
|
||||
0x00002000, 0x04002000, 0x00042000, 0x04042000,
|
||||
0x00002002, 0x04002002, 0x00042002, 0x04042002,
|
||||
0x00000020, 0x04000020, 0x00040020, 0x04040020,
|
||||
0x00000022, 0x04000022, 0x00040022, 0x04040022,
|
||||
0x00002020, 0x04002020, 0x00042020, 0x04042020,
|
||||
0x00002022, 0x04002022, 0x00042022, 0x04042022,
|
||||
0x00000800, 0x04000800, 0x00040800, 0x04040800,
|
||||
0x00000802, 0x04000802, 0x00040802, 0x04040802,
|
||||
0x00002800, 0x04002800, 0x00042800, 0x04042800,
|
||||
0x00002802, 0x04002802, 0x00042802, 0x04042802,
|
||||
0x00000820, 0x04000820, 0x00040820, 0x04040820,
|
||||
0x00000822, 0x04000822, 0x00040822, 0x04040822,
|
||||
0x00002820, 0x04002820, 0x00042820, 0x04042820,
|
||||
0x00002822, 0x04002822, 0x00042822, 0x04042822
|
||||
}
|
||||
};
|
||||
|
||||
#if VECT_SIZE == 1
|
||||
#define BOX(i,n,S) (S)[(n)][(i)]
|
||||
#elif VECT_SIZE == 2
|
||||
#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1])
|
||||
#elif VECT_SIZE == 4
|
||||
#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3])
|
||||
#elif VECT_SIZE == 8
|
||||
#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7])
|
||||
#elif VECT_SIZE == 16
|
||||
#define BOX(i,n,S) make_u32x ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
DECLSPEC void _des_crypt_encrypt (u32 *iv, u32 *data, u32 *Kc, u32 *Kd, SHM_TYPE u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32 r = data[0];
|
||||
u32 l = data[1];
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32 u;
|
||||
u32 t;
|
||||
|
||||
u = Kc[i + 0] ^ hc_rotl32 (r, 30u);
|
||||
t = Kd[i + 0] ^ hc_rotl32 (r, 26u);
|
||||
|
||||
l ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
|
||||
|
||||
u = Kc[i + 1] ^ hc_rotl32 (l, 30u);
|
||||
t = Kd[i + 1] ^ hc_rotl32 (l, 26u);
|
||||
|
||||
r ^= BOX (((u >> 0) & 0x3f), 0, s_SPtrans)
|
||||
| BOX (((u >> 8) & 0x3f), 2, s_SPtrans)
|
||||
| BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
|
||||
| BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
|
||||
| BOX (((t >> 0) & 0x3f), 1, s_SPtrans)
|
||||
| BOX (((t >> 8) & 0x3f), 3, s_SPtrans)
|
||||
| BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
|
||||
| BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
|
||||
}
|
||||
|
||||
iv[0] = l;
|
||||
iv[1] = r;
|
||||
}
|
||||
|
||||
DECLSPEC void _des_crypt_keysetup (u32 c, u32 d, u32 *Kc, u32 *Kd, SHM_TYPE u32 (*s_skb)[64])
|
||||
{
|
||||
u32 tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
HPERM_OP (d, tt, 2, 0xcccc0000);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
PERM_OP (c, d, tt, 8, 0x00ff00ff);
|
||||
PERM_OP (d, c, tt, 1, 0x55555555);
|
||||
|
||||
d = ((d & 0x000000ff) << 16)
|
||||
| ((d & 0x0000ff00) << 0)
|
||||
| ((d & 0x00ff0000) >> 16)
|
||||
| ((c & 0xf0000000) >> 4);
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i++)
|
||||
{
|
||||
if ((i < 2) || (i == 8) || (i == 15))
|
||||
{
|
||||
c = ((c >> 1) | (c << 27));
|
||||
d = ((d >> 1) | (d << 27));
|
||||
}
|
||||
else
|
||||
{
|
||||
c = ((c >> 2) | (c << 26));
|
||||
d = ((d >> 2) | (d << 26));
|
||||
}
|
||||
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32 c00 = (c >> 0) & 0x0000003f;
|
||||
const u32 c06 = (c >> 6) & 0x00383003;
|
||||
const u32 c07 = (c >> 7) & 0x0000003c;
|
||||
const u32 c13 = (c >> 13) & 0x0000060f;
|
||||
const u32 c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
|((c06 >> 8) & 0xff), 2, s_skb)
|
||||
| BOX (((c20 >> 0) & 0xff)
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32 d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32 d07 = (d >> 7) & 0x00003f03;
|
||||
const u32 d21 = (d >> 21) & 0x0000000f;
|
||||
const u32 d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
| BOX (((d21 >> 0) & 0xff)
|
||||
|((d22 >> 0) & 0xff), 7, s_skb);
|
||||
|
||||
Kc[i] = ((t << 16) | (s & 0x0000ffff));
|
||||
Kd[i] = ((s >> 16) | (t & 0xffff0000));
|
||||
}
|
||||
}
|
||||
|
||||
DECLSPEC void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 *out)
|
||||
{
|
||||
u32 t[8];
|
||||
|
||||
t[0] = (w0 >> 0) & 0xff;
|
||||
t[1] = (w0 >> 8) & 0xff;
|
||||
t[2] = (w0 >> 16) & 0xff;
|
||||
t[3] = (w0 >> 24) & 0xff;
|
||||
t[4] = (w1 >> 0) & 0xff;
|
||||
t[5] = (w1 >> 8) & 0xff;
|
||||
t[6] = (w1 >> 16) & 0xff;
|
||||
t[7] = (w1 >> 24) & 0xff;
|
||||
|
||||
u32 k[8];
|
||||
|
||||
k[0] = (t[0] >> 0);
|
||||
k[1] = (t[0] << 7) | (t[1] >> 1);
|
||||
k[2] = (t[1] << 6) | (t[2] >> 2);
|
||||
k[3] = (t[2] << 5) | (t[3] >> 3);
|
||||
k[4] = (t[3] << 4) | (t[4] >> 4);
|
||||
k[5] = (t[4] << 3) | (t[5] >> 5);
|
||||
k[6] = (t[5] << 2) | (t[6] >> 6);
|
||||
k[7] = (t[6] << 1);
|
||||
|
||||
out[0] = ((k[0] & 0xff) << 0)
|
||||
| ((k[1] & 0xff) << 8)
|
||||
| ((k[2] & 0xff) << 16)
|
||||
| ((k[3] & 0xff) << 24);
|
||||
|
||||
out[1] = ((k[4] & 0xff) << 0)
|
||||
| ((k[5] & 0xff) << 8)
|
||||
| ((k[6] & 0xff) << 16)
|
||||
| ((k[7] & 0xff) << 24);
|
||||
}
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
DECLSPEC u8 hex_convert (const u8 c)
|
||||
{
|
||||
return (c & 15) + (c >> 6) * 9;
|
||||
}
|
||||
|
||||
DECLSPEC u8 hex_to_u8 (const u8 *hex)
|
||||
{
|
||||
u8 v = 0;
|
||||
|
||||
v |= ((u8) hex_convert (hex[1]) << 0);
|
||||
v |= ((u8) hex_convert (hex[0]) << 4);
|
||||
|
||||
return (v);
|
||||
}
|
||||
#endif
|
||||
|
||||
typedef struct netntlm
|
||||
{
|
||||
u32 user_len;
|
||||
u32 domain_len;
|
||||
u32 srvchall_len;
|
||||
u32 clichall_len;
|
||||
|
||||
u32 userdomain_buf[64];
|
||||
u32 chall_buf[256];
|
||||
|
||||
} netntlm_t;
|
||||
|
||||
typedef struct netntlm_tmp
|
||||
{
|
||||
u32 digest_buf[4];
|
||||
|
||||
} netntlm_tmp_t;
|
||||
|
||||
KERNEL_FQ void m27000_init (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 in[16];
|
||||
|
||||
in[ 0] = pws[gid].i[ 0];
|
||||
in[ 1] = pws[gid].i[ 1];
|
||||
in[ 2] = pws[gid].i[ 2];
|
||||
in[ 3] = pws[gid].i[ 3];
|
||||
in[ 4] = pws[gid].i[ 4];
|
||||
in[ 5] = pws[gid].i[ 5];
|
||||
in[ 6] = pws[gid].i[ 6];
|
||||
in[ 7] = pws[gid].i[ 7];
|
||||
|
||||
u8 *in_ptr = (u8 *) in;
|
||||
|
||||
u32 out[4];
|
||||
|
||||
u8 *out_ptr = (u8 *) out;
|
||||
|
||||
for (int i = 0, j = 0; i < 16; i += 1, j += 2)
|
||||
{
|
||||
out_ptr[i] = hex_to_u8 (in_ptr + j);
|
||||
}
|
||||
|
||||
tmps[gid].digest_buf[0] = out[ 0];
|
||||
tmps[gid].digest_buf[1] = out[ 1];
|
||||
tmps[gid].digest_buf[2] = out[ 2];
|
||||
tmps[gid].digest_buf[3] = out[ 3];
|
||||
|
||||
}
|
||||
|
||||
KERNEL_FQ void m27000_loop (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
|
||||
}
|
||||
|
||||
KERNEL_FQ void m27000_comp (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 lsz = get_local_size (0);
|
||||
|
||||
/**
|
||||
* sbox, kbox
|
||||
*/
|
||||
|
||||
#ifdef REAL_SHM
|
||||
|
||||
LOCAL_VK u32 s_SPtrans[8][64];
|
||||
LOCAL_VK u32 s_skb[8][64];
|
||||
|
||||
for (u32 i = lid; i < 64; i += lsz)
|
||||
{
|
||||
s_SPtrans[0][i] = c_SPtrans[0][i];
|
||||
s_SPtrans[1][i] = c_SPtrans[1][i];
|
||||
s_SPtrans[2][i] = c_SPtrans[2][i];
|
||||
s_SPtrans[3][i] = c_SPtrans[3][i];
|
||||
s_SPtrans[4][i] = c_SPtrans[4][i];
|
||||
s_SPtrans[5][i] = c_SPtrans[5][i];
|
||||
s_SPtrans[6][i] = c_SPtrans[6][i];
|
||||
s_SPtrans[7][i] = c_SPtrans[7][i];
|
||||
|
||||
s_skb[0][i] = c_skb[0][i];
|
||||
s_skb[1][i] = c_skb[1][i];
|
||||
s_skb[2][i] = c_skb[2][i];
|
||||
s_skb[3][i] = c_skb[3][i];
|
||||
s_skb[4][i] = c_skb[4][i];
|
||||
s_skb[5][i] = c_skb[5][i];
|
||||
s_skb[6][i] = c_skb[6][i];
|
||||
s_skb[7][i] = c_skb[7][i];
|
||||
}
|
||||
|
||||
SYNC_THREADS ();
|
||||
|
||||
#else
|
||||
|
||||
CONSTANT_AS u32a (*s_SPtrans)[64] = c_SPtrans;
|
||||
CONSTANT_AS u32a (*s_skb)[64] = c_skb;
|
||||
|
||||
#endif
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 s0 = salt_bufs[SALT_POS].salt_buf[0];
|
||||
const u32 s1 = salt_bufs[SALT_POS].salt_buf[1];
|
||||
const u32 s2 = salt_bufs[SALT_POS].salt_buf[2];
|
||||
|
||||
const u32 a = tmps[gid].digest_buf[0];
|
||||
const u32 b = tmps[gid].digest_buf[1];
|
||||
const u32 c = tmps[gid].digest_buf[2];
|
||||
const u32 d = tmps[gid].digest_buf[3];
|
||||
|
||||
// I believe this matches the last 2 bytes and throws away.
|
||||
// Taken from 5500.
|
||||
if ((d >> 16) != s2) return;
|
||||
|
||||
/**
|
||||
* DES1
|
||||
*/
|
||||
|
||||
u32 key[2];
|
||||
|
||||
transform_netntlmv1_key (a, b, key);
|
||||
|
||||
u32 Kc[16];
|
||||
u32 Kd[16];
|
||||
|
||||
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
|
||||
|
||||
u32 data[2];
|
||||
|
||||
data[0] = s0;
|
||||
data[1] = s1;
|
||||
|
||||
u32 out1[2];
|
||||
|
||||
_des_crypt_encrypt (out1, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
/**
|
||||
* DES2
|
||||
*/
|
||||
|
||||
transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
|
||||
|
||||
_des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
|
||||
|
||||
u32 out2[2];
|
||||
|
||||
_des_crypt_encrypt (out2, data, Kc, Kd, s_SPtrans);
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 r0 = out1[0];
|
||||
const u32 r1 = out1[1];
|
||||
const u32 r2 = out2[0];
|
||||
const u32 r3 = out2[1];
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
#include COMPARE_M
|
||||
#endif
|
||||
}
|
||||
197
OpenCL/m27100-pure.cl
Normal file
197
OpenCL/m27100-pure.cl
Normal file
@ -0,0 +1,197 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
// #define NEW_SIMD_CODE
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
#include "inc_vendor.h"
|
||||
#include "inc_types.h"
|
||||
#include "inc_platform.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_md4.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
#endif
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
DECLSPEC u8 hex_convert (const u8 c)
|
||||
{
|
||||
return (c & 15) + (c >> 6) * 9;
|
||||
}
|
||||
|
||||
DECLSPEC u8 hex_to_u8 (const u8 *hex)
|
||||
{
|
||||
u8 v = 0;
|
||||
|
||||
v |= ((u8) hex_convert (hex[1]) << 0);
|
||||
v |= ((u8) hex_convert (hex[0]) << 4);
|
||||
|
||||
return (v);
|
||||
}
|
||||
#endif
|
||||
|
||||
typedef struct netntlm
|
||||
{
|
||||
u32 user_len;
|
||||
u32 domain_len;
|
||||
u32 srvchall_len;
|
||||
u32 clichall_len;
|
||||
|
||||
u32 userdomain_buf[64];
|
||||
u32 chall_buf[256];
|
||||
|
||||
} netntlm_t;
|
||||
|
||||
typedef struct netntlmv2_tmp
|
||||
{
|
||||
u32 digest_buf[4];
|
||||
|
||||
} netntlm_tmp_t;
|
||||
|
||||
|
||||
KERNEL_FQ void m27100_init (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 in[16];
|
||||
|
||||
in[ 0] = pws[gid].i[ 0];
|
||||
in[ 1] = pws[gid].i[ 1];
|
||||
in[ 2] = pws[gid].i[ 2];
|
||||
in[ 3] = pws[gid].i[ 3];
|
||||
in[ 4] = pws[gid].i[ 4];
|
||||
in[ 5] = pws[gid].i[ 5];
|
||||
in[ 6] = pws[gid].i[ 6];
|
||||
in[ 7] = pws[gid].i[ 7];
|
||||
|
||||
u8 *in_ptr = (u8 *) in;
|
||||
|
||||
u32 out[4];
|
||||
|
||||
u8 *out_ptr = (u8 *) out;
|
||||
|
||||
for (int i = 0, j = 0; i < 16; i += 1, j += 2)
|
||||
{
|
||||
out_ptr[i] = hex_to_u8 (in_ptr + j);
|
||||
}
|
||||
|
||||
tmps[gid].digest_buf[0] = out[ 0];
|
||||
tmps[gid].digest_buf[1] = out[ 1];
|
||||
tmps[gid].digest_buf[2] = out[ 2];
|
||||
tmps[gid].digest_buf[3] = out[ 3];
|
||||
|
||||
}
|
||||
|
||||
|
||||
KERNEL_FQ void m27100_loop (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
|
||||
|
||||
}
|
||||
|
||||
KERNEL_FQ void m27100_comp (KERN_ATTR_TMPS_ESALT (netntlm_tmp_t, netntlm_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
w0[0] = tmps[gid].digest_buf[0];
|
||||
w0[1] = tmps[gid].digest_buf[1];
|
||||
w0[2] = tmps[gid].digest_buf[2];
|
||||
w0[3] = tmps[gid].digest_buf[3];
|
||||
w1[0] = 0;
|
||||
w1[1] = 0;
|
||||
w1[2] = 0;
|
||||
w1[3] = 0;
|
||||
w2[0] = 0;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = 0;
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
|
||||
md5_hmac_init_64 (&ctx0, w0, w1, w2, w3);
|
||||
|
||||
md5_hmac_update_global (&ctx0, esalt_bufs[DIGESTS_OFFSET].userdomain_buf, esalt_bufs[DIGESTS_OFFSET].user_len + esalt_bufs[DIGESTS_OFFSET].domain_len);
|
||||
|
||||
md5_hmac_final (&ctx0);
|
||||
|
||||
w0[0] = ctx0.opad.h[0];
|
||||
w0[1] = ctx0.opad.h[1];
|
||||
w0[2] = ctx0.opad.h[2];
|
||||
w0[3] = ctx0.opad.h[3];
|
||||
w1[0] = 0;
|
||||
w1[1] = 0;
|
||||
w1[2] = 0;
|
||||
w1[3] = 0;
|
||||
w2[0] = 0;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = 0;
|
||||
|
||||
md5_hmac_ctx_t ctx;
|
||||
|
||||
md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
|
||||
|
||||
md5_hmac_update_global (&ctx, esalt_bufs[DIGESTS_OFFSET].chall_buf, esalt_bufs[DIGESTS_OFFSET].srvchall_len + esalt_bufs[DIGESTS_OFFSET].clichall_len);
|
||||
|
||||
md5_hmac_final (&ctx);
|
||||
|
||||
tmps[gid].digest_buf[0] = ctx.opad.h[0];
|
||||
tmps[gid].digest_buf[1] = ctx.opad.h[1];
|
||||
tmps[gid].digest_buf[2] = ctx.opad.h[2];
|
||||
tmps[gid].digest_buf[3] = ctx.opad.h[3];
|
||||
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
const u32 r2 = ctx.opad.h[DGST_R2];
|
||||
const u32 r3 = ctx.opad.h[DGST_R3];
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#ifdef KERNEL_STATIC
|
||||
#include COMPARE_M
|
||||
#endif
|
||||
}
|
||||
@ -11,10 +11,13 @@
|
||||
## Bugs
|
||||
##
|
||||
|
||||
- Fixed buffer overflow in Stargazer Stellar Wallet XLM module in hash_encode() if a hash was cracked
|
||||
- Fixed autotune unitialized tmps variable for slow hashes by calling _init kernel before calling _loop kernel
|
||||
- Fixed datatype in function sha384_hmac_init_vector_128() that could come into effect if vector datatype was manually set
|
||||
- Fixed false negative in all VeraCrypt hash-modes if both conditions are met: 1. use CPU for cracking and 2. PIM range was used
|
||||
- Fixed multiple buffer overflow in DPAPI masterkey file v1 module
|
||||
- Fixed multiple buffer overflow in DPAPI masterkey file v1 and v2 module
|
||||
- Fixed out-of-boundary read in input_tokenizer() if the signature in the hash is longer than the length of the plugins' signature constant
|
||||
- Fixed out-of-boundary read in Stuffit5 module in hash_decode()
|
||||
|
||||
##
|
||||
## Improvements
|
||||
@ -42,6 +45,7 @@
|
||||
- Kernel Cache: Add kernel threads into hash computation which is later used in the kernel cache filename
|
||||
- Memory Management: Refactored the code responsible for limiting kernel accel in order to avoid out of -host- memory situations
|
||||
- SCRYPT Kernels: Add more optimized values for some new NV/AMD GPUs
|
||||
- 7-Zip Hook: Increase supported data length from 320kb to 8mb
|
||||
|
||||
##
|
||||
## Algorithms
|
||||
|
||||
@ -5,6 +5,7 @@
|
||||
#define RULE_OP_MANGLE_UREST_LFIRST 'C' // upper case all chars, lower case 1st
|
||||
#define RULE_OP_MANGLE_TREST 't' // switch the case of each char
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT 'T' // switch the case of each char on pos N
|
||||
#define RULE_OP_MANGLE_TOGGLE_AT_SEP '3' // switch the case of the first letter after occurrence N of char X
|
||||
#define RULE_OP_MANGLE_REVERSE 'r' // reverse word
|
||||
#define RULE_OP_MANGLE_DUPEWORD 'd' // append word to itself
|
||||
#define RULE_OP_MANGLE_DUPEWORD_TIMES 'p' // append word to itself N times
|
||||
|
||||
180
hashcat.hctune
180
hashcat.hctune
@ -386,189 +386,13 @@ GeForce_GTX_TITAN 3 9900 2 A
|
||||
DEVICE_TYPE_CPU * 3200 1 N A
|
||||
|
||||
##
|
||||
## SCRYPT
|
||||
## SCRYPT: Tunings for SCRYPT based hash-modes can be found inside the plugin source
|
||||
## See function module_extra_tuningdb_block()
|
||||
##
|
||||
|
||||
DEVICE_TYPE_CPU * 8900 1 N A
|
||||
DEVICE_TYPE_CPU * 9300 1 N A
|
||||
DEVICE_TYPE_CPU * 15700 1 N A
|
||||
DEVICE_TYPE_CPU * 22700 1 N A
|
||||
|
||||
DEVICE_TYPE_GPU * 8900 1 N A
|
||||
DEVICE_TYPE_GPU * 9300 1 N A
|
||||
DEVICE_TYPE_GPU * 15700 1 1 A
|
||||
DEVICE_TYPE_GPU * 22700 1 N A
|
||||
|
||||
##
|
||||
## CryptoAPI
|
||||
##
|
||||
|
||||
DEVICE_TYPE_CPU * 14500 1 A A
|
||||
DEVICE_TYPE_GPU * 14500 1 A A
|
||||
|
||||
## Here's an example of how to manually tune SCRYPT algorithm kernels for your hardware.
|
||||
## Manually tuning the GPU will yield increased performance. There is typically no noticeable change to CPU performance.
|
||||
##
|
||||
## First, you need to know the parameters of your SCRYPT hash: N, r and p.
|
||||
##
|
||||
## The reference SCRYPT parameter values are N=14, r=8 and p=1, but these will likely not match the parameters used by real-world applications.
|
||||
## For reference, the N value represents an exponent (2^N, which we calculate by bit shifting 1 left by N bits).
|
||||
## Hashcat expects this N value in decimal format: 1 << 14 = 16384
|
||||
##
|
||||
## Now that you have the 3 configuration items in decimal format, multiply them by 128 (underlaying crypto primitive block size).
|
||||
## For example: 128 * 16384 * 8 * 1 = 16777216 = 16MB
|
||||
## This is the amount of memory required for the GPU to compute the hash of one password candidate.
|
||||
##
|
||||
## Hashcat computes multiple password candidates in parallel - this is what allows for full utilization of the device.
|
||||
## The number of password candidates that Hashcat can run in parallel is VRAM limited and depends on:
|
||||
##
|
||||
## 1. Compute devices' native compute units
|
||||
## 2. Compute devices' native thread count
|
||||
## 3. Artificial multiplier (--kernel-accel aka -n)
|
||||
##
|
||||
## In order to find these values:
|
||||
##
|
||||
## 1. On startup Hashcat will show: * Device #1: GeForce GTX 980, 3963/4043 MB, 16MCU. The 16 MCU is the number of compute units on that device.
|
||||
## 2. Native thread counts are fixed values: CPU=1, GPU-Intel=8, GPU-AMD=64 (wavefronts), GPU-NVIDIA=32 (warps)
|
||||
##
|
||||
## Now multiply them together. For my GTX980: 16 * 32 * 16777216 = 8589934592 = 8GB
|
||||
##
|
||||
## If we want to actually make use of all computing resources, this GPU would require 8GB of GPU RAM.
|
||||
## However, it doesn't have that:
|
||||
##
|
||||
## Device #1: GeForce GTX 980, 3963/4043 MB, 16MCU. We only have 4043 MB (4GB minus some overhead from the OS).
|
||||
##
|
||||
## How do we deal with this? This is where SCRYPT TMTO(time-memory trde off) kicks in. The SCRYPT algorithm is designed in such a way that we
|
||||
## can pre-compute that 16MB buffer from a self-choosen offset. Details on how this actually works are not important for this process.
|
||||
##
|
||||
## What's relevant to us is that we can halve the buffer size, but we pay with twice the computation time.
|
||||
## We can repeat this as often as we want. That's why it's a trade-off.
|
||||
##
|
||||
## This mechanic can be manually set using --scrypt-tmto on the commandline, but this is not the best way.
|
||||
##
|
||||
## Back to our problem. We need 8GB of memory but have only ~4GB.
|
||||
## It's not a full 4GB. The OS needs some of it and Hashcat needs some of it to store password candidates and other things.
|
||||
## If you run a headless server it should be safe to subtract a fixed value of 200MB from whatever you have in your GPU.
|
||||
##
|
||||
## So lets divide our required memory(8GB) by 2 until it fits in our VRAM - 200MB.
|
||||
##
|
||||
## (8GB >> 0) = 8GB < 3.8GB = No, Does not fit
|
||||
## (8GB >> 1) = 4GB < 3.8GB = No, Does not fit
|
||||
## (8GB >> 2) = 2GB < 3.8GB = Yes!
|
||||
##
|
||||
## This process is automated in Hashcat, but it is important to understand what's happening here.
|
||||
## Because of the light overhead from the OS and Hashcat, we pay a very high price.
|
||||
## Even though it is just 200MB, it forces us to increase the TMTO by another step.
|
||||
## In terms of speed, the speed is now only 1/4 of what we could archieve on that same GPU if it had only 8.2GB ram.
|
||||
## But now we end up in a situation that we waste 1.8GB RAM which costs us ((1.8GB/16MB)>>1) candidates/second.
|
||||
##
|
||||
## This is where manual tuning can come into play.
|
||||
## If we know that the resources we need are close to what we have (in this case 3.8GB <-> 4.0GB)
|
||||
## We could decide to throw away some of our compute units so that we will no longer need 4.0GB but only 3.8GB.
|
||||
## Therefore, we do not need to increase the TMTO by another step to fit in VRAM.
|
||||
##
|
||||
## If we cut down our 16 MCU to only 15 MCU or 14 MCU using --kernel-accel(-n), we end up with:
|
||||
##
|
||||
## 16 * 32 * 16777216 = 8589934592 / 2 = 4294967296 = 4.00GB < 3.80GB = Nope, next
|
||||
## 15 * 32 * 16777216 = 8053063680 / 2 = 4026531840 = 3.84GB < 3.80GB = Nope, next
|
||||
## 14 * 32 * 16777216 = 7516192768 / 2 = 3758096384 = 3.58GB < 3.80GB = Yes!
|
||||
##
|
||||
## So we can throw away 2/16 compute units, but save half of the computation trade-off on the rest of the compute device.
|
||||
## On my GTX980, this improves the performance from 163 H/s to 201 H/s.
|
||||
## You don't need to control --scrypt-tmto manually because now that the multiplier (-n) is smaller than the native value
|
||||
## Hashcat will automatically realize it can decrease the TMTO by one.
|
||||
##
|
||||
## At this point, you found the optimal base value for your compute device. In this case: 14.
|
||||
##
|
||||
## Depending on our hardware, especially hardware with very slow memory access like a GPU
|
||||
## there's a good chance that it's cheaper (faster) to compute an extra step on the GPU register.
|
||||
## So if we increase the TMTO again by one, this gives an extra speed boost.
|
||||
##
|
||||
## On my GTX980, this improves the performance from 201 H/s to 255 H/s.
|
||||
## Again, there's no need to control this with --scrypt-tmto. Hashcat will realize it has to increase the TMTO again.
|
||||
##
|
||||
## All together, you can control all of this by using the -n parameter in the command line.
|
||||
## This is not ideal in a production environment because you must use the --force flag.
|
||||
## The best way to set this is by using this Hashcat.hctune file to store it. This avoids the need to bypass any warnings.
|
||||
##
|
||||
## Find the ideal -n value, then store it here along with the proper compute device name.
|
||||
## Formatting guidelines are availabe at the top of this document.
|
||||
##
|
||||
## -------------------------------------------------
|
||||
##
|
||||
## You can also ignore all theoretical derivations and semi-automate the process in the real scenario (I prefer this approach):
|
||||
##
|
||||
## 1. For example, to find the value for 8900, first create a valid hash for 8900 as follows:
|
||||
##
|
||||
## $ ./hashcat --example-hashes -m 8900 | grep Example.Hash | grep -v Format | cut -b 25- > tmp.hash.8900
|
||||
##
|
||||
## 2. Now let it iterate through all -n values to a certain point. In this case, I'm using 200, but in general it's a value that is at least twice that of the multiprocessor. If you don't mind you can just leave it as it is, it just runs a little longer.
|
||||
##
|
||||
## $ export i=1; while [ $i -ne 201 ]; do echo $i; ./hashcat --quiet tmp.hash.8900 --keep-guessing --self-test-disable --markov-disable --restore-disable --outfile-autohex-disable --wordlist-autohex-disable --potfile-disable --logfile-disable --hwmon-disable --status --status-timer 1 --runtime 28 --machine-readable --optimized-kernel-enable --workload-profile 3 --hash-type 8900 --attack-mode 3 ?b?b?b?b?b?b?b --backend-devices 1 --force -n $i; i=$(($i+1)); done | tee x
|
||||
##
|
||||
## 3. Determine the highest measured H/s speed. But don't just use the highest value. Instead, use the number that seems most stable, usually at the beginning.
|
||||
##
|
||||
## $ grep "$(printf 'STATUS\t3')" x | cut -f4 -d$'\t' | sort -n | tail
|
||||
##
|
||||
## 4. To match the speed you have chosen to the correct value in the "x" file, simply search for it in it. Then go up a little on the block where you found him. The value -n is the single value that begins before the block start. If you have multiple blocks at the same speed, choose the lowest value for -n
|
||||
##
|
||||
|
||||
## 4GB
|
||||
GeForce_GTX_980 * 8900 1 29 A
|
||||
GeForce_GTX_980 * 9300 1 128 A
|
||||
GeForce_GTX_980 * 15700 1 24 A
|
||||
GeForce_GTX_980 * 22700 1 29 A
|
||||
|
||||
## 8GB
|
||||
GeForce_GTX_1080 * 8900 1 15 A
|
||||
GeForce_GTX_1080 * 9300 1 256 A
|
||||
GeForce_GTX_1080 * 15700 1 28 A
|
||||
GeForce_GTX_1080 * 22700 1 15 A
|
||||
|
||||
## 11GB
|
||||
GeForce_RTX_2080_Ti * 8900 1 68 A
|
||||
GeForce_RTX_2080_Ti * 9300 1 528 A
|
||||
GeForce_RTX_2080_Ti * 15700 1 68 A
|
||||
GeForce_RTX_2080_Ti * 22700 1 68 A
|
||||
|
||||
## 8GB
|
||||
GeForce_RTX_3060_Ti * 8900 1 51 A
|
||||
GeForce_RTX_3060_Ti * 9300 1 256 A
|
||||
GeForce_RTX_3060_Ti * 15700 1 11 A
|
||||
GeForce_RTX_3060_Ti * 22700 1 51 A
|
||||
|
||||
## 8GB
|
||||
GeForce_RTX_3070 * 8900 1 46 A
|
||||
GeForce_RTX_3070 * 9300 1 368 A
|
||||
GeForce_RTX_3070 * 15700 1 22 A
|
||||
GeForce_RTX_3070 * 22700 1 46 A
|
||||
|
||||
## 24GB
|
||||
GeForce_RTX_3090 * 8900 1 82 A
|
||||
GeForce_RTX_3090 * 9300 1 984 A
|
||||
GeForce_RTX_3090 * 15700 1 82 A
|
||||
GeForce_RTX_3090 * 22700 1 82 A
|
||||
|
||||
## 4GB
|
||||
ALIAS_AMD_RX480 * 8900 1 15 A
|
||||
ALIAS_AMD_RX480 * 9300 1 232 A
|
||||
ALIAS_AMD_RX480 * 15700 1 58 A
|
||||
ALIAS_AMD_RX480 * 22700 1 15 A
|
||||
|
||||
## 8GB
|
||||
ALIAS_AMD_Vega64 * 8900 1 31 A
|
||||
ALIAS_AMD_Vega64 * 9300 1 440 A
|
||||
ALIAS_AMD_Vega64 * 15700 1 53 A
|
||||
ALIAS_AMD_Vega64 * 22700 1 31 A
|
||||
|
||||
## 32GB
|
||||
ALIAS_AMD_MI100 * 8900 1 79 A
|
||||
ALIAS_AMD_MI100 * 9300 1 1000 A
|
||||
ALIAS_AMD_MI100 * 15700 1 120 A
|
||||
ALIAS_AMD_MI100 * 22700 1 79 A
|
||||
|
||||
## 16GB
|
||||
ALIAS_AMD_RX6900XT * 8900 1 59 A
|
||||
ALIAS_AMD_RX6900XT * 9300 1 720 A
|
||||
ALIAS_AMD_RX6900XT * 15700 1 56 A
|
||||
ALIAS_AMD_RX6900XT * 22700 1 59 A
|
||||
|
||||
@ -69,13 +69,12 @@ int hc_cuFuncSetAttribute (hashcat_ctx_t *hashcat_ctx, CUfunction hfunc,
|
||||
int hc_cuInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags);
|
||||
int hc_cuLaunchKernel (hashcat_ctx_t *hashcat_ctx, CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra);
|
||||
int hc_cuMemAlloc (hashcat_ctx_t *hashcat_ctx, CUdeviceptr *dptr, size_t bytesize);
|
||||
int hc_cuMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount);
|
||||
int hc_cuMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream);
|
||||
int hc_cuMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount);
|
||||
int hc_cuMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, CUdeviceptr srcDevice, size_t ByteCount, CUstream hStream);
|
||||
int hc_cuMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount);
|
||||
int hc_cuMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount, CUstream hStream);
|
||||
int hc_cuMemFree (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dptr);
|
||||
int hc_cuMemsetD32Async (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, unsigned int ui, size_t N, CUstream hStream);
|
||||
int hc_cuMemsetD8Async (hashcat_ctx_t *hashcat_ctx, CUdeviceptr dstDevice, unsigned char uc, size_t N, CUstream hStream);
|
||||
int hc_cuModuleGetFunction (hashcat_ctx_t *hashcat_ctx, CUfunction *hfunc, CUmodule hmod, const char *name);
|
||||
int hc_cuModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues);
|
||||
int hc_cuModuleUnload (hashcat_ctx_t *hashcat_ctx, CUmodule hmod);
|
||||
@ -120,12 +119,11 @@ int hc_hipInit (hashcat_ctx_t *hashcat_ctx, unsigned int Flags
|
||||
int hc_hipLaunchKernel (hashcat_ctx_t *hashcat_ctx, hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, hipStream_t hStream, void **kernelParams, void **extra);
|
||||
int hc_hipMemAlloc (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t bytesize);
|
||||
int hc_hipMemFree (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dptr);
|
||||
int hc_hipMemcpyDtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount);
|
||||
int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemcpyDtoH (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount);
|
||||
int hc_hipMemcpyDtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemcpyHtoD (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount);
|
||||
int hc_hipMemcpyDtoHAsync (hashcat_ctx_t *hashcat_ctx, void *dstHost, hipDeviceptr_t srcDevice, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemcpyHtoDAsync (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, const void *srcHost, size_t ByteCount, hipStream_t hStream);
|
||||
int hc_hipMemsetD32Async (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned int ui, size_t N, hipStream_t hStream);
|
||||
int hc_hipMemsetD8Async (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t dstDevice, unsigned char uc, size_t N, hipStream_t hStream);
|
||||
int hc_hipModuleGetFunction (hashcat_ctx_t *hashcat_ctx, hipFunction_t *hfunc, hipModule_t hmod, const char *name);
|
||||
int hc_hipModuleGetGlobal (hashcat_ctx_t *hashcat_ctx, hipDeviceptr_t *dptr, size_t *bytes, hipModule_t hmod, const char *name);
|
||||
int hc_hipModuleLoadDataEx (hashcat_ctx_t *hashcat_ctx, hipModule_t *module, const void *image, unsigned int numOptions, hipJitOption *options, void **optionValues);
|
||||
@ -143,6 +141,7 @@ int hc_clCreateKernel (hashcat_ctx_t *hashcat_ctx, cl_program program
|
||||
int hc_clCreateProgramWithBinary (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_uint num_devices, const cl_device_id *device_list, const size_t *lengths, const unsigned char **binaries, cl_int *binary_status, cl_program *program);
|
||||
int hc_clCreateProgramWithSource (hashcat_ctx_t *hashcat_ctx, cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_program *program);
|
||||
int hc_clEnqueueCopyBuffer (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_buffer, size_t src_offset, size_t dst_offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
|
||||
int hc_clEnqueueFillBuffer (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, const void *pattern, size_t pattern_size, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
|
||||
int hc_clEnqueueMapBuffer (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_map, cl_map_flags map_flags, size_t offset, size_t size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event, void **buf);
|
||||
int hc_clEnqueueNDRangeKernel (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim, const size_t *global_work_offset, const size_t *global_work_size, const size_t *local_work_size, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
|
||||
int hc_clEnqueueReadBuffer (hashcat_ctx_t *hashcat_ctx, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_read, size_t offset, size_t size, void *ptr, cl_uint num_events_in_wait_list, const cl_event *event_wait_list, cl_event *event);
|
||||
@ -178,17 +177,20 @@ void rebuild_pws_compressed_append (hc_device_param_t *device_param, const u64 p
|
||||
|
||||
int run_cuda_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num);
|
||||
int run_cuda_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 num);
|
||||
int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u32 value, const u64 size);
|
||||
int run_cuda_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u8 value, const u64 size);
|
||||
int run_cuda_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 offset, const u32 value, const u64 size);
|
||||
int run_cuda_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, CUdeviceptr buf, const u64 size);
|
||||
|
||||
int run_hip_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num);
|
||||
int run_hip_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 num);
|
||||
int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u32 value, const u64 size);
|
||||
int run_hip_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u8 value, const u64 size);
|
||||
int run_hip_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 offset, const u32 value, const u64 size);
|
||||
int run_hip_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, hipDeviceptr_t buf, const u64 size);
|
||||
|
||||
int run_opencl_kernel_atinit (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);
|
||||
int run_opencl_kernel_utf8toutf16le (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 num);
|
||||
int run_opencl_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u32 value, const u64 size);
|
||||
int run_opencl_kernel_memset (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 offset, const u8 value, const u64 size);
|
||||
int run_opencl_kernel_memset32 (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 offset, const u32 value, const u64 size);
|
||||
int run_opencl_kernel_bzero (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, cl_mem buf, const u64 size);
|
||||
|
||||
int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kern_run, const u64 pws_pos, const u64 num, const u32 event_update, const u32 iteration);
|
||||
|
||||
@ -63,6 +63,14 @@
|
||||
#define HC_API_CALL
|
||||
#endif
|
||||
|
||||
#if defined (__GNUC__)
|
||||
#define HC_ALIGN(x) __attribute__((aligned(x)))
|
||||
#elif defined (_MSC_VER)
|
||||
#define HC_ALIGN(x) __declspec(align(x))
|
||||
#else
|
||||
#define HC_ALIGN(x)
|
||||
#endif
|
||||
|
||||
#if defined (_WIN)
|
||||
#define WIN32_LEAN_AND_MEAN
|
||||
#endif
|
||||
|
||||
@ -46,6 +46,7 @@ typedef cl_context (CL_API_CALL *OCL_CLCREATECONTEXT) (const cl_
|
||||
typedef cl_kernel (CL_API_CALL *OCL_CLCREATEKERNEL) (cl_program, const char *, cl_int *);
|
||||
typedef cl_program (CL_API_CALL *OCL_CLCREATEPROGRAMWITHBINARY) (cl_context, cl_uint, const cl_device_id *, const size_t *, const unsigned char **, cl_int *, cl_int *);
|
||||
typedef cl_program (CL_API_CALL *OCL_CLCREATEPROGRAMWITHSOURCE) (cl_context, cl_uint, const char **, const size_t *, cl_int *);
|
||||
typedef cl_int (CL_API_CALL *OCL_CLENQUEUEFILLBUFFER) (cl_command_queue, cl_mem, const void *, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *);
|
||||
typedef cl_int (CL_API_CALL *OCL_CLENQUEUECOPYBUFFER) (cl_command_queue, cl_mem, cl_mem, size_t, size_t, size_t, cl_uint, const cl_event *, cl_event *);
|
||||
typedef void * (CL_API_CALL *OCL_CLENQUEUEMAPBUFFER) (cl_command_queue, cl_mem, cl_bool, cl_map_flags, size_t, size_t, cl_uint, const cl_event *, cl_event *, cl_int *);
|
||||
typedef cl_int (CL_API_CALL *OCL_CLENQUEUENDRANGEKERNEL) (cl_command_queue, cl_kernel, cl_uint, const size_t *, const size_t *, const size_t *, cl_uint, const cl_event *, cl_event *);
|
||||
@ -87,6 +88,7 @@ typedef struct hc_opencl_lib
|
||||
OCL_CLCREATEPROGRAMWITHBINARY clCreateProgramWithBinary;
|
||||
OCL_CLCREATEPROGRAMWITHSOURCE clCreateProgramWithSource;
|
||||
OCL_CLENQUEUECOPYBUFFER clEnqueueCopyBuffer;
|
||||
OCL_CLENQUEUEFILLBUFFER clEnqueueFillBuffer;
|
||||
OCL_CLENQUEUEMAPBUFFER clEnqueueMapBuffer;
|
||||
OCL_CLENQUEUENDRANGEKERNEL clEnqueueNDRangeKernel;
|
||||
OCL_CLENQUEUEREADBUFFER clEnqueueReadBuffer;
|
||||
|
||||
@ -1028,17 +1028,14 @@ typedef CUresult (CUDA_API_CALL *CUDA_CUINIT) (unsigned int);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CULAUNCHKERNEL) (CUfunction, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, unsigned int, CUstream, void **, void **);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOC) (CUdeviceptr *, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMALLOCHOST) (void **, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOD) (CUdeviceptr, CUdeviceptr, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTODASYNC) (CUdeviceptr, CUdeviceptr, size_t, CUstream);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOH) (void *, CUdeviceptr, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYDTOHASYNC) (void *, CUdeviceptr, size_t, CUstream);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTOD) (CUdeviceptr, const void *, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMCPYHTODASYNC) (CUdeviceptr, const void *, size_t, CUstream);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREE) (CUdeviceptr);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMFREEHOST) (void *);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMGETINFO) (size_t *, size_t *);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32) (CUdeviceptr, unsigned int, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8) (CUdeviceptr, unsigned char, size_t);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD32ASYNC) (CUdeviceptr, unsigned int, size_t, CUstream);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMEMSETD8ASYNC) (CUdeviceptr, unsigned char, size_t, CUstream);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETFUNCTION) (CUfunction *, CUmodule, const char *);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMODULEGETGLOBAL) (CUdeviceptr *, size_t *, CUmodule, const char *);
|
||||
typedef CUresult (CUDA_API_CALL *CUDA_CUMODULELOAD) (CUmodule *, const char *);
|
||||
@ -1093,17 +1090,14 @@ typedef struct hc_cuda_lib
|
||||
CUDA_CULAUNCHKERNEL cuLaunchKernel;
|
||||
CUDA_CUMEMALLOC cuMemAlloc;
|
||||
CUDA_CUMEMALLOCHOST cuMemAllocHost;
|
||||
CUDA_CUMEMCPYDTOD cuMemcpyDtoD;
|
||||
CUDA_CUMEMCPYDTODASYNC cuMemcpyDtoDAsync;
|
||||
CUDA_CUMEMCPYDTOH cuMemcpyDtoH;
|
||||
CUDA_CUMEMCPYDTOHASYNC cuMemcpyDtoHAsync;
|
||||
CUDA_CUMEMCPYHTOD cuMemcpyHtoD;
|
||||
CUDA_CUMEMCPYHTODASYNC cuMemcpyHtoDAsync;
|
||||
CUDA_CUMEMFREE cuMemFree;
|
||||
CUDA_CUMEMFREEHOST cuMemFreeHost;
|
||||
CUDA_CUMEMGETINFO cuMemGetInfo;
|
||||
CUDA_CUMEMSETD32 cuMemsetD32;
|
||||
CUDA_CUMEMSETD8 cuMemsetD8;
|
||||
CUDA_CUMEMSETD32ASYNC cuMemsetD32Async;
|
||||
CUDA_CUMEMSETD8ASYNC cuMemsetD8Async;
|
||||
CUDA_CUMODULEGETFUNCTION cuModuleGetFunction;
|
||||
CUDA_CUMODULEGETGLOBAL cuModuleGetGlobal;
|
||||
CUDA_CUMODULELOAD cuModuleLoad;
|
||||
|
||||
@ -378,12 +378,11 @@ typedef hipError_t (HIP_API_CALL *HIP_HIPLAUNCHKERNEL) (hipFunction_t,
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMALLOC) (hipDeviceptr_t *, size_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMFREE) (hipDeviceptr_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMGETINFO) (size_t *, size_t *);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOD) (hipDeviceptr_t, hipDeviceptr_t, size_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTODASYNC) (hipDeviceptr_t, hipDeviceptr_t, size_t, hipStream_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOH) (void *, hipDeviceptr_t, size_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYDTOHASYNC) (void *, hipDeviceptr_t, size_t, hipStream_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTOD) (hipDeviceptr_t, const void *, size_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMCPYHTODASYNC) (hipDeviceptr_t, const void *, size_t, hipStream_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD32ASYNC) (hipDeviceptr_t, unsigned int, size_t, hipStream_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMEMSETD8ASYNC) (hipDeviceptr_t, unsigned char, size_t, hipStream_t);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETFUNCTION) (hipFunction_t *, hipModule_t, const char *);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMODULEGETGLOBAL) (hipDeviceptr_t *, size_t *, hipModule_t, const char *);
|
||||
typedef hipError_t (HIP_API_CALL *HIP_HIPMODULELOADDATAEX) (hipModule_t *, const void *, unsigned int, hipJitOption *, void **);
|
||||
@ -421,12 +420,11 @@ typedef struct hc_hip_lib
|
||||
HIP_HIPMEMALLOC hipMemAlloc;
|
||||
HIP_HIPMEMFREE hipMemFree;
|
||||
HIP_HIPMEMGETINFO hipMemGetInfo;
|
||||
HIP_HIPMEMCPYDTOD hipMemcpyDtoD;
|
||||
HIP_HIPMEMCPYDTODASYNC hipMemcpyDtoDAsync;
|
||||
HIP_HIPMEMCPYDTOH hipMemcpyDtoH;
|
||||
HIP_HIPMEMCPYDTOHASYNC hipMemcpyDtoHAsync;
|
||||
HIP_HIPMEMCPYHTOD hipMemcpyHtoD;
|
||||
HIP_HIPMEMCPYHTODASYNC hipMemcpyHtoDAsync;
|
||||
HIP_HIPMEMSETD32ASYNC hipMemsetD32Async;
|
||||
HIP_HIPMEMSETD8ASYNC hipMemsetD8Async;
|
||||
HIP_HIPMODULEGETFUNCTION hipModuleGetFunction;
|
||||
HIP_HIPMODULEGETGLOBAL hipModuleGetGlobal;
|
||||
HIP_HIPMODULELOADDATAEX hipModuleLoadDataEx;
|
||||
|
||||
@ -15,7 +15,7 @@ int hash_encode (const hashconfig_t *hashconfig, const hashes_t *hashes, const m
|
||||
|
||||
int save_hash (hashcat_ctx_t *hashcat_ctx);
|
||||
|
||||
void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain);
|
||||
int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain);
|
||||
|
||||
//int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 salt_pos);
|
||||
int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param);
|
||||
|
||||
@ -13,7 +13,7 @@
|
||||
#include <limits.h>
|
||||
#include <inttypes.h>
|
||||
|
||||
static const int MODULE_INTERFACE_VERSION_MINIMUM = 520;
|
||||
static const int MODULE_INTERFACE_VERSION_MINIMUM = 630;
|
||||
|
||||
static const int MODULE_HASH_MODES_MAXIMUM = 100000;
|
||||
|
||||
|
||||
@ -18,6 +18,7 @@ u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *ha
|
||||
u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
u64 module_esalt_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
const char *module_extra_tuningdb_block (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
u32 module_forced_outfile_format (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
u32 module_hash_category (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra);
|
||||
|
||||
@ -11,9 +11,13 @@
|
||||
|
||||
#define TUNING_DB_FILE "hashcat.hctune"
|
||||
|
||||
int sort_by_tuning_db_alias (const void *v1, const void *v2);
|
||||
int sort_by_tuning_db_entry (const void *v1, const void *v2);
|
||||
|
||||
int tuning_db_init (hashcat_ctx_t *hashcat_ctx);
|
||||
void tuning_db_destroy (hashcat_ctx_t *hashcat_ctx);
|
||||
|
||||
bool tuning_db_process_line (hashcat_ctx_t *hashcat_ctx, const char *line_buf, const int line_num);
|
||||
tuning_db_entry_t *tuning_db_search (hashcat_ctx_t *hashcat_ctx, const char *device_name, const cl_device_type device_type, int attack_mode, const int hash_mode);
|
||||
|
||||
#endif // _TUNINGDB_H
|
||||
|
||||
@ -295,6 +295,7 @@ typedef enum rule_functions
|
||||
RULE_OP_MANGLE_UREST_LFIRST = 'C',
|
||||
RULE_OP_MANGLE_TREST = 't',
|
||||
RULE_OP_MANGLE_TOGGLE_AT = 'T',
|
||||
RULE_OP_MANGLE_TOGGLE_AT_SEP = '3',
|
||||
RULE_OP_MANGLE_REVERSE = 'r',
|
||||
RULE_OP_MANGLE_DUPEWORD = 'd',
|
||||
RULE_OP_MANGLE_DUPEWORD_TIMES = 'p',
|
||||
@ -926,7 +927,6 @@ typedef struct hashes
|
||||
|
||||
void *digests_buf;
|
||||
u32 *digests_shown;
|
||||
u32 *digests_shown_tmp;
|
||||
|
||||
u32 salts_cnt;
|
||||
u32 salts_done;
|
||||
@ -1427,6 +1427,7 @@ typedef struct hc_device_param
|
||||
|
||||
CUevent cuda_event1;
|
||||
CUevent cuda_event2;
|
||||
CUevent cuda_event3;
|
||||
|
||||
CUmodule cuda_module;
|
||||
CUmodule cuda_module_shared;
|
||||
@ -1508,6 +1509,7 @@ typedef struct hc_device_param
|
||||
|
||||
hipEvent_t hip_event1;
|
||||
hipEvent_t hip_event2;
|
||||
hipEvent_t hip_event3;
|
||||
|
||||
hipModule_t hip_module;
|
||||
hipModule_t hip_module_shared;
|
||||
@ -2042,9 +2044,11 @@ typedef struct tuning_db
|
||||
|
||||
tuning_db_alias_t *alias_buf;
|
||||
int alias_cnt;
|
||||
int alias_alloc;
|
||||
|
||||
tuning_db_entry_t *entry_buf;
|
||||
int entry_cnt;
|
||||
int entry_alloc;
|
||||
|
||||
} tuning_db_t;
|
||||
|
||||
@ -2648,6 +2652,7 @@ typedef struct module_ctx
|
||||
u32 (*module_dgst_pos3) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
u32 (*module_dgst_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
u64 (*module_esalt_size) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
const char *(*module_extra_tuningdb_block) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
u32 (*module_forced_outfile_format) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
u32 (*module_hash_category) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
const char *(*module_hash_name) (const hashconfig_t *, const user_options_t *, const user_options_extra_t *);
|
||||
|
||||
@ -53,7 +53,7 @@ endif
|
||||
## Do not modify
|
||||
##
|
||||
|
||||
MODULE_INTERFACE_VERSION := 620
|
||||
MODULE_INTERFACE_VERSION := 630
|
||||
|
||||
##
|
||||
## Native compiler paths
|
||||
@ -203,12 +203,11 @@ endif
|
||||
## because UNRAR
|
||||
ifeq ($(ENABLE_UNRAR),1)
|
||||
ifeq ($(USE_SYSTEM_UNRAR),0)
|
||||
ifneq ($(UNAME),Darwin)
|
||||
CFLAGS_UNRAR += -Wno-misleading-indentation
|
||||
ifneq ($(CC),clang)
|
||||
CFLAGS_UNRAR += -Wno-class-memaccess
|
||||
else
|
||||
CFLAGS_UNRAR += -Wno-missing-braces
|
||||
CFLAGS_UNRAR += -Wno-misleading-indentation
|
||||
endif
|
||||
CFLAGS_UNRAR += -Wno-missing-braces
|
||||
CFLAGS_UNRAR += -Wno-unused-variable
|
||||
CFLAGS_UNRAR += -Wno-unused-parameter
|
||||
CFLAGS_UNRAR += -Wno-unused-function
|
||||
|
||||
358
src/autotune.c
358
src/autotune.c
@ -10,7 +10,7 @@
|
||||
#include "status.h"
|
||||
#include "autotune.h"
|
||||
|
||||
static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops)
|
||||
static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops, const u32 kernel_threads)
|
||||
{
|
||||
hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
|
||||
user_options_t *user_options = hashcat_ctx->user_options;
|
||||
@ -19,7 +19,9 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
|
||||
device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
|
||||
device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
|
||||
|
||||
u32 kernel_power_try = device_param->hardware_power * kernel_accel;
|
||||
const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * kernel_threads;
|
||||
|
||||
u32 kernel_power_try = hardware_power * kernel_accel;
|
||||
|
||||
if (user_options->attack_mode == ATTACK_MODE_ASSOCIATION)
|
||||
{
|
||||
@ -33,6 +35,10 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
|
||||
}
|
||||
}
|
||||
|
||||
const u32 kernel_threads_sav = device_param->kernel_threads;
|
||||
|
||||
device_param->kernel_threads = kernel_threads;
|
||||
|
||||
const double spin_damp_sav = device_param->spin_damp;
|
||||
|
||||
device_param->spin_damp = 0;
|
||||
@ -50,71 +56,51 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par
|
||||
}
|
||||
else
|
||||
{
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0);
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE)
|
||||
{
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_try, true, 0);
|
||||
}
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0);
|
||||
}
|
||||
|
||||
device_param->spin_damp = spin_damp_sav;
|
||||
|
||||
const double exec_msec_prev = get_avg_exec_time (device_param, 1);
|
||||
|
||||
return exec_msec_prev;
|
||||
}
|
||||
|
||||
/*
|
||||
static double try_run_preferred (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops)
|
||||
{
|
||||
hashconfig_t *hashconfig = hashcat_ctx->hashconfig;
|
||||
|
||||
device_param->kernel_params_buf32[28] = 0;
|
||||
device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set
|
||||
device_param->kernel_params_buf32[30] = kernel_loops; // because there's two variables for inner iters for slow and fast hashes
|
||||
|
||||
const u32 kernel_power_try = device_param->hardware_power * kernel_accel;
|
||||
|
||||
const u32 kernel_threads_sav = device_param->kernel_threads;
|
||||
|
||||
const double spin_damp_sav = device_param->spin_damp;
|
||||
|
||||
device_param->spin_damp = 0;
|
||||
|
||||
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
|
||||
{
|
||||
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
|
||||
{
|
||||
device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple1;
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_try, true, 0);
|
||||
}
|
||||
else
|
||||
{
|
||||
device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple4;
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 0, kernel_power_try, true, 0);
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
device_param->kernel_threads = device_param->kernel_preferred_wgs_multiple2;
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 0, kernel_power_try, true, 0);
|
||||
}
|
||||
|
||||
device_param->kernel_threads = kernel_threads_sav;
|
||||
|
||||
device_param->spin_damp = spin_damp_sav;
|
||||
|
||||
const double exec_msec_prev = get_avg_exec_time (device_param, 1);
|
||||
|
||||
return exec_msec_prev;
|
||||
}
|
||||
*/
|
||||
|
||||
static double try_run_times (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops, const u32 kernel_threads, const int times)
|
||||
{
|
||||
double exec_msec_best = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
|
||||
for (int i = 1; i < times; i++)
|
||||
{
|
||||
double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
|
||||
if (exec_msec > exec_msec_best) continue;
|
||||
|
||||
exec_msec_best = exec_msec;
|
||||
}
|
||||
|
||||
return exec_msec_best;
|
||||
}
|
||||
|
||||
static u32 previous_power_of_two (const u32 x)
|
||||
{
|
||||
// https://stackoverflow.com/questions/2679815/previous-power-of-2
|
||||
// really cool!
|
||||
|
||||
if (x == 0) return 0;
|
||||
|
||||
u32 r = x;
|
||||
|
||||
r |= (r >> 1);
|
||||
r |= (r >> 2);
|
||||
r |= (r >> 4);
|
||||
r |= (r >> 8);
|
||||
r |= (r >> 16);
|
||||
|
||||
return r - (r >> 1);
|
||||
}
|
||||
|
||||
static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
|
||||
{
|
||||
@ -131,9 +117,57 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
const u32 kernel_loops_min = device_param->kernel_loops_min;
|
||||
const u32 kernel_loops_max = device_param->kernel_loops_max;
|
||||
|
||||
const u32 kernel_threads_min = device_param->kernel_threads_min;
|
||||
const u32 kernel_threads_max = device_param->kernel_threads_max;
|
||||
|
||||
u32 kernel_accel = kernel_accel_min;
|
||||
u32 kernel_loops = kernel_loops_min;
|
||||
|
||||
// for the threads we take as initial value what we receive from the runtime
|
||||
// but is only to start with something, we will fine tune this value as soon as we have our workload specified
|
||||
// this thread limiting is also performed insinde run_kernel() so we need to redo it here, too
|
||||
|
||||
u32 kernel_wgs = 0;
|
||||
u32 kernel_wgs_multiple = 0;
|
||||
|
||||
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
|
||||
{
|
||||
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
|
||||
{
|
||||
kernel_wgs = device_param->kernel_wgs1;
|
||||
|
||||
kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple1;
|
||||
}
|
||||
else
|
||||
{
|
||||
kernel_wgs = device_param->kernel_wgs4;
|
||||
|
||||
kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple4;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
kernel_wgs = device_param->kernel_wgs2;
|
||||
|
||||
kernel_wgs_multiple = device_param->kernel_preferred_wgs_multiple2;
|
||||
}
|
||||
|
||||
u32 kernel_threads = kernel_threads_max;
|
||||
|
||||
if ((kernel_wgs >= kernel_threads_min) && (kernel_wgs <= kernel_threads_max))
|
||||
{
|
||||
kernel_threads = kernel_wgs;
|
||||
}
|
||||
|
||||
// having a value power of 2 makes it easier to divide
|
||||
|
||||
const u32 kernel_threads_p2 = previous_power_of_two (kernel_threads);
|
||||
|
||||
if ((kernel_threads_p2 >= kernel_threads_min) && (kernel_threads_p2 <= kernel_threads_max))
|
||||
{
|
||||
kernel_threads = kernel_threads_p2;
|
||||
}
|
||||
|
||||
// in this case the user specified a fixed -n and -u on the commandline
|
||||
// no way to tune anything
|
||||
// but we need to run a few caching rounds
|
||||
@ -149,10 +183,10 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
|
||||
if (hashconfig->warmup_disable == false)
|
||||
{
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads);
|
||||
}
|
||||
|
||||
#endif
|
||||
@ -164,29 +198,19 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
|
||||
const u32 kernel_power_max = device_param->hardware_power * kernel_accel_max;
|
||||
|
||||
int CU_rc;
|
||||
int HIP_rc;
|
||||
int CL_rc;
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
CU_rc = run_cuda_kernel_atinit (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, kernel_power_max);
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
if (run_cuda_kernel_atinit (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, kernel_power_max) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
HIP_rc = run_hip_kernel_atinit (hashcat_ctx, device_param, device_param->hip_d_pws_buf, kernel_power_max);
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
if (run_hip_kernel_atinit (hashcat_ctx, device_param, device_param->hip_d_pws_buf, kernel_power_max) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
CL_rc = run_opencl_kernel_atinit (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, kernel_power_max);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
if (run_opencl_kernel_atinit (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, kernel_power_max) == -1) return -1;
|
||||
}
|
||||
|
||||
if (user_options->slow_candidates == true)
|
||||
@ -200,35 +224,53 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
{
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
CU_rc = hc_cuMemcpyDtoD (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t));
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
if (hc_cuMemcpyDtoDAsync (hashcat_ctx, device_param->cuda_d_rules_c, device_param->cuda_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->cuda_stream) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
HIP_rc = hc_hipMemcpyDtoD (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t));
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
if (hc_hipMemcpyDtoDAsync (hashcat_ctx, device_param->hip_d_rules_c, device_param->hip_d_rules, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), device_param->hip_stream) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
CL_rc = hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
if (hc_clEnqueueCopyBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_rules, device_param->opencl_d_rules_c, 0, 0, MIN (kernel_loops_max, KERNEL_RULES) * sizeof (kernel_rule_t), 0, NULL, NULL) == -1) return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// we also need to initialize some values using kernels
|
||||
|
||||
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
|
||||
{
|
||||
// nothing to do
|
||||
}
|
||||
else
|
||||
{
|
||||
const u32 kernel_threads_sav = device_param->kernel_threads;
|
||||
|
||||
device_param->kernel_threads = device_param->kernel_wgs1;
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 0, kernel_power_max, false, 0);
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_LOOP_PREPARE)
|
||||
{
|
||||
device_param->kernel_threads = device_param->kernel_wgs2p;
|
||||
|
||||
run_kernel (hashcat_ctx, device_param, KERN_RUN_2P, 0, kernel_power_max, false, 0);
|
||||
}
|
||||
|
||||
device_param->kernel_threads = kernel_threads_sav;
|
||||
}
|
||||
|
||||
// Do a pre-autotune test run to find out if kernel runtime is above some TDR limit
|
||||
|
||||
u32 kernel_loops_max_reduced = kernel_loops_max;
|
||||
|
||||
if (true)
|
||||
{
|
||||
double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min);
|
||||
double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min, kernel_threads);
|
||||
|
||||
if (exec_msec > 2000)
|
||||
{
|
||||
@ -237,7 +279,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
return -1;
|
||||
}
|
||||
|
||||
exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min);
|
||||
exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops_min, kernel_threads);
|
||||
|
||||
const u32 mm = kernel_loops_max / kernel_loops_min;
|
||||
|
||||
@ -257,16 +299,16 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
{
|
||||
if (kernel_loops > kernel_loops_max_reduced) continue;
|
||||
|
||||
double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_min, kernel_loops);
|
||||
double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_min, kernel_loops, kernel_threads, 1);
|
||||
|
||||
if (exec_msec < target_msec) break;
|
||||
}
|
||||
}
|
||||
|
||||
// now the same for kernel-accel but with the new kernel-loops from previous loop set
|
||||
|
||||
#define STEPS_CNT 16
|
||||
|
||||
// now the same for kernel-accel but with the new kernel-loops from previous loop set
|
||||
|
||||
if (kernel_accel_min < kernel_accel_max)
|
||||
{
|
||||
for (int i = 0; i < STEPS_CNT; i++)
|
||||
@ -276,7 +318,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
if (kernel_accel_try < kernel_accel_min) continue;
|
||||
if (kernel_accel_try > kernel_accel_max) break;
|
||||
|
||||
double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_try, kernel_loops);
|
||||
double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads, 1);
|
||||
|
||||
if (exec_msec > target_msec) break;
|
||||
|
||||
@ -292,7 +334,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
const u32 kernel_accel_orig = kernel_accel;
|
||||
const u32 kernel_loops_orig = kernel_loops;
|
||||
|
||||
double exec_msec_prev = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
double exec_msec_prev = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 1);
|
||||
|
||||
for (int i = 1; i < STEPS_CNT; i++)
|
||||
{
|
||||
@ -307,7 +349,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
|
||||
// do a real test
|
||||
|
||||
const double exec_msec = try_run (hashcat_ctx, device_param, kernel_accel_try, kernel_loops_try);
|
||||
const double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops_try, kernel_threads, 1);
|
||||
|
||||
if (exec_msec_prev < exec_msec) break;
|
||||
|
||||
@ -324,7 +366,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
}
|
||||
}
|
||||
|
||||
double exec_msec_pre_final = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
double exec_msec_pre_final = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 1);
|
||||
|
||||
const u32 exec_left = (const u32) (target_msec / exec_msec_pre_final);
|
||||
|
||||
@ -339,46 +381,43 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
kernel_accel *= exec_accel_min;
|
||||
}
|
||||
|
||||
// start finding best thread count is easier.
|
||||
// it's either the preferred or the maximum thread count
|
||||
// v6.2.4 new section: find thread count
|
||||
// This is not as effective as it could be because of inaccurate kernel return timers
|
||||
// But is better than fixed values
|
||||
// Timers in this section are critical, so we rerun meassurements 3 times
|
||||
|
||||
/*
|
||||
const u32 kernel_threads_min = device_param->kernel_threads_min;
|
||||
const u32 kernel_threads_max = device_param->kernel_threads_max;
|
||||
|
||||
if (kernel_threads_min < kernel_threads_max)
|
||||
if (kernel_threads_max > kernel_threads_min)
|
||||
{
|
||||
const double exec_msec_max = try_run (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
const u32 kernel_accel_orig = kernel_accel;
|
||||
const u32 kernel_threads_orig = kernel_threads;
|
||||
|
||||
u32 preferred_threads = 0;
|
||||
double exec_msec_prev = try_run_times (hashcat_ctx, device_param, kernel_accel, kernel_loops, kernel_threads, 3);
|
||||
|
||||
if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
|
||||
for (int i = 1; i < STEPS_CNT; i++)
|
||||
{
|
||||
if (hashconfig->opti_type & OPTI_TYPE_OPTIMIZED_KERNEL)
|
||||
{
|
||||
preferred_threads = device_param->kernel_preferred_wgs_multiple1;
|
||||
}
|
||||
else
|
||||
{
|
||||
preferred_threads = device_param->kernel_preferred_wgs_multiple4;
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
preferred_threads = device_param->kernel_preferred_wgs_multiple2;
|
||||
}
|
||||
const u32 kernel_accel_try = kernel_accel_orig * (1U << i);
|
||||
const u32 kernel_threads_try = kernel_threads_orig / (1U << i);
|
||||
|
||||
if ((preferred_threads >= kernel_threads_min) && (preferred_threads <= kernel_threads_max))
|
||||
{
|
||||
const double exec_msec_preferred = try_run_preferred (hashcat_ctx, device_param, kernel_accel, kernel_loops);
|
||||
// since we do not modify total amount of workitems, we can (and need) to do increase kernel_accel_max
|
||||
|
||||
if (exec_msec_preferred < exec_msec_max)
|
||||
{
|
||||
device_param->kernel_threads = preferred_threads;
|
||||
}
|
||||
const u32 kernel_accel_max_try = kernel_accel_max * (1U << i);
|
||||
|
||||
if (kernel_accel_try > kernel_accel_max_try) break;
|
||||
|
||||
if (kernel_threads_try < kernel_threads_min) break;
|
||||
|
||||
if (kernel_threads_try % kernel_wgs_multiple) break; // this would just be waste of time
|
||||
|
||||
double exec_msec = try_run_times (hashcat_ctx, device_param, kernel_accel_try, kernel_loops, kernel_threads_try, 3);
|
||||
|
||||
if (exec_msec > exec_msec_prev) continue;
|
||||
|
||||
exec_msec_prev = exec_msec;
|
||||
|
||||
kernel_accel = kernel_accel_try;
|
||||
kernel_threads = kernel_threads_try;
|
||||
}
|
||||
}
|
||||
*/
|
||||
}
|
||||
|
||||
// reset them fake words
|
||||
@ -386,77 +425,43 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
int CU_rc;
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, device_param->size_pws) == -1) return -1;
|
||||
|
||||
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_pws_buf, device_param->size_pws);
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, device_param->size_plains) == -1) return -1;
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, device_param->size_shown) == -1) return -1;
|
||||
|
||||
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_plain_bufs, device_param->size_plains);
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, device_param->size_results) == -1) return -1;
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
|
||||
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown, device_param->size_shown);
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
|
||||
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, device_param->size_results);
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
|
||||
CU_rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tmps, device_param->size_tmps);
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_tmps, device_param->size_tmps) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
int HIP_rc;
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws) == -1) return -1;
|
||||
|
||||
HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_pws_buf, device_param->size_pws);
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains) == -1) return -1;
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown) == -1) return -1;
|
||||
|
||||
HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_plain_bufs, device_param->size_plains);
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results) == -1) return -1;
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
|
||||
HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown, device_param->size_shown);
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
|
||||
HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, device_param->size_results);
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
|
||||
HIP_rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps);
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_tmps, device_param->size_tmps) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
int CL_rc;
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, device_param->size_pws) == -1) return -1;
|
||||
|
||||
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_pws_buf, device_param->size_pws);
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_plain_bufs, device_param->size_plains) == -1) return -1;
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, device_param->size_shown) == -1) return -1;
|
||||
|
||||
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_plain_bufs, device_param->size_plains);
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, device_param->size_results) == -1) return -1;
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tmps, device_param->size_tmps) == -1) return -1;
|
||||
|
||||
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, device_param->size_shown);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
|
||||
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, device_param->size_results);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
|
||||
CL_rc = run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_tmps, device_param->size_tmps);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1;
|
||||
}
|
||||
|
||||
// reset timer
|
||||
@ -478,8 +483,13 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param
|
||||
|
||||
// store
|
||||
|
||||
device_param->kernel_accel = kernel_accel;
|
||||
device_param->kernel_loops = kernel_loops;
|
||||
device_param->kernel_accel = kernel_accel;
|
||||
device_param->kernel_loops = kernel_loops;
|
||||
device_param->kernel_threads = kernel_threads;
|
||||
|
||||
const u32 hardware_power = ((hashconfig->opts_type & OPTS_TYPE_MP_MULTI_DISABLE) ? 1 : device_param->device_processors) * device_param->kernel_threads;
|
||||
|
||||
device_param->hardware_power = hardware_power;
|
||||
|
||||
const u32 kernel_power = device_param->hardware_power * device_param->kernel_accel;
|
||||
|
||||
|
||||
1086
src/backend.c
1086
src/backend.c
File diff suppressed because it is too large
Load Diff
382
src/hashes.c
382
src/hashes.c
@ -300,7 +300,7 @@ int save_hash (hashcat_ctx_t *hashcat_ctx)
|
||||
return 0;
|
||||
}
|
||||
|
||||
void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain)
|
||||
int check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, plain_t *plain)
|
||||
{
|
||||
const debugfile_ctx_t *debugfile_ctx = hashcat_ctx->debugfile_ctx;
|
||||
const hashes_t *hashes = hashcat_ctx->hashes;
|
||||
@ -313,23 +313,62 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
|
||||
void *tmps = NULL;
|
||||
|
||||
cl_event opencl_event;
|
||||
int rc;
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
|
||||
{
|
||||
tmps = hcmalloc (hashconfig->tmp_size);
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
hc_cuMemcpyDtoH (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size);
|
||||
rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->cuda_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->cuda_stream);
|
||||
|
||||
if (rc == 0)
|
||||
{
|
||||
rc = hc_cuEventRecord (hashcat_ctx, device_param->cuda_event3, device_param->cuda_stream);
|
||||
}
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
hcfree (tmps);
|
||||
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
hc_hipMemcpyDtoH (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size);
|
||||
rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, tmps, device_param->hip_d_tmps + (plain->gidvid * hashconfig->tmp_size), hashconfig->tmp_size, device_param->hip_stream);
|
||||
|
||||
if (rc == 0)
|
||||
{
|
||||
rc = hc_hipEventRecord (hashcat_ctx, device_param->hip_event3, device_param->hip_stream);
|
||||
}
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
hcfree (tmps);
|
||||
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_TRUE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, NULL);
|
||||
rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_tmps, CL_FALSE, plain->gidvid * hashconfig->tmp_size, hashconfig->tmp_size, tmps, 0, NULL, &opencl_event);
|
||||
|
||||
if (rc == 0)
|
||||
{
|
||||
rc = hc_clFlush (hashcat_ctx, device_param->opencl_command_queue);
|
||||
}
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
hcfree (tmps);
|
||||
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -337,15 +376,14 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
|
||||
u8 *out_buf = hashes->out_buf;
|
||||
|
||||
int out_len = hash_encode (hashcat_ctx->hashconfig, hashcat_ctx->hashes, hashcat_ctx->module_ctx, (char *) out_buf, HCBUFSIZ_LARGE, salt_pos, digest_pos);
|
||||
int out_len = hash_encode (hashconfig, hashes, module_ctx, (char *) out_buf, HCBUFSIZ_LARGE, salt_pos, digest_pos);
|
||||
|
||||
out_buf[out_len] = 0;
|
||||
|
||||
// plain
|
||||
|
||||
u8 plain_buf[0x1000]; // while the password itself can have only length 256, the module could encode it with something like base64 which inflates the requires buffer size
|
||||
|
||||
memset (plain_buf, 0, sizeof (plain_buf));
|
||||
u8 plain_buf[HCBUFSIZ_TINY] = { 0 }; // while the password itself can have only length 256, the module could encode it with something like base64 which inflates the requires buffer size
|
||||
u8 postprocess_buf[HCBUFSIZ_TINY] = { 0 };
|
||||
|
||||
u8 *plain_ptr = plain_buf;
|
||||
|
||||
@ -355,18 +393,27 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
|
||||
if (module_ctx->module_build_plain_postprocess != MODULE_DEFAULT)
|
||||
{
|
||||
u8 temp_buf[0x1000];
|
||||
|
||||
memset (temp_buf, 0, sizeof (temp_buf));
|
||||
|
||||
const int temp_len = module_ctx->module_build_plain_postprocess (hashcat_ctx->hashconfig, hashcat_ctx->hashes, tmps, (u32 *) plain_buf, sizeof (plain_buf), plain_len, (u32 *)temp_buf, sizeof (temp_buf));
|
||||
|
||||
if (temp_len < (int) sizeof (plain_buf))
|
||||
if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
|
||||
{
|
||||
memcpy (plain_buf, temp_buf, temp_len);
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
if (hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event3) == -1) return -1;
|
||||
}
|
||||
|
||||
plain_len = temp_len;
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
if (hc_hipEventSynchronize (hashcat_ctx, device_param->hip_event3) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1;
|
||||
}
|
||||
}
|
||||
|
||||
plain_len = module_ctx->module_build_plain_postprocess (hashconfig, hashes, tmps, (u32 *) plain_buf, sizeof (plain_buf), plain_len, (u32 *) postprocess_buf, sizeof (postprocess_buf));
|
||||
|
||||
plain_ptr = postprocess_buf;
|
||||
}
|
||||
|
||||
// crackpos
|
||||
@ -407,6 +454,24 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
|
||||
if (module_ctx->module_hash_encode_potfile != MODULE_DEFAULT)
|
||||
{
|
||||
if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
|
||||
{
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
if (hc_cuEventSynchronize (hashcat_ctx, device_param->cuda_event3) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
if (hc_hipEventSynchronize (hashcat_ctx, device_param->hip_event3) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
if (hc_clWaitForEvents (hashcat_ctx, 1, &opencl_event) == -1) return -1;
|
||||
}
|
||||
}
|
||||
|
||||
salt_t *salts_buf = hashes->salts_buf;
|
||||
|
||||
salts_buf += salt_pos;
|
||||
@ -471,7 +536,14 @@ void check_hash (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
if (hashconfig->opts_type & OPTS_TYPE_COPY_TMPS)
|
||||
{
|
||||
hcfree (tmps);
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
if (hc_clReleaseEvent (hashcat_ctx, opencl_event) == -1) return -1;
|
||||
}
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
//int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 salt_pos)
|
||||
@ -484,173 +556,204 @@ int check_cracked (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param)
|
||||
user_options_t *user_options = hashcat_ctx->user_options;
|
||||
|
||||
u32 num_cracked = 0;
|
||||
|
||||
int CU_rc;
|
||||
int HIP_rc;
|
||||
int CL_rc;
|
||||
int rc;
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32));
|
||||
if (hc_cuMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->cuda_d_result, sizeof (u32), device_param->cuda_stream) == -1) return -1;
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
if (hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32));
|
||||
if (hc_hipMemcpyDtoHAsync (hashcat_ctx, &num_cracked, device_param->hip_d_result, sizeof (u32), device_param->hip_stream) == -1) return -1;
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
if (hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream) == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
/* blocking */
|
||||
if (hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL) == -1) return -1;
|
||||
}
|
||||
|
||||
if (user_options->speed_only == true)
|
||||
if (num_cracked == 0 || user_options->speed_only == true)
|
||||
{
|
||||
// we want the hc_clEnqueueReadBuffer to run in benchmark mode because it has an influence in performance
|
||||
// we want to get the num_cracked in benchmark mode because it has an influence in performance
|
||||
// however if the benchmark cracks the artificial hash used for benchmarks we don't want to see that!
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
if (num_cracked)
|
||||
plain_t *cracked = (plain_t *) hcmalloc (num_cracked * sizeof (plain_t));
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
plain_t *cracked = (plain_t *) hccalloc (num_cracked, sizeof (plain_t));
|
||||
rc = hc_cuMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->cuda_stream);
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
if (rc == 0)
|
||||
{
|
||||
CU_rc = hc_cuMemcpyDtoH (hashcat_ctx, cracked, device_param->cuda_d_plain_bufs, num_cracked * sizeof (plain_t));
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
rc = hc_cuStreamSynchronize (hashcat_ctx, device_param->cuda_stream);
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
if (rc == -1)
|
||||
{
|
||||
HIP_rc = hc_hipMemcpyDtoH (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t));
|
||||
hcfree (cracked);
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
rc = hc_hipMemcpyDtoHAsync (hashcat_ctx, cracked, device_param->hip_d_plain_bufs, num_cracked * sizeof (plain_t), device_param->hip_stream);
|
||||
|
||||
if (rc == 0)
|
||||
{
|
||||
rc = hc_hipStreamSynchronize (hashcat_ctx, device_param->hip_stream);
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
if (rc == -1)
|
||||
{
|
||||
CL_rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
|
||||
hcfree (cracked);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
/* blocking */
|
||||
rc = hc_clEnqueueReadBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
hcfree (cracked);
|
||||
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
u32 cpt_cracked = 0;
|
||||
|
||||
hc_thread_mutex_lock (status_ctx->mux_display);
|
||||
|
||||
for (u32 i = 0; i < num_cracked; i++)
|
||||
{
|
||||
const u32 hash_pos = cracked[i].hash_pos;
|
||||
|
||||
if (hashes->digests_shown[hash_pos] == 1) continue;
|
||||
|
||||
const u32 salt_pos = cracked[i].salt_pos;
|
||||
salt_t *salt_buf = &hashes->salts_buf[salt_pos];
|
||||
|
||||
if ((hashconfig->opts_type & OPTS_TYPE_PT_NEVERCRACK) == 0)
|
||||
{
|
||||
hashes->digests_shown[hash_pos] = 1;
|
||||
|
||||
hashes->digests_done++;
|
||||
|
||||
cpt_cracked++;
|
||||
|
||||
salt_buf->digests_done++;
|
||||
|
||||
if (salt_buf->digests_done == salt_buf->digests_cnt)
|
||||
{
|
||||
hashes->salts_shown[salt_pos] = 1;
|
||||
|
||||
hashes->salts_done++;
|
||||
}
|
||||
}
|
||||
|
||||
u32 cpt_cracked = 0;
|
||||
if (hashes->salts_done == hashes->salts_cnt) mycracked (hashcat_ctx);
|
||||
|
||||
rc = check_hash (hashcat_ctx, device_param, &cracked[i]);
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
break;
|
||||
}
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_PT_NEVERCRACK)
|
||||
{
|
||||
// we need to reset cracked state on the device
|
||||
// otherwise host thinks again and again the hash was cracked
|
||||
// and returns invalid password each time
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
rc = run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
rc = run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
/* NOTE: run_opencl_kernel_bzero() does not handle buffer offset */
|
||||
rc = run_opencl_kernel_memset32 (hashcat_ctx, device_param, device_param->opencl_d_digests_shown, salt_buf->digests_offset * sizeof (u32), 0, salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
hc_thread_mutex_unlock (status_ctx->mux_display);
|
||||
|
||||
hcfree (cracked);
|
||||
|
||||
if (rc == -1)
|
||||
{
|
||||
return -1;
|
||||
}
|
||||
|
||||
if (cpt_cracked > 0)
|
||||
{
|
||||
hc_thread_mutex_lock (status_ctx->mux_display);
|
||||
|
||||
for (u32 i = 0; i < num_cracked; i++)
|
||||
{
|
||||
const u32 hash_pos = cracked[i].hash_pos;
|
||||
cpt_ctx->cpt_buf[cpt_ctx->cpt_pos].timestamp = time (NULL);
|
||||
cpt_ctx->cpt_buf[cpt_ctx->cpt_pos].cracked = cpt_cracked;
|
||||
|
||||
if (hashes->digests_shown[hash_pos] == 1) continue;
|
||||
cpt_ctx->cpt_pos++;
|
||||
|
||||
const u32 salt_pos = cracked[i].salt_pos;
|
||||
salt_t *salt_buf = &hashes->salts_buf[salt_pos];
|
||||
cpt_ctx->cpt_total += cpt_cracked;
|
||||
|
||||
if ((hashconfig->opts_type & OPTS_TYPE_PT_NEVERCRACK) == 0)
|
||||
{
|
||||
hashes->digests_shown[hash_pos] = 1;
|
||||
|
||||
hashes->digests_done++;
|
||||
|
||||
cpt_cracked++;
|
||||
|
||||
salt_buf->digests_done++;
|
||||
|
||||
if (salt_buf->digests_done == salt_buf->digests_cnt)
|
||||
{
|
||||
hashes->salts_shown[salt_pos] = 1;
|
||||
|
||||
hashes->salts_done++;
|
||||
}
|
||||
}
|
||||
|
||||
if (hashes->salts_done == hashes->salts_cnt) mycracked (hashcat_ctx);
|
||||
|
||||
check_hash (hashcat_ctx, device_param, &cracked[i]);
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_PT_NEVERCRACK)
|
||||
{
|
||||
// we need to reset cracked state on the device
|
||||
// otherwise host thinks again and again the hash was cracked
|
||||
// and returns invalid password each time
|
||||
|
||||
memset (hashes->digests_shown_tmp, 0, salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_digests_shown + (salt_buf->digests_offset * sizeof (u32)), &hashes->digests_shown_tmp[salt_buf->digests_offset], salt_buf->digests_cnt * sizeof (u32));
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (u32), salt_buf->digests_cnt * sizeof (u32), &hashes->digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (cpt_ctx->cpt_pos == CPT_CACHE) cpt_ctx->cpt_pos = 0;
|
||||
|
||||
hc_thread_mutex_unlock (status_ctx->mux_display);
|
||||
}
|
||||
|
||||
hcfree (cracked);
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
if (run_cuda_kernel_bzero (hashcat_ctx, device_param, device_param->cuda_d_result, sizeof (u32)) == -1) return -1;
|
||||
}
|
||||
|
||||
if (cpt_cracked > 0)
|
||||
{
|
||||
hc_thread_mutex_lock (status_ctx->mux_display);
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
if (run_hip_kernel_bzero (hashcat_ctx, device_param, device_param->hip_d_result, sizeof (u32)) == -1) return -1;
|
||||
}
|
||||
|
||||
cpt_ctx->cpt_buf[cpt_ctx->cpt_pos].timestamp = time (NULL);
|
||||
cpt_ctx->cpt_buf[cpt_ctx->cpt_pos].cracked = cpt_cracked;
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
if (run_opencl_kernel_bzero (hashcat_ctx, device_param, device_param->opencl_d_result, sizeof (u32)) == -1) return -1;
|
||||
|
||||
cpt_ctx->cpt_pos++;
|
||||
|
||||
cpt_ctx->cpt_total += cpt_cracked;
|
||||
|
||||
if (cpt_ctx->cpt_pos == CPT_CACHE) cpt_ctx->cpt_pos = 0;
|
||||
|
||||
hc_thread_mutex_unlock (status_ctx->mux_display);
|
||||
}
|
||||
|
||||
num_cracked = 0;
|
||||
|
||||
if (device_param->is_cuda == true)
|
||||
{
|
||||
CU_rc = hc_cuMemcpyHtoD (hashcat_ctx, device_param->cuda_d_result, &num_cracked, sizeof (u32));
|
||||
|
||||
if (CU_rc == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_hip == true)
|
||||
{
|
||||
HIP_rc = hc_hipMemcpyHtoD (hashcat_ctx, device_param->hip_d_result, &num_cracked, sizeof (u32));
|
||||
|
||||
if (HIP_rc == -1) return -1;
|
||||
}
|
||||
|
||||
if (device_param->is_opencl == true)
|
||||
{
|
||||
CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->opencl_command_queue, device_param->opencl_d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
|
||||
|
||||
if (CL_rc == -1) return -1;
|
||||
}
|
||||
if (hc_clFlush (hashcat_ctx, device_param->opencl_command_queue) == -1) return -1;
|
||||
}
|
||||
|
||||
return 0;
|
||||
@ -1569,7 +1672,6 @@ int hashes_init_stage2 (hashcat_ctx_t *hashcat_ctx)
|
||||
u32 digests_done = 0;
|
||||
|
||||
u32 *digests_shown = (u32 *) hccalloc (digests_cnt, sizeof (u32));
|
||||
u32 *digests_shown_tmp = (u32 *) hccalloc (digests_cnt, sizeof (u32));
|
||||
|
||||
u32 salts_cnt = 0;
|
||||
u32 salts_done = 0;
|
||||
@ -1706,7 +1808,6 @@ int hashes_init_stage2 (hashcat_ctx_t *hashcat_ctx)
|
||||
hashes->digests_done = digests_done;
|
||||
hashes->digests_buf = digests_buf_new;
|
||||
hashes->digests_shown = digests_shown;
|
||||
hashes->digests_shown_tmp = digests_shown_tmp;
|
||||
|
||||
hashes->salts_cnt = salts_cnt;
|
||||
hashes->salts_done = salts_done;
|
||||
@ -2193,7 +2294,6 @@ void hashes_destroy (hashcat_ctx_t *hashcat_ctx)
|
||||
|
||||
hcfree (hashes->digests_buf);
|
||||
hcfree (hashes->digests_shown);
|
||||
hcfree (hashes->digests_shown_tmp);
|
||||
|
||||
hcfree (hashes->salts_buf);
|
||||
hcfree (hashes->salts_shown);
|
||||
|
||||
@ -150,6 +150,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
CHECK_DEFINED (module_ctx->module_esalt_size);
|
||||
CHECK_DEFINED (module_ctx->module_extra_buffer_size);
|
||||
CHECK_DEFINED (module_ctx->module_extra_tmp_size);
|
||||
CHECK_DEFINED (module_ctx->module_extra_tuningdb_block);
|
||||
CHECK_DEFINED (module_ctx->module_forced_outfile_format);
|
||||
CHECK_DEFINED (module_ctx->module_hash_binary_count);
|
||||
CHECK_DEFINED (module_ctx->module_hash_binary_parse);
|
||||
|
||||
@ -138,6 +138,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -161,6 +161,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -164,6 +164,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -156,6 +156,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -160,6 +160,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -155,6 +155,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -268,6 +268,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -170,6 +170,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -154,6 +154,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -162,6 +162,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -161,6 +161,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -140,6 +140,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -140,6 +140,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -136,6 +136,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -154,6 +154,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -180,6 +180,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -184,6 +184,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -180,6 +180,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -160,6 +160,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -179,6 +179,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -162,6 +162,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -176,6 +176,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -170,6 +170,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -153,6 +153,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -181,6 +181,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -188,6 +188,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -158,6 +158,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -155,6 +155,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -122,6 +122,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -153,6 +153,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -249,6 +249,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -262,6 +262,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -341,6 +341,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -127,6 +127,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -138,6 +138,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -141,6 +141,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -165,6 +165,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -168,6 +168,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -175,6 +175,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -201,6 +201,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -197,6 +197,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -201,6 +201,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -177,6 +177,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -172,6 +172,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -204,6 +204,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -176,6 +176,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -95,13 +95,6 @@ int module_build_plain_postprocess (MAYBE_UNUSED const hashconfig_t *hashconfig,
|
||||
return src_len;
|
||||
}
|
||||
|
||||
u32 module_kernel_threads_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
|
||||
{
|
||||
const u32 kernel_threads_max = 64; // performance only optimization
|
||||
|
||||
return kernel_threads_max;
|
||||
}
|
||||
|
||||
u32 module_kernel_loops_max (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
|
||||
{
|
||||
u32 kernel_loops_max = KERNEL_LOOPS_MAX;
|
||||
@ -315,6 +308,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
@ -345,7 +339,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_kernel_accel_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_max = module_kernel_loops_max;
|
||||
module_ctx->module_kernel_loops_min = module_kernel_loops_min;
|
||||
module_ctx->module_kernel_threads_max = module_kernel_threads_max;
|
||||
module_ctx->module_kernel_threads_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kern_type = module_kern_type;
|
||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||
|
||||
@ -262,6 +262,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -202,6 +202,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -198,6 +198,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -227,6 +227,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -204,6 +204,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -189,6 +189,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -203,6 +203,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -173,6 +173,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -174,6 +174,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -191,6 +191,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
@ -567,6 +567,7 @@ void module_init (module_ctx_t *module_ctx)
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tuningdb_block = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user