mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-22 16:18:09 +00:00
Prepare pure kernel rule engine for performance optimization
This commit is contained in:
jsteube
parent
e47506c610
commit
0b0abb5c12
138
OpenCL/inc_rp.cl
138
OpenCL/inc_rp.cl
@ -3,65 +3,45 @@
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
int class_digit (const u8 c)
|
||||
{
|
||||
if ((c >= '0') && (c <= '9')) return 1;
|
||||
#define MAYBE_UNUSED
|
||||
|
||||
return 0;
|
||||
static void upper_at (u8 *buf, const int pos)
|
||||
{
|
||||
const u8 c = buf[pos];
|
||||
|
||||
if ((c >= 'a') && (c <= 'z')) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
int class_lower (const u8 c)
|
||||
static void lower_at (u8 *buf, const int pos)
|
||||
{
|
||||
if ((c >= 'a') && (c <= 'z')) return 1;
|
||||
const u8 c = buf[pos];
|
||||
|
||||
return 0;
|
||||
if ((c >= 'A') && (c <= 'Z')) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
int class_upper (const u8 c)
|
||||
static void toggle_at (u8 *buf, const int pos)
|
||||
{
|
||||
if ((c >= 'A') && (c <= 'Z')) return 1;
|
||||
const u8 c = buf[pos];
|
||||
|
||||
return 0;
|
||||
if ((c >= 'a') && (c <= 'z')) buf[pos] ^= 0x20;
|
||||
if ((c >= 'A') && (c <= 'Z')) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
int class_alpha (const u8 c)
|
||||
{
|
||||
if (class_lower (c) == 1) return 1;
|
||||
if (class_upper (c) == 1) return 1;
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
void upper_at (u8 *buf, const int pos)
|
||||
{
|
||||
if (class_lower (buf[pos])) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
void lower_at (u8 *buf, const int pos)
|
||||
{
|
||||
if (class_upper (buf[pos])) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
void toggle_at (u8 *buf, const int pos)
|
||||
{
|
||||
if (class_alpha (buf[pos])) buf[pos] ^= 0x20;
|
||||
}
|
||||
|
||||
void mangle_switch (u8 *buf, const int l, const int r)
|
||||
static void mangle_switch (u8 *buf, const int l, const int r)
|
||||
{
|
||||
const u8 c = buf[r];
|
||||
buf[r] = buf[l];
|
||||
buf[l] = c;
|
||||
}
|
||||
|
||||
int mangle_lrest (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++) lower_at (buf, pos);
|
||||
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_lrest_ufirst (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++) lower_at (buf, pos);
|
||||
|
||||
@ -70,14 +50,14 @@ int mangle_lrest_ufirst (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_urest (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++) upper_at (buf, pos);
|
||||
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_urest_lfirst (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++) upper_at (buf, pos);
|
||||
|
||||
@ -86,14 +66,14 @@ int mangle_urest_lfirst (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_trest (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++) toggle_at (buf, pos);
|
||||
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_toggle_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_toggle_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -102,7 +82,7 @@ int mangle_toggle_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_reverse (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_reverse (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int l = 0; l < len / 2; l++)
|
||||
{
|
||||
@ -114,7 +94,7 @@ int mangle_reverse (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_dupeword (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupeword (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len * 2;
|
||||
|
||||
@ -127,7 +107,7 @@ int mangle_dupeword (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_dupeword_times (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupeword_times (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = (len * p0) + len;
|
||||
|
||||
@ -140,7 +120,7 @@ int mangle_dupeword_times (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_reflect (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_reflect (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len * 2;
|
||||
|
||||
@ -153,7 +133,7 @@ int mangle_reflect (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return out_len;
|
||||
}
|
||||
|
||||
int mangle_append (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_append (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len + 1;
|
||||
|
||||
@ -164,7 +144,7 @@ int mangle_append (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_prepend (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_prepend (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len + 1;
|
||||
|
||||
@ -178,7 +158,7 @@ int mangle_prepend (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_rotate_left (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_rotate_left (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int l = 0, r = len - 1; r > l; r--)
|
||||
{
|
||||
@ -188,7 +168,7 @@ int mangle_rotate_left (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_rotate_right (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_rotate_right (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int l = 0, r = len - 1; l < r; l++)
|
||||
{
|
||||
@ -198,7 +178,7 @@ int mangle_rotate_right (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_delete_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_delete_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -212,19 +192,19 @@ int mangle_delete_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len - 1);
|
||||
}
|
||||
|
||||
int mangle_delete_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_delete_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
return mangle_delete_at (0, p1, buf, len);
|
||||
}
|
||||
|
||||
int mangle_delete_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_delete_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (len == 0) return 0;
|
||||
|
||||
return mangle_delete_at (len - 1, p1, buf, len);
|
||||
}
|
||||
|
||||
int mangle_extract (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_extract (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -243,7 +223,7 @@ int mangle_extract (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (p1);
|
||||
}
|
||||
|
||||
int mangle_omit (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_omit (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -262,7 +242,7 @@ int mangle_omit (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len - p1);
|
||||
}
|
||||
|
||||
int mangle_insert (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_insert (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len + 1) return (len);
|
||||
|
||||
@ -280,7 +260,7 @@ int mangle_insert (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_overstrike (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_overstrike (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -289,7 +269,7 @@ int mangle_overstrike (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_truncate_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_truncate_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -301,7 +281,7 @@ int mangle_truncate_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (p0);
|
||||
}
|
||||
|
||||
int mangle_replace (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_replace (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
for (int pos = 0; pos < len; pos++)
|
||||
{
|
||||
@ -313,7 +293,7 @@ int mangle_replace (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_purgechar (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_purgechar (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
int out_len = 0;
|
||||
|
||||
@ -334,7 +314,7 @@ int mangle_purgechar (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_dupechar_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupechar_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len + p0;
|
||||
|
||||
@ -350,7 +330,7 @@ int mangle_dupechar_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_dupechar_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupechar_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len + p0;
|
||||
|
||||
@ -366,7 +346,7 @@ int mangle_dupechar_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_dupechar_all (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupechar_all (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
const int out_len = len + len;
|
||||
|
||||
@ -384,7 +364,7 @@ int mangle_dupechar_all (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_switch_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_switch_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (len < 2) return (len);
|
||||
|
||||
@ -393,7 +373,7 @@ int mangle_switch_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_switch_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_switch_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (len < 2) return (len);
|
||||
|
||||
@ -402,7 +382,7 @@ int mangle_switch_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_switch_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_switch_at (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
if (p1 >= len) return (len);
|
||||
@ -412,7 +392,7 @@ int mangle_switch_at (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_chr_shiftl (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_chr_shiftl (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -421,7 +401,7 @@ int mangle_chr_shiftl (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_chr_shiftr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_chr_shiftr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -430,7 +410,7 @@ int mangle_chr_shiftr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_chr_incr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_chr_incr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -439,7 +419,7 @@ int mangle_chr_incr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_chr_decr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_chr_decr (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -448,7 +428,7 @@ int mangle_chr_decr (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_replace_np1 (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_replace_np1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if ((p0 + 1) >= len) return (len);
|
||||
|
||||
@ -457,7 +437,7 @@ int mangle_replace_np1 (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_replace_nm1 (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_replace_nm1 (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 == 0) return (len);
|
||||
|
||||
@ -468,7 +448,7 @@ int mangle_replace_nm1 (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int mangle_dupeblock_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupeblock_first (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -486,7 +466,7 @@ int mangle_dupeblock_first (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_dupeblock_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_dupeblock_last (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
if (p0 >= len) return (len);
|
||||
|
||||
@ -504,7 +484,7 @@ int mangle_dupeblock_last (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (out_len);
|
||||
}
|
||||
|
||||
int mangle_title_sep (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
static int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int len)
|
||||
{
|
||||
int upper_next = 1;
|
||||
|
||||
@ -532,7 +512,7 @@ int mangle_title_sep (const u8 p0, const u8 p1, u8 *buf, const int len)
|
||||
return (len);
|
||||
}
|
||||
|
||||
int apply_rule (const u32 name, const u8 p0, const u8 p1, u8 *buf, const int in_len)
|
||||
static int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u8 *buf, const int in_len)
|
||||
{
|
||||
int out_len = in_len;
|
||||
|
||||
@ -583,12 +563,8 @@ int apply_rule (const u32 name, const u8 p0, const u8 p1, u8 *buf, const int in_
|
||||
return out_len;
|
||||
}
|
||||
|
||||
u32 apply_rules (__global const u32 *cmds, u32 in_buf[64], const int in_len, u32 out_buf[64])
|
||||
int apply_rules (__global const u32 *cmds, u32 buf[64], const int in_len)
|
||||
{
|
||||
const int in_lenv = ceil ((float) in_len / 4);
|
||||
|
||||
for (int i = 0; i < in_lenv; i++) out_buf[i] = in_buf[i];
|
||||
|
||||
int out_len = in_len;
|
||||
|
||||
for (u32 i = 0; cmds[i] != 0; i++)
|
||||
@ -599,8 +575,8 @@ u32 apply_rules (__global const u32 *cmds, u32 in_buf[64], const int in_len, u32
|
||||
const u8 p0 = (cmd >> 8) & 0xff;
|
||||
const u8 p1 = (cmd >> 16) & 0xff;
|
||||
|
||||
out_len = apply_rule (name, p0, p1, (u8 *) out_buf, out_len);
|
||||
out_len = apply_rule (name, p0, p1, (u8 *) buf, out_len);
|
||||
}
|
||||
|
||||
return (u32) out_len;
|
||||
return out_len;
|
||||
}
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update_utf16le (&ctx, out_buf, out_len);
|
||||
md5_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update_utf16le (&ctx, out_buf, out_len);
|
||||
md5_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update_utf16le (&ctx, out_buf, out_len);
|
||||
md5_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update_utf16le (&ctx, out_buf, out_len);
|
||||
md5_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,13 +49,13 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx;
|
||||
|
||||
md5_hmac_init (&ctx, out_buf, out_len);
|
||||
md5_hmac_init (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
@ -106,16 +97,7 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -134,13 +116,13 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx;
|
||||
|
||||
md5_hmac_init (&ctx, out_buf, out_len);
|
||||
md5_hmac_init (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -62,13 +53,13 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
md5_hmac_update (&ctx, out_buf, out_len);
|
||||
md5_hmac_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_final (&ctx);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -140,13 +122,13 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
md5_hmac_update (&ctx, out_buf, out_len);
|
||||
md5_hmac_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,13 +49,13 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
@ -106,16 +97,7 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -134,13 +116,13 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -62,13 +53,13 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -140,13 +122,13 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
@ -118,16 +109,7 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -135,15 +117,15 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update (&ctx, out_buf, out_len);
|
||||
md4_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update (&ctx, out_buf, out_len);
|
||||
md4_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx0;
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_utf16le (&ctx0, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx0);
|
||||
|
||||
@ -121,16 +112,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -149,15 +131,15 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx0;
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_utf16le (&ctx0, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx0);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha224_ctx_t ctx;
|
||||
|
||||
sha224_init (&ctx);
|
||||
|
||||
sha224_update_swap (&ctx, out_buf, out_len);
|
||||
sha224_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha224_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha224_ctx_t ctx;
|
||||
|
||||
sha224_init (&ctx);
|
||||
|
||||
sha224_update_swap (&ctx, out_buf, out_len);
|
||||
sha224_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha224_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,13 +49,13 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_ctx_t ctx;
|
||||
|
||||
sha256_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha256_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
@ -106,16 +97,7 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -134,13 +116,13 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_ctx_t ctx;
|
||||
|
||||
sha256_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha256_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -62,13 +53,13 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_final (&ctx);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -140,13 +122,13 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha256_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha256_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_hmac_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,13 +49,13 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_ctx_t ctx;
|
||||
|
||||
sha512_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha512_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
@ -106,16 +97,7 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -134,13 +116,13 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_ctx_t ctx;
|
||||
|
||||
sha512_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha512_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -62,13 +53,13 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_final (&ctx);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -140,13 +122,13 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_ctx_t ctx = ctx0;
|
||||
|
||||
sha512_hmac_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_hmac_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_hmac_final (&ctx);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -88,15 +79,15 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -186,16 +177,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -214,15 +196,15 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = 32;
|
||||
|
||||
@ -88,15 +79,15 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -209,16 +200,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = 32;
|
||||
|
||||
@ -237,15 +219,15 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -88,15 +79,15 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -199,16 +190,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -227,15 +209,15 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -64,13 +55,13 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
@ -112,16 +103,7 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -146,13 +128,13 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = 32;
|
||||
|
||||
@ -88,15 +79,15 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -209,16 +200,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = 32;
|
||||
|
||||
@ -237,15 +219,15 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m04010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -83,13 +74,13 @@ __kernel void m04010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1 = ctx0;
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx1);
|
||||
|
||||
@ -188,16 +179,7 @@ __kernel void m04010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -211,13 +193,13 @@ __kernel void m04010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1 = ctx0;
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx1);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -94,15 +85,15 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1;
|
||||
|
||||
md5_init (&ctx1);
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx1, s, salt_len);
|
||||
|
||||
@ -203,16 +194,7 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -237,15 +219,15 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1;
|
||||
|
||||
md5_init (&ctx1);
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx1, s, salt_len);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -88,15 +79,15 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -186,16 +177,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -214,15 +196,15 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -61,16 +61,7 @@ __kernel void m04400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -78,15 +69,15 @@ __kernel void m04400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
@ -179,16 +170,7 @@ __kernel void m04400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -196,15 +178,15 @@ __kernel void m04400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m04500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -77,15 +68,15 @@ __kernel void m04500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
@ -178,16 +169,7 @@ __kernel void m04500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -195,15 +177,15 @@ __kernel void m04500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m04520_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -83,15 +74,15 @@ __kernel void m04520_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
@ -193,16 +184,7 @@ __kernel void m04520_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -216,15 +198,15 @@ __kernel void m04520_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
|
||||
@ -61,16 +61,7 @@ __kernel void m04700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -78,15 +69,15 @@ __kernel void m04700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
@ -175,16 +166,7 @@ __kernel void m04700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -192,15 +174,15 @@ __kernel void m04700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update (&ctx0, out_buf, out_len);
|
||||
md5_update (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx0);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len - 1;
|
||||
|
||||
@ -64,13 +55,13 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
@ -112,16 +103,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len - 1;
|
||||
|
||||
@ -146,13 +128,13 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -64,13 +55,13 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
@ -112,16 +103,7 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -146,13 +128,13 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m05100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m05100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m05100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -116,15 +98,15 @@ __kernel void m05100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m05300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,13 +38,13 @@ __kernel void m05300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
|
||||
md5_hmac_init (&ctx0, out_buf, out_len);
|
||||
md5_hmac_init (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_update_global (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
|
||||
|
||||
@ -125,16 +116,7 @@ __kernel void m05300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -142,13 +124,13 @@ __kernel void m05300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
|
||||
md5_hmac_init (&ctx0, out_buf, out_len);
|
||||
md5_hmac_init (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_hmac_update_global (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,13 +38,13 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
|
||||
sha1_hmac_init_swap (&ctx0, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update_global_swap (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
|
||||
|
||||
@ -125,16 +116,7 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -142,13 +124,13 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
|
||||
sha1_hmac_init_swap (&ctx0, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update_global_swap (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
|
||||
|
||||
|
||||
@ -544,16 +544,7 @@ __kernel void m05500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -561,15 +552,15 @@ __kernel void m05500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
@ -689,16 +680,7 @@ __kernel void m05500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -706,15 +688,15 @@ __kernel void m05500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
|
||||
@ -31,16 +31,7 @@ __kernel void m05600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -48,15 +39,15 @@ __kernel void m05600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx1;
|
||||
|
||||
md4_init (&ctx1);
|
||||
|
||||
md4_update_utf16le (&ctx1, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx1);
|
||||
|
||||
@ -151,16 +142,7 @@ __kernel void m05600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -168,15 +150,15 @@ __kernel void m05600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx1;
|
||||
|
||||
md4_init (&ctx1);
|
||||
|
||||
md4_update_utf16le (&ctx1, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx1);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m06000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m06000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
ripemd160_ctx_t ctx;
|
||||
|
||||
ripemd160_init (&ctx);
|
||||
|
||||
ripemd160_update (&ctx, out_buf, out_len);
|
||||
ripemd160_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
ripemd160_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m06000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m06000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
ripemd160_ctx_t ctx;
|
||||
|
||||
ripemd160_init (&ctx);
|
||||
|
||||
ripemd160_update (&ctx, out_buf, out_len);
|
||||
ripemd160_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
ripemd160_final (&ctx);
|
||||
|
||||
|
||||
@ -78,15 +78,15 @@ __kernel void m06100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
whirlpool_ctx_t ctx;
|
||||
|
||||
whirlpool_init (&ctx, s_Ch, s_Cl);
|
||||
|
||||
whirlpool_update (&ctx, out_buf, out_len);
|
||||
whirlpool_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
whirlpool_final (&ctx);
|
||||
|
||||
@ -174,15 +174,15 @@ __kernel void m06100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
whirlpool_ctx_t ctx;
|
||||
|
||||
whirlpool_init (&ctx, s_Ch, s_Cl);
|
||||
|
||||
whirlpool_update (&ctx, out_buf, out_len);
|
||||
whirlpool_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
whirlpool_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m07000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m07000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
/**
|
||||
* pepper
|
||||
@ -128,16 +119,7 @@ __kernel void m07000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -151,13 +133,13 @@ __kernel void m07000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
/**
|
||||
* pepper
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m07300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,13 +38,13 @@ __kernel void m07300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update_global (&ctx, rakp_bufs[digests_offset].salt_buf, rakp_bufs[digests_offset].salt_len);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m07300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,13 +94,13 @@ __kernel void m07300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_swap (&ctx, out_buf, out_len);
|
||||
sha1_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_hmac_update_global (&ctx, rakp_bufs[digests_offset].salt_buf, rakp_bufs[digests_offset].salt_len);
|
||||
|
||||
|
||||
@ -283,16 +283,7 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
__local RC4_KEY rc4_keys[64];
|
||||
|
||||
@ -320,15 +311,15 @@ __kernel void m07500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
@ -361,16 +352,7 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
__local RC4_KEY rc4_keys[64];
|
||||
|
||||
@ -398,15 +380,15 @@ __kernel void m07500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m08100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m08100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, w, pw_len + 1);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m08100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m08100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_swap (&ctx, w, pw_len + 1);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len + 1);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -71,9 +62,9 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
@ -83,7 +74,7 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
ctx1.len = 1;
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx1, s_pc, salt_len_pc + 1);
|
||||
|
||||
@ -162,16 +153,7 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -203,9 +185,9 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
@ -215,7 +197,7 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
ctx1.len = 1;
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx1, s_pc, salt_len_pc + 1);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m08400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -83,15 +74,15 @@ __kernel void m08400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
@ -232,16 +223,7 @@ __kernel void m08400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -255,15 +237,15 @@ __kernel void m08400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m09900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m09900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, w, 100);
|
||||
md5_update (&ctx, tmp.i, 100);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -95,16 +86,7 @@ __kernel void m09900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -112,15 +94,15 @@ __kernel void m09900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx;
|
||||
|
||||
md5_init (&ctx);
|
||||
|
||||
md5_update (&ctx, w, 100);
|
||||
md5_update (&ctx, tmp.i, 100);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -47,15 +47,15 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha384_ctx_t ctx;
|
||||
|
||||
sha384_init (&ctx);
|
||||
|
||||
sha384_update (&ctx, out_buf, out_len);
|
||||
sha384_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
@ -112,15 +112,15 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha384_ctx_t ctx;
|
||||
|
||||
sha384_init (&ctx);
|
||||
|
||||
sha384_update (&ctx, out_buf, out_len);
|
||||
sha384_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m11000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -53,13 +44,13 @@ __kernel void m11000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -99,16 +90,7 @@ __kernel void m11000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -122,13 +104,13 @@ __kernel void m11000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx = ctx0;
|
||||
|
||||
md5_update (&ctx, out_buf, out_len);
|
||||
md5_update (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -82,16 +82,7 @@ __kernel void m11100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -99,15 +90,15 @@ __kernel void m11100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1;
|
||||
|
||||
md5_init (&ctx1);
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
u32 s0[4];
|
||||
u32 s1[4];
|
||||
@ -249,16 +240,7 @@ __kernel void m11100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -266,15 +248,15 @@ __kernel void m11100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1;
|
||||
|
||||
md5_init (&ctx1);
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
u32 s0[4];
|
||||
u32 s1[4];
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m11200_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -53,15 +44,15 @@ __kernel void m11200_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx2;
|
||||
|
||||
sha1_init (&ctx2);
|
||||
|
||||
sha1_update_swap (&ctx2, out_buf, out_len);
|
||||
sha1_update_swap (&ctx2, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx2);
|
||||
|
||||
@ -167,16 +158,7 @@ __kernel void m11200_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -190,15 +172,15 @@ __kernel void m11200_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx2;
|
||||
|
||||
sha1_init (&ctx2);
|
||||
|
||||
sha1_update_swap (&ctx2, out_buf, out_len);
|
||||
sha1_update_swap (&ctx2, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx2);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -83,13 +74,13 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1 = ctx0;
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx1);
|
||||
|
||||
@ -179,16 +170,7 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -202,13 +184,13 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_ctx_t ctx1 = ctx0;
|
||||
|
||||
md5_update (&ctx1, out_buf, out_len);
|
||||
md5_update (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
md5_final (&ctx1);
|
||||
|
||||
|
||||
@ -76,16 +76,7 @@ __kernel void m12600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -93,15 +84,15 @@ __kernel void m12600_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
@ -234,16 +225,7 @@ __kernel void m12600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -251,15 +233,15 @@ __kernel void m12600_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_swap (&ctx0, out_buf, out_len);
|
||||
sha1_update_swap (&ctx0, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx0);
|
||||
|
||||
|
||||
@ -392,16 +392,7 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
__local RC4_KEY rc4_keys[64];
|
||||
|
||||
@ -418,15 +409,15 @@ __kernel void m13100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
@ -461,16 +452,7 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
__local RC4_KEY rc4_keys[64];
|
||||
|
||||
@ -487,15 +469,15 @@ __kernel void m13100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_ctx_t ctx;
|
||||
|
||||
md4_init (&ctx);
|
||||
|
||||
md4_update_utf16le (&ctx, out_buf, out_len);
|
||||
md4_update_utf16le (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
md4_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m13300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m13300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -97,16 +88,7 @@ __kernel void m13300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -114,15 +96,15 @@ __kernel void m13300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -63,16 +63,7 @@ __kernel void m13500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -80,13 +71,13 @@ __kernel void m13500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
@ -159,16 +150,7 @@ __kernel void m13500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -176,13 +158,13 @@ __kernel void m13500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx = ctx0;
|
||||
|
||||
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m13800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -47,15 +38,15 @@ __kernel void m13800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
|
||||
|
||||
@ -97,16 +88,7 @@ __kernel void m13800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -114,15 +96,15 @@ __kernel void m13800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_ctx_t ctx;
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
|
||||
sha256_update_utf16le_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m13900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -83,15 +74,15 @@ __kernel void m13900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
@ -232,16 +223,7 @@ __kernel void m13900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -255,15 +237,15 @@ __kernel void m13900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1;
|
||||
|
||||
sha1_init (&ctx1);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_final (&ctx1);
|
||||
|
||||
|
||||
@ -60,16 +60,7 @@ __kernel void m14400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -131,9 +122,9 @@ __kernel void m14400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1 = ctx0;
|
||||
|
||||
@ -156,7 +147,7 @@ __kernel void m14400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_update_64 (&ctx1, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
d40[0] = 0x2d2d2d2d;
|
||||
d40[1] = 0;
|
||||
@ -244,7 +235,7 @@ __kernel void m14400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_update_64 (&ctx, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
d40[0] = 0x2d2d2d2d;
|
||||
d40[1] = 0;
|
||||
@ -328,16 +319,7 @@ __kernel void m14400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -399,9 +381,9 @@ __kernel void m14400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx1 = ctx0;
|
||||
|
||||
@ -424,7 +406,7 @@ __kernel void m14400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_update_64 (&ctx1, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_swap (&ctx1, out_buf, out_len);
|
||||
sha1_update_swap (&ctx1, tmp.i, tmp.pw_len);
|
||||
|
||||
d40[0] = 0x2d2d2d2d;
|
||||
d40[1] = 0;
|
||||
@ -512,7 +494,7 @@ __kernel void m14400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_update_64 (&ctx, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
d40[0] = 0x2d2d2d2d;
|
||||
d40[1] = 0;
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m15000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m15000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
@ -108,16 +99,7 @@ __kernel void m15000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -136,15 +118,15 @@ __kernel void m15000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_ctx_t ctx;
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_swap (&ctx, out_buf, out_len);
|
||||
sha512_update_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha512_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -30,16 +30,7 @@ __kernel void m15500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -58,15 +49,15 @@ __kernel void m15500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_utf16be_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16be_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
@ -114,16 +105,7 @@ __kernel void m15500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
pw_t pw = pws[gid];
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
@ -142,15 +124,15 @@ __kernel void m15500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
u32 out_buf[64] = { 0 };
|
||||
pw_t tmp = pw;
|
||||
|
||||
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_ctx_t ctx;
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_utf16be_swap (&ctx, out_buf, out_len);
|
||||
sha1_update_utf16be_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
sha1_update (&ctx, s, salt_len);
|
||||
|
||||
|
||||
@ -6,6 +6,6 @@
|
||||
#ifndef _RP_KERNEL_ON_CPU_H
|
||||
#define _RP_KERNEL_ON_CPU_H
|
||||
|
||||
int apply_rules (const u32 *cmds, u32 in_buf[64], const int in_len, u32 out_buf[64]);
|
||||
int apply_rules (const u32 *cmds, u32 buf[64], const int in_len);
|
||||
|
||||
#endif // _RP_KERNEL_ON_CPU_H
|
||||
|
||||
@ -56,7 +56,12 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
|
||||
}
|
||||
else
|
||||
{
|
||||
plain_len = (int) apply_rules (straight_ctx->kernel_rules_buf[off].cmds, pw.i, pw.pw_len, plain_buf);
|
||||
for (int i = 0; i < 64; i++)
|
||||
{
|
||||
plain_buf[i] = pw.i[i];
|
||||
}
|
||||
|
||||
plain_len = (int) apply_rules (straight_ctx->kernel_rules_buf[off].cmds, plain_buf, pw.pw_len);
|
||||
}
|
||||
|
||||
if (plain_len > (int) hashconfig->pw_max) plain_len = (int) hashconfig->pw_max;
|
||||
|
||||
@ -567,14 +567,8 @@ static int apply_rule (const u32 name, MAYBE_UNUSED const u8 p0, MAYBE_UNUSED co
|
||||
return out_len;
|
||||
}
|
||||
|
||||
int apply_rules (const u32 *cmds, u32 in_buf[64], const int in_len, u32 out_buf[64])
|
||||
int apply_rules (const u32 *cmds, u32 buf[64], const int in_len)
|
||||
{
|
||||
//const int in_lenv = ceil ((float) in_len / 4);
|
||||
|
||||
//for (int i = 0; i < in_lenv; i++) out_buf[i] = in_buf[i];
|
||||
|
||||
for (int i = 0; i < 64; i++) out_buf[i] = in_buf[i];
|
||||
|
||||
int out_len = in_len;
|
||||
|
||||
for (u32 i = 0; cmds[i] != 0; i++)
|
||||
@ -585,7 +579,7 @@ int apply_rules (const u32 *cmds, u32 in_buf[64], const int in_len, u32 out_buf[
|
||||
const u8 p0 = (cmd >> 8) & 0xff;
|
||||
const u8 p1 = (cmd >> 16) & 0xff;
|
||||
|
||||
out_len = apply_rule (name, p0, p1, (u8 *) out_buf, out_len);
|
||||
out_len = apply_rule (name, p0, p1, (u8 *) buf, out_len);
|
||||
}
|
||||
|
||||
return out_len;
|
||||
|
||||
@ -128,7 +128,12 @@ int process_stdout (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
|
||||
}
|
||||
else
|
||||
{
|
||||
plain_len = (int) apply_rules (straight_ctx->kernel_rules_buf[off].cmds, pw.i, pw.pw_len, plain_buf);
|
||||
for (int i = 0; i < 64; i++)
|
||||
{
|
||||
plain_buf[i] = pw.i[i];
|
||||
}
|
||||
|
||||
plain_len = (int) apply_rules (straight_ctx->kernel_rules_buf[off].cmds, plain_buf, pw.pw_len);
|
||||
}
|
||||
|
||||
if (plain_len > hashconfig->pw_max) plain_len = hashconfig->pw_max;
|
||||
|
||||
Loading…
Reference in New Issue
Block a user