arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-26 09:58:16 +00:00
Code Issues Releases Wiki Activity

Prepare for pure kernel rule engine support

Browse Source
This commit is contained in:
jsteube 2017-08-11 16:09:12 +02:00
parent 34d882a116
commit 8a0d21360b
78 changed files with 956 additions and 310 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

296
OpenCL/inc_rp.cl Normal file
View File

@ -0,0 +1,296 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
int class_digit (const u8 c)
{
if ((c >= '0') && (c <= '9')) return 1;
return 0;
}
int class_lower (const u8 c)
{
if ((c >= 'a') && (c <= 'z')) return 1;
return 0;
}
int class_upper (const u8 c)
{
if ((c >= 'A') && (c <= 'Z')) return 1;
return 0;
}
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)
{
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)
{
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)
{
for (int pos = 0; pos < len; pos++) lower_at (buf, pos);
upper_at (buf, 0);
return (len);
}
int mangle_urest (const u8 p0, 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)
{
for (int pos = 0; pos < len; pos++) upper_at (buf, pos);
lower_at (buf, 0);
return (len);
}
int mangle_trest (const u8 p0, 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)
{
if (p0 >= len) return (len);
toggle_at (buf, p0);
return (len);
}
int mangle_reverse (const u8 p0, const u8 p1, u8 *buf, const int len)
{
for (int l = 0; l < len / 2; l++)
{
const int r = len - 1 - l;
mangle_switch (buf, l, r);
}
return (len);
}
int mangle_dupeword (const u8 p0, const u8 p1, u8 *buf, const int len)
{
const int out_len = len * 2;
if (out_len >= RP_PASSWORD_SIZE) return (len);
u8 *out = buf + len;
for (int i = 0; i < len; i++) *out++ = *buf++;
return (out_len);
}
int mangle_dupeword_times (const u8 p0, const u8 p1, u8 *buf, const int len)
{
const int out_len = (len * p0) + len;
if (out_len >= RP_PASSWORD_SIZE) return (len);
u8 *out = buf + len;
for (int t = 0; t < p0; t++) for (int i = 0; i < len; i++) *out++ = *buf++;
return (out_len);
}
int mangle_reflect (const u8 p0, const u8 p1, u8 *buf, const int len)
{
const int out_len = len * 2;
if (out_len >= RP_PASSWORD_SIZE) return (len);
mangle_dupeword (p0, p1, buf, len);
mangle_reverse (p0, p1, buf + len, len);
return out_len;
}
int mangle_append (const u8 p0, const u8 p1, u8 *buf, const int len)
{
const int out_len = len + 1;
if (out_len >= RP_PASSWORD_SIZE) return (len);
buf[len] = p0;
return (out_len);
}
int mangle_prepend (const u8 p0, const u8 p1, u8 *buf, const int len)
{
const int out_len = len + 1;
for (int pos = len - 1; pos >= 0; pos--)
{
buf[pos + 1] = buf[pos];
}
buf[0] = p0;
return (out_len);
}
int mangle_rotate_left (const u8 p0, const u8 p1, u8 *buf, const int len)
{
for (int l = 0, r = len - 1; r > l; r--)
{
mangle_switch (buf, l, r);
}
return (len);
}
int mangle_rotate_right (const u8 p0, const u8 p1, u8 *buf, const int len)
{
for (int l = 0, r = len - 1; l < r; l++)
{
mangle_switch (buf, l, r);
}
return (len);
}
int mangle_delete_at (const u8 p0, const u8 p1, u8 *buf, const int len)
{
if (p0 >= len) return (len);
for (int pos = p0; pos < len - 1; pos++)
{
buf[pos] = buf[pos + 1];
}
buf[len - 1] = 0;
return (len - 1);
}
int mangle_delete_first (const u8 p0, 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)
{
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)
{
if (p0 >= len) return (len);
if ((p0 + p1) > len) return (len);
for (int pos = 0; pos < p1; pos++)
{
buf[pos] = buf[p0 + pos];
}
for (int pos = p1; pos < len; pos++)
{
buf[pos] = 0;
}
return (p1);
}
int apply_rule (const u32 name, const u8 p0, const u8 p1, u8 *buf, const int in_len)
{
int out_len = in_len;
switch (name)
{
case RULE_OP_MANGLE_LREST: out_len = mangle_lrest (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_LREST_UFIRST: out_len = mangle_lrest_ufirst (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_UREST: out_len = mangle_urest (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_UREST_LFIRST: out_len = mangle_urest_lfirst (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_TREST: out_len = mangle_trest (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_TOGGLE_AT: out_len = mangle_toggle_at (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_REVERSE: out_len = mangle_reverse (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_DUPEWORD: out_len = mangle_dupeword (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_DUPEWORD_TIMES: out_len = mangle_dupeword_times (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_REFLECT: out_len = mangle_reflect (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_APPEND: out_len = mangle_append (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_PREPEND: out_len = mangle_prepend (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_ROTATE_LEFT: out_len = mangle_rotate_left (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_ROTATE_RIGHT: out_len = mangle_rotate_right (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_DELETE_FIRST: out_len = mangle_delete_first (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_DELETE_LAST: out_len = mangle_delete_last (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_DELETE_AT: out_len = mangle_delete_at (p0, p1, buf, out_len); break;
case RULE_OP_MANGLE_EXTRACT: out_len = mangle_extract (p0, p1, buf, out_len); break;
}
return out_len;
}
u32 apply_rules (__global const u32 *cmds, u32 in_buf[64], const int in_len, u32 out_buf[64])
{
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++)
{
const u32 cmd = cmds[i];
const u8 name = (cmd >> 0) & 0xff;
const u8 p0 = (cmd >> 8) & 0xff;
const u8 p1 = (cmd >> 16) & 0xff;
out_len = apply_rule (name, p0, p1, (u8 *) out_buf, out_len);
}
return (u32) out_len;
}

60
OpenCL/inc_rp.h Normal file
View File

@ -0,0 +1,60 @@
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#define RULE_OP_MANGLE_NOOP ':'
#define RULE_OP_MANGLE_LREST 'l'
#define RULE_OP_MANGLE_UREST 'u'
#define RULE_OP_MANGLE_LREST_UFIRST 'c'
#define RULE_OP_MANGLE_UREST_LFIRST 'C'
#define RULE_OP_MANGLE_TREST 't'
#define RULE_OP_MANGLE_TOGGLE_AT 'T'
#define RULE_OP_MANGLE_REVERSE 'r'
#define RULE_OP_MANGLE_DUPEWORD 'd'
#define RULE_OP_MANGLE_DUPEWORD_TIMES 'p'
#define RULE_OP_MANGLE_REFLECT 'f'
#define RULE_OP_MANGLE_ROTATE_LEFT '{'
#define RULE_OP_MANGLE_ROTATE_RIGHT '}'
#define RULE_OP_MANGLE_APPEND '$'
#define RULE_OP_MANGLE_PREPEND '^'
#define RULE_OP_MANGLE_DELETE_FIRST '['
#define RULE_OP_MANGLE_DELETE_LAST ']'
#define RULE_OP_MANGLE_DELETE_AT 'D'
#define RULE_OP_MANGLE_EXTRACT 'x'
#define RULE_OP_MANGLE_OMIT 'O'
#define RULE_OP_MANGLE_INSERT 'i'
#define RULE_OP_MANGLE_OVERSTRIKE 'o'
#define RULE_OP_MANGLE_TRUNCATE_AT '\''
#define RULE_OP_MANGLE_REPLACE 's'
#define RULE_OP_MANGLE_PURGECHAR '@'
#define RULE_OP_MANGLE_TOGGLECASE_REC 'a'
#define RULE_OP_MANGLE_DUPECHAR_FIRST 'z'
#define RULE_OP_MANGLE_DUPECHAR_LAST 'Z'
#define RULE_OP_MANGLE_DUPECHAR_ALL 'q'
#define RULE_OP_MANGLE_TITLE_SEP 'e'
#define RULE_OP_REJECT_LESS '<'
#define RULE_OP_REJECT_GREATER '>'
#define RULE_OP_REJECT_CONTAIN '!'
#define RULE_OP_REJECT_NOT_CONTAIN '/'
#define RULE_OP_REJECT_EQUAL_FIRST '('
#define RULE_OP_REJECT_EQUAL_LAST ')'
#define RULE_OP_REJECT_EQUAL_AT '='
#define RULE_OP_REJECT_CONTAINS '%'
/* hashcat only */
#define RULE_OP_MANGLE_SWITCH_FIRST 'k'
#define RULE_OP_MANGLE_SWITCH_LAST 'K'
#define RULE_OP_MANGLE_SWITCH_AT '*'
#define RULE_OP_MANGLE_CHR_SHIFTL 'L'
#define RULE_OP_MANGLE_CHR_SHIFTR 'R'
#define RULE_OP_MANGLE_CHR_INCR '+'
#define RULE_OP_MANGLE_CHR_DECR '-'
#define RULE_OP_MANGLE_REPLACE_NP1 '.'
#define RULE_OP_MANGLE_REPLACE_NM1 ','
#define RULE_OP_MANGLE_DUPEBLOCK_FIRST 'y'
#define RULE_OP_MANGLE_DUPEBLOCK_LAST 'Y'
#define RULE_OP_MANGLE_TITLE 'E'
#define RP_PASSWORD_SIZE 256

16
OpenCL/m00000_a0.cl
View File

@ -10,8 +10,8 @@
#include "inc_hash_functions.cl"
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_rp_optimized.h"
#include "inc_rp_optimized.cl"
#include "inc_rp.h"
#include "inc_rp.cl"
#include "inc_scalar.cl"
#include "inc_hash_md5.cl"
@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);

12
OpenCL/m00010_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);

12
OpenCL/m00020_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);

12
OpenCL/m00030_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update_utf16le (&ctx, w, pw_len);
md5_update_utf16le (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update_utf16le (&ctx, w, pw_len);
md5_update_utf16le (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);

12
OpenCL/m00040_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update_utf16le (&ctx, w, pw_len);
md5_update_utf16le (&ctx, out_buf, out_len);
md5_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update_utf16le (&ctx, w, pw_len);
md5_update_utf16le (&ctx, out_buf, out_len);
md5_final (&ctx);

12
OpenCL/m00050_a0.cl
View File

@ -58,11 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx;
md5_hmac_init (&ctx, w, pw_len);
md5_hmac_init (&ctx, out_buf, out_len);
md5_hmac_update (&ctx, s, salt_len);
@ -132,11 +134,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx;
md5_hmac_init (&ctx, w, pw_len);
md5_hmac_init (&ctx, out_buf, out_len);
md5_hmac_update (&ctx, s, salt_len);

12
OpenCL/m00060_a0.cl
View File

@ -62,11 +62,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx = ctx0;
md5_hmac_update (&ctx, w, pw_len);
md5_hmac_update (&ctx, out_buf, out_len);
md5_hmac_final (&ctx);
@ -138,11 +140,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx = ctx0;
md5_hmac_update (&ctx, w, pw_len);
md5_hmac_update (&ctx, out_buf, out_len);
md5_hmac_final (&ctx);

12
OpenCL/m00100_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);

12
OpenCL/m00110_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);

12
OpenCL/m00120_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);

12
OpenCL/m00130_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);

12
OpenCL/m00140_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);

12
OpenCL/m00150_a0.cl
View File

@ -58,11 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx;
sha1_hmac_init_swap (&ctx, w, pw_len);
sha1_hmac_init_swap (&ctx, out_buf, out_len);
sha1_hmac_update (&ctx, s, salt_len);
@ -132,11 +134,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx;
sha1_hmac_init_swap (&ctx, w, pw_len);
sha1_hmac_init_swap (&ctx, out_buf, out_len);
sha1_hmac_update (&ctx, s, salt_len);

12
OpenCL/m00160_a0.cl
View File

@ -62,11 +62,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx = ctx0;
sha1_hmac_update_swap (&ctx, w, pw_len);
sha1_hmac_update_swap (&ctx, out_buf, out_len);
sha1_hmac_final (&ctx);
@ -138,11 +140,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx = ctx0;
sha1_hmac_update_swap (&ctx, w, pw_len);
sha1_hmac_update_swap (&ctx, out_buf, out_len);
sha1_hmac_final (&ctx);

12
OpenCL/m00300_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);
@ -133,13 +135,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);

12
OpenCL/m00900_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update (&ctx, w, pw_len);
md4_update (&ctx, out_buf, out_len);
md4_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update (&ctx, w, pw_len);
md4_update (&ctx, out_buf, out_len);
md4_final (&ctx);

12
OpenCL/m01000_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);

12
OpenCL/m01100_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx0;
md4_init (&ctx0);
md4_update_utf16le (&ctx0, w, pw_len);
md4_update_utf16le (&ctx0, out_buf, out_len);
md4_final (&ctx0);
@ -147,13 +149,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx0;
md4_init (&ctx0);
md4_update_utf16le (&ctx0, w, pw_len);
md4_update_utf16le (&ctx0, out_buf, out_len);
md4_final (&ctx0);

12
OpenCL/m01300_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha224_ctx_t ctx;
sha224_init (&ctx);
sha224_update_swap (&ctx, w, pw_len);
sha224_update_swap (&ctx, out_buf, out_len);
sha224_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha224_ctx_t ctx;
sha224_init (&ctx);
sha224_update_swap (&ctx, w, pw_len);
sha224_update_swap (&ctx, out_buf, out_len);
sha224_final (&ctx);

12
OpenCL/m01400_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);

12
OpenCL/m01410_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_update (&ctx, s, salt_len);

12
OpenCL/m01420_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx = ctx0;
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx = ctx0;
sha256_update_swap (&ctx, w, pw_len);
sha256_update_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);

12
OpenCL/m01430_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_update (&ctx, s, salt_len);

12
OpenCL/m01440_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx = ctx0;
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx = ctx0;
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_final (&ctx);

12
OpenCL/m01450_a0.cl
View File

@ -58,11 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_hmac_ctx_t ctx;
sha256_hmac_init_swap (&ctx, w, pw_len);
sha256_hmac_init_swap (&ctx, out_buf, out_len);
sha256_hmac_update (&ctx, s, salt_len);
@ -132,11 +134,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_hmac_ctx_t ctx;
sha256_hmac_init_swap (&ctx, w, pw_len);
sha256_hmac_init_swap (&ctx, out_buf, out_len);
sha256_hmac_update (&ctx, s, salt_len);

12
OpenCL/m01460_a0.cl
View File

@ -62,11 +62,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_hmac_ctx_t ctx = ctx0;
sha256_hmac_update_swap (&ctx, w, pw_len);
sha256_hmac_update_swap (&ctx, out_buf, out_len);
sha256_hmac_final (&ctx);
@ -138,11 +140,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_hmac_ctx_t ctx = ctx0;
sha256_hmac_update_swap (&ctx, w, pw_len);
sha256_hmac_update_swap (&ctx, out_buf, out_len);
sha256_hmac_final (&ctx);

12
OpenCL/m01700_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);

12
OpenCL/m01710_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);

12
OpenCL/m01720_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx = ctx0;
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx = ctx0;
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);

12
OpenCL/m01730_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_utf16le_swap (&ctx, w, pw_len);
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_utf16le_swap (&ctx, w, pw_len);
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);

12
OpenCL/m01740_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx = ctx0;
sha512_update_utf16le_swap (&ctx, w, pw_len);
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx = ctx0;
sha512_update_utf16le_swap (&ctx, w, pw_len);
sha512_update_utf16le_swap (&ctx, out_buf, out_len);
sha512_final (&ctx);

12
OpenCL/m01750_a0.cl
View File

@ -58,11 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_hmac_ctx_t ctx;
sha512_hmac_init_swap (&ctx, w, pw_len);
sha512_hmac_init_swap (&ctx, out_buf, out_len);
sha512_hmac_update (&ctx, s, salt_len);
@ -132,11 +134,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_hmac_ctx_t ctx;
sha512_hmac_init_swap (&ctx, w, pw_len);
sha512_hmac_init_swap (&ctx, out_buf, out_len);
sha512_hmac_update (&ctx, s, salt_len);

12
OpenCL/m01760_a0.cl
View File

@ -62,11 +62,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_hmac_ctx_t ctx = ctx0;
sha512_hmac_update_swap (&ctx, w, pw_len);
sha512_hmac_update_swap (&ctx, out_buf, out_len);
sha512_hmac_final (&ctx);
@ -138,11 +140,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_hmac_ctx_t ctx = ctx0;
sha512_hmac_update_swap (&ctx, w, pw_len);
sha512_hmac_update_swap (&ctx, out_buf, out_len);
sha512_hmac_final (&ctx);

12
OpenCL/m02610_a0.cl
View File

@ -88,13 +88,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -212,13 +214,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m02810_a0.cl
View File

@ -88,13 +88,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -235,13 +237,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m03710_a0.cl
View File

@ -88,13 +88,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -225,13 +227,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m03800_a0.cl
View File

@ -64,11 +64,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);
@ -144,11 +146,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);

12
OpenCL/m03910_a0.cl
View File

@ -88,13 +88,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -235,13 +237,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m04010_a0.cl
View File

@ -83,11 +83,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1 = ctx0;
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_final (&ctx1);
@ -209,11 +211,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1 = ctx0;
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_final (&ctx1);

12
OpenCL/m04110_a0.cl
View File

@ -94,13 +94,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1;
md5_init (&ctx1);
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_update (&ctx1, s, salt_len);
@ -235,13 +237,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1;
md5_init (&ctx1);
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_update (&ctx1, s, salt_len);

12
OpenCL/m04310_a0.cl
View File

@ -88,13 +88,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -212,13 +214,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m04400_a0.cl
View File

@ -78,13 +78,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);
@ -194,13 +196,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);

12
OpenCL/m04500_a0.cl
View File

@ -77,13 +77,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);
@ -193,13 +195,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);

12
OpenCL/m04520_a0.cl
View File

@ -83,13 +83,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);
@ -214,13 +216,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);

12
OpenCL/m04700_a0.cl
View File

@ -78,13 +78,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);
@ -190,13 +192,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx0;
md5_init (&ctx0);
md5_update (&ctx0, w, pw_len);
md5_update (&ctx0, out_buf, out_len);
md5_final (&ctx0);

12
OpenCL/m04800_a0.cl
View File

@ -64,11 +64,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);
@ -144,11 +146,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_update (&ctx, s, salt_len);

12
OpenCL/m04900_a0.cl
View File

@ -64,11 +64,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);
@ -144,11 +146,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);

12
OpenCL/m05100_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);
@ -114,13 +116,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
md5_init (&ctx);
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);

12
OpenCL/m05300_a0.cl
View File

@ -47,11 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx0;
md5_hmac_init (&ctx0, w, pw_len);
md5_hmac_init (&ctx0, out_buf, out_len);
md5_hmac_update_global (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
@ -140,11 +142,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_hmac_ctx_t ctx0;
md5_hmac_init (&ctx0, w, pw_len);
md5_hmac_init (&ctx0, out_buf, out_len);
md5_hmac_update_global (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);

12
OpenCL/m05400_a0.cl
View File

@ -47,11 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx0;
sha1_hmac_init_swap (&ctx0, w, pw_len);
sha1_hmac_init_swap (&ctx0, out_buf, out_len);
sha1_hmac_update_global_swap (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);
@ -140,11 +142,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx0;
sha1_hmac_init_swap (&ctx0, w, pw_len);
sha1_hmac_init_swap (&ctx0, out_buf, out_len);
sha1_hmac_update_global_swap (&ctx0, ikepsk_bufs[digests_offset].nr_buf, ikepsk_bufs[digests_offset].nr_len);

12
OpenCL/m05500_a0.cl
View File

@ -561,13 +561,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);
@ -704,13 +706,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);

12
OpenCL/m05600_a0.cl
View File

@ -48,13 +48,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx1;
md4_init (&ctx1);
md4_update_utf16le (&ctx1, w, pw_len);
md4_update_utf16le (&ctx1, out_buf, out_len);
md4_final (&ctx1);
@ -166,13 +168,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx1;
md4_init (&ctx1);
md4_update_utf16le (&ctx1, w, pw_len);
md4_update_utf16le (&ctx1, out_buf, out_len);
md4_final (&ctx1);

12
OpenCL/m06000_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
ripemd160_ctx_t ctx;
ripemd160_init (&ctx);
ripemd160_update (&ctx, w, pw_len);
ripemd160_update (&ctx, out_buf, out_len);
ripemd160_final (&ctx);
@ -110,13 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
ripemd160_ctx_t ctx;
ripemd160_init (&ctx);
ripemd160_update (&ctx, w, pw_len);
ripemd160_update (&ctx, out_buf, out_len);
ripemd160_final (&ctx);

12
OpenCL/m06100_a0.cl
View File

@ -78,13 +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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
whirlpool_ctx_t ctx;
whirlpool_init (&ctx, s_Ch, s_Cl);
whirlpool_update (&ctx, w, pw_len);
whirlpool_update (&ctx, out_buf, out_len);
whirlpool_final (&ctx);
@ -172,13 +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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
whirlpool_ctx_t ctx;
whirlpool_init (&ctx, s_Ch, s_Cl);
whirlpool_update (&ctx, w, pw_len);
whirlpool_update (&ctx, out_buf, out_len);
whirlpool_final (&ctx);

12
OpenCL/m07000_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
/**
* pepper
@ -149,11 +151,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
/**
* pepper

12
OpenCL/m07300_a0.cl
View File

@ -47,11 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx;
sha1_hmac_init_swap (&ctx, w, pw_len);
sha1_hmac_init_swap (&ctx, out_buf, out_len);
sha1_hmac_update_global (&ctx, rakp_bufs[digests_offset].salt_buf, rakp_bufs[digests_offset].salt_len);
@ -110,11 +112,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_hmac_ctx_t ctx;
sha1_hmac_init_swap (&ctx, w, pw_len);
sha1_hmac_init_swap (&ctx, out_buf, out_len);
sha1_hmac_update_global (&ctx, rakp_bufs[digests_offset].salt_buf, rakp_bufs[digests_offset].salt_len);

12
OpenCL/m07500_a0.cl
View File

@ -320,13 +320,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);
@ -396,13 +398,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);

8
OpenCL/m08100_a0.cl
View File

@ -53,7 +53,9 @@ __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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
@ -120,7 +122,9 @@ __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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;

12
OpenCL/m08300_a0.cl
View File

@ -71,7 +71,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
@ -81,7 +83,7 @@ __kernel void m08300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
ctx1.len = 1;
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_update (&ctx1, s_pc, salt_len_pc + 1);
@ -201,7 +203,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
@ -211,7 +215,7 @@ __kernel void m08300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
ctx1.len = 1;
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_update (&ctx1, s_pc, salt_len_pc + 1);

12
OpenCL/m08400_a0.cl
View File

@ -83,13 +83,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);
@ -253,13 +255,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);

8
OpenCL/m09900_a0.cl
View File

@ -47,7 +47,9 @@ __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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;
@ -110,7 +112,9 @@ __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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx;

12
OpenCL/m10800_a0.cl
View File

@ -47,13 +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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha384_ctx_t ctx;
sha384_init (&ctx);
sha384_update (&ctx, w, pw_len);
sha384_update (&ctx, out_buf, out_len);
sha384_final (&ctx);
@ -110,13 +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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha384_ctx_t ctx;
sha384_init (&ctx);
sha384_update (&ctx, w, pw_len);
sha384_update (&ctx, out_buf, out_len);
sha384_final (&ctx);

12
OpenCL/m11000_a0.cl
View File

@ -53,11 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);
@ -120,11 +122,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx = ctx0;
md5_update (&ctx, w, pw_len);
md5_update (&ctx, out_buf, out_len);
md5_final (&ctx);

12
OpenCL/m11100_a0.cl
View File

@ -99,13 +99,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1;
md5_init (&ctx1);
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
u32 s0[4];
u32 s1[4];
@ -264,13 +266,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1;
md5_init (&ctx1);
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
u32 s0[4];
u32 s1[4];

12
OpenCL/m11200_a0.cl
View File

@ -53,13 +53,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx2;
sha1_init (&ctx2);
sha1_update_swap (&ctx2, w, pw_len);
sha1_update_swap (&ctx2, out_buf, out_len);
sha1_final (&ctx2);
@ -188,13 +190,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx2;
sha1_init (&ctx2);
sha1_update_swap (&ctx2, w, pw_len);
sha1_update_swap (&ctx2, out_buf, out_len);
sha1_final (&ctx2);

12
OpenCL/m11400_a0.cl
View File

@ -83,11 +83,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1 = ctx0;
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_final (&ctx1);
@ -200,11 +202,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md5_ctx_t ctx1 = ctx0;
md5_update (&ctx1, w, pw_len);
md5_update (&ctx1, out_buf, out_len);
md5_final (&ctx1);

12
OpenCL/m12600_a0.cl
View File

@ -93,13 +93,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);
@ -249,13 +251,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx0;
sha1_init (&ctx0);
sha1_update_swap (&ctx0, w, pw_len);
sha1_update_swap (&ctx0, out_buf, out_len);
sha1_final (&ctx0);

12
OpenCL/m13100_a0.cl
View File

@ -418,13 +418,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);
@ -485,13 +487,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
md4_ctx_t ctx;
md4_init (&ctx);
md4_update_utf16le (&ctx, w, pw_len);
md4_update_utf16le (&ctx, out_buf, out_len);
md4_final (&ctx);

12
OpenCL/m13300_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);
@ -112,13 +114,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_swap (&ctx, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);

12
OpenCL/m13500_a0.cl
View File

@ -80,11 +80,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);
@ -174,11 +176,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx = ctx0;
sha1_update_utf16le_swap (&ctx, w, pw_len);
sha1_update_utf16le_swap (&ctx, out_buf, out_len);
sha1_final (&ctx);

12
OpenCL/m13800_a0.cl
View File

@ -47,13 +47,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
@ -112,13 +114,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha256_ctx_t ctx;
sha256_init (&ctx);
sha256_update_utf16le_swap (&ctx, w, pw_len);
sha256_update_utf16le_swap (&ctx, out_buf, out_len);
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);

12
OpenCL/m13900_a0.cl
View File

@ -83,13 +83,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);
@ -253,13 +255,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1;
sha1_init (&ctx1);
sha1_update_swap (&ctx1, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
sha1_final (&ctx1);

16
OpenCL/m14400_a0.cl
View File

@ -131,7 +131,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1 = ctx0;
@ -154,7 +156,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, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -242,7 +244,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, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -397,7 +399,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx1 = ctx0;
@ -420,7 +424,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, w, pw_len);
sha1_update_swap (&ctx1, out_buf, out_len);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -508,7 +512,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, w, pw_len);
sha1_update_swap (&ctx, out_buf, out_len);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;

12
OpenCL/m15000_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);
@ -134,13 +136,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha512_ctx_t ctx;
sha512_init (&ctx);
sha512_update_swap (&ctx, w, pw_len);
sha512_update_swap (&ctx, out_buf, out_len);
sha512_update (&ctx, s, salt_len);

12
OpenCL/m15500_a0.cl
View File

@ -58,13 +58,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_utf16be_swap (&ctx, w, pw_len);
sha1_update_utf16be_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);
@ -140,13 +142,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++)
{
// todo: add rules engine
u32 out_buf[64] = { 0 };
const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w, pw_len, out_buf);
sha1_ctx_t ctx;
sha1_init (&ctx);
sha1_update_utf16be_swap (&ctx, w, pw_len);
sha1_update_utf16be_swap (&ctx, out_buf, out_len);
sha1_update (&ctx, s, salt_len);

10
src/user_options.c
View File

@ -1422,16 +1422,6 @@ void user_options_preprocess (hashcat_ctx_t *hashcat_ctx)
}
}
}
// temp
if ((user_options->rp_files_cnt > 0) || (user_options->rp_gen > 0))
{
// pure (unoptimized) kernels do not yet have support for rules, switch to opimized mode
// we can remove this after rule engine has been converted
user_options->optimized_kernel_enable = true;
}
}
void user_options_postprocess (hashcat_ctx_t *hashcat_ctx)