mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-26 18:08:20 +00:00
Fix -m 5100 multihash and benchmark
This commit is contained in:
parent
98a6aff4a8
commit
c0a38846e3
@ -52,18 +52,6 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
|
|
||||||
const u32 pw_len = pws[gid].pw_len;
|
const u32 pw_len = pws[gid].pw_len;
|
||||||
|
|
||||||
/**
|
|
||||||
* digest
|
|
||||||
*/
|
|
||||||
|
|
||||||
const u32 search[4] =
|
|
||||||
{
|
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
|
||||||
};
|
|
||||||
|
|
||||||
/**
|
/**
|
||||||
* loop
|
* loop
|
||||||
*/
|
*/
|
||||||
@ -164,14 +152,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_M_SIMD (a, b, e, f);
|
COMPARE_M_SIMD (a, b, z, z);
|
||||||
|
COMPARE_M_SIMD (b, c, z, z);
|
||||||
COMPARE_M_SIMD (b, c, e, f);
|
COMPARE_M_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_M_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -221,8 +206,8 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
{
|
{
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
0,
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
0
|
||||||
};
|
};
|
||||||
|
|
||||||
/**
|
/**
|
||||||
@ -325,14 +310,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_S_SIMD (a, b, e, f);
|
COMPARE_S_SIMD (a, b, z, z);
|
||||||
|
COMPARE_S_SIMD (b, c, z, z);
|
||||||
COMPARE_S_SIMD (b, c, e, f);
|
COMPARE_S_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_S_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -205,14 +205,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_M_SIMD (a, b, e, f);
|
COMPARE_M_SIMD (a, b, z, z);
|
||||||
|
COMPARE_M_SIMD (b, c, z, z);
|
||||||
COMPARE_M_SIMD (b, c, e, f);
|
COMPARE_M_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_M_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -262,8 +259,8 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
{
|
{
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
0,
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
0
|
||||||
};
|
};
|
||||||
|
|
||||||
/**
|
/**
|
||||||
@ -421,14 +418,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_S_SIMD (a, b, e, f);
|
COMPARE_S_SIMD (a, b, z, z);
|
||||||
|
COMPARE_S_SIMD (b, c, z, z);
|
||||||
COMPARE_S_SIMD (b, c, e, f);
|
COMPARE_S_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_S_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
@ -145,14 +145,11 @@ void m05100m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_M_SIMD (a, b, e, f);
|
COMPARE_M_SIMD (a, b, z, z);
|
||||||
|
COMPARE_M_SIMD (b, c, z, z);
|
||||||
COMPARE_M_SIMD (b, c, e, f);
|
COMPARE_M_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_M_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
@ -173,8 +170,8 @@ void m05100s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
|||||||
{
|
{
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
0,
|
||||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
0
|
||||||
};
|
};
|
||||||
|
|
||||||
/**
|
/**
|
||||||
@ -293,14 +290,11 @@ void m05100s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
|
|||||||
c += MD5M_C;
|
c += MD5M_C;
|
||||||
d += MD5M_D;
|
d += MD5M_D;
|
||||||
|
|
||||||
u32x e = 0;
|
u32x z = 0;
|
||||||
u32x f = 0;
|
|
||||||
|
|
||||||
COMPARE_S_SIMD (a, b, e, f);
|
COMPARE_S_SIMD (a, b, z, z);
|
||||||
|
COMPARE_S_SIMD (b, c, z, z);
|
||||||
COMPARE_S_SIMD (b, c, e, f);
|
COMPARE_S_SIMD (c, d, z, z);
|
||||||
|
|
||||||
COMPARE_S_SIMD (c, d, e, f);
|
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
Loading…
Reference in New Issue
Block a user