From 8a0d21360b7ff4eeaacb4d35b5cb1bada8d43b08 Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 11 Aug 2017 16:09:12 +0200 Subject: [PATCH] Prepare for pure kernel rule engine support --- OpenCL/inc_rp.cl | 296 ++++++++++++++++++++++++++++++++++++++++++++ OpenCL/inc_rp.h | 60 +++++++++ OpenCL/m00000_a0.cl | 16 ++- OpenCL/m00010_a0.cl | 12 +- OpenCL/m00020_a0.cl | 12 +- OpenCL/m00030_a0.cl | 12 +- OpenCL/m00040_a0.cl | 12 +- OpenCL/m00050_a0.cl | 12 +- OpenCL/m00060_a0.cl | 12 +- OpenCL/m00100_a0.cl | 12 +- OpenCL/m00110_a0.cl | 12 +- OpenCL/m00120_a0.cl | 12 +- OpenCL/m00130_a0.cl | 12 +- OpenCL/m00140_a0.cl | 12 +- OpenCL/m00150_a0.cl | 12 +- OpenCL/m00160_a0.cl | 12 +- OpenCL/m00300_a0.cl | 12 +- OpenCL/m00900_a0.cl | 12 +- OpenCL/m01000_a0.cl | 12 +- OpenCL/m01100_a0.cl | 12 +- OpenCL/m01300_a0.cl | 12 +- OpenCL/m01400_a0.cl | 12 +- OpenCL/m01410_a0.cl | 12 +- OpenCL/m01420_a0.cl | 12 +- OpenCL/m01430_a0.cl | 12 +- OpenCL/m01440_a0.cl | 12 +- OpenCL/m01450_a0.cl | 12 +- OpenCL/m01460_a0.cl | 12 +- OpenCL/m01700_a0.cl | 12 +- OpenCL/m01710_a0.cl | 12 +- OpenCL/m01720_a0.cl | 12 +- OpenCL/m01730_a0.cl | 12 +- OpenCL/m01740_a0.cl | 12 +- OpenCL/m01750_a0.cl | 12 +- OpenCL/m01760_a0.cl | 12 +- OpenCL/m02610_a0.cl | 12 +- OpenCL/m02810_a0.cl | 12 +- OpenCL/m03710_a0.cl | 12 +- OpenCL/m03800_a0.cl | 12 +- OpenCL/m03910_a0.cl | 12 +- OpenCL/m04010_a0.cl | 12 +- OpenCL/m04110_a0.cl | 12 +- OpenCL/m04310_a0.cl | 12 +- OpenCL/m04400_a0.cl | 12 +- OpenCL/m04500_a0.cl | 12 +- OpenCL/m04520_a0.cl | 12 +- OpenCL/m04700_a0.cl | 12 +- OpenCL/m04800_a0.cl | 12 +- OpenCL/m04900_a0.cl | 12 +- OpenCL/m05100_a0.cl | 12 +- OpenCL/m05300_a0.cl | 12 +- OpenCL/m05400_a0.cl | 12 +- OpenCL/m05500_a0.cl | 12 +- OpenCL/m05600_a0.cl | 12 +- OpenCL/m06000_a0.cl | 12 +- OpenCL/m06100_a0.cl | 12 +- OpenCL/m07000_a0.cl | 12 +- OpenCL/m07300_a0.cl | 12 +- OpenCL/m07500_a0.cl | 12 +- OpenCL/m08100_a0.cl | 8 +- OpenCL/m08300_a0.cl | 12 +- OpenCL/m08400_a0.cl | 12 +- OpenCL/m09900_a0.cl | 8 +- OpenCL/m10800_a0.cl | 12 +- OpenCL/m11000_a0.cl | 12 +- OpenCL/m11100_a0.cl | 12 +- OpenCL/m11200_a0.cl | 12 +- OpenCL/m11400_a0.cl | 12 +- OpenCL/m12600_a0.cl | 12 +- OpenCL/m13100_a0.cl | 12 +- OpenCL/m13300_a0.cl | 12 +- OpenCL/m13500_a0.cl | 12 +- OpenCL/m13800_a0.cl | 12 +- OpenCL/m13900_a0.cl | 12 +- OpenCL/m14400_a0.cl | 16 ++- OpenCL/m15000_a0.cl | 12 +- OpenCL/m15500_a0.cl | 12 +- src/user_options.c | 10 -- 78 files changed, 956 insertions(+), 310 deletions(-) create mode 100644 OpenCL/inc_rp.cl create mode 100644 OpenCL/inc_rp.h diff --git a/OpenCL/inc_rp.cl b/OpenCL/inc_rp.cl new file mode 100644 index 000000000..798e61ab2 --- /dev/null +++ b/OpenCL/inc_rp.cl @@ -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; +} diff --git a/OpenCL/inc_rp.h b/OpenCL/inc_rp.h new file mode 100644 index 000000000..4694fb19e --- /dev/null +++ b/OpenCL/inc_rp.h @@ -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 diff --git a/OpenCL/m00000_a0.cl b/OpenCL/m00000_a0.cl index 00d06138a..9ece53d60 100644 --- a/OpenCL/m00000_a0.cl +++ b/OpenCL/m00000_a0.cl @@ -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); diff --git a/OpenCL/m00010_a0.cl b/OpenCL/m00010_a0.cl index c6451cc24..e87d56c56 100644 --- a/OpenCL/m00010_a0.cl +++ b/OpenCL/m00010_a0.cl @@ -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); diff --git a/OpenCL/m00020_a0.cl b/OpenCL/m00020_a0.cl index 1aeb8f9c7..b0bce11fa 100644 --- a/OpenCL/m00020_a0.cl +++ b/OpenCL/m00020_a0.cl @@ -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); diff --git a/OpenCL/m00030_a0.cl b/OpenCL/m00030_a0.cl index e9c5af749..32efe1331 100644 --- a/OpenCL/m00030_a0.cl +++ b/OpenCL/m00030_a0.cl @@ -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); diff --git a/OpenCL/m00040_a0.cl b/OpenCL/m00040_a0.cl index c52d7735b..ec08a4137 100644 --- a/OpenCL/m00040_a0.cl +++ b/OpenCL/m00040_a0.cl @@ -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); diff --git a/OpenCL/m00050_a0.cl b/OpenCL/m00050_a0.cl index 09021f5bc..c0f0153c9 100644 --- a/OpenCL/m00050_a0.cl +++ b/OpenCL/m00050_a0.cl @@ -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); diff --git a/OpenCL/m00060_a0.cl b/OpenCL/m00060_a0.cl index 5a1aa73d7..26893ac01 100644 --- a/OpenCL/m00060_a0.cl +++ b/OpenCL/m00060_a0.cl @@ -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); diff --git a/OpenCL/m00100_a0.cl b/OpenCL/m00100_a0.cl index e632c593d..efc077743 100644 --- a/OpenCL/m00100_a0.cl +++ b/OpenCL/m00100_a0.cl @@ -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); diff --git a/OpenCL/m00110_a0.cl b/OpenCL/m00110_a0.cl index bd17fc935..064eb960e 100644 --- a/OpenCL/m00110_a0.cl +++ b/OpenCL/m00110_a0.cl @@ -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); diff --git a/OpenCL/m00120_a0.cl b/OpenCL/m00120_a0.cl index 9efeaee78..986840608 100644 --- a/OpenCL/m00120_a0.cl +++ b/OpenCL/m00120_a0.cl @@ -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); diff --git a/OpenCL/m00130_a0.cl b/OpenCL/m00130_a0.cl index 81bffa707..520d8a51e 100644 --- a/OpenCL/m00130_a0.cl +++ b/OpenCL/m00130_a0.cl @@ -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); diff --git a/OpenCL/m00140_a0.cl b/OpenCL/m00140_a0.cl index 24747dab6..619b311a8 100644 --- a/OpenCL/m00140_a0.cl +++ b/OpenCL/m00140_a0.cl @@ -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); diff --git a/OpenCL/m00150_a0.cl b/OpenCL/m00150_a0.cl index 39361a681..2c4afb28b 100644 --- a/OpenCL/m00150_a0.cl +++ b/OpenCL/m00150_a0.cl @@ -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); diff --git a/OpenCL/m00160_a0.cl b/OpenCL/m00160_a0.cl index 457d90020..5e3826290 100644 --- a/OpenCL/m00160_a0.cl +++ b/OpenCL/m00160_a0.cl @@ -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); diff --git a/OpenCL/m00300_a0.cl b/OpenCL/m00300_a0.cl index fbbe3c585..8e84a4f1a 100644 --- a/OpenCL/m00300_a0.cl +++ b/OpenCL/m00300_a0.cl @@ -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); diff --git a/OpenCL/m00900_a0.cl b/OpenCL/m00900_a0.cl index 9a14c10ee..5355e511d 100644 --- a/OpenCL/m00900_a0.cl +++ b/OpenCL/m00900_a0.cl @@ -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); diff --git a/OpenCL/m01000_a0.cl b/OpenCL/m01000_a0.cl index 5263d98d0..8f7342760 100644 --- a/OpenCL/m01000_a0.cl +++ b/OpenCL/m01000_a0.cl @@ -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); diff --git a/OpenCL/m01100_a0.cl b/OpenCL/m01100_a0.cl index 5ff66907e..5a20902ad 100644 --- a/OpenCL/m01100_a0.cl +++ b/OpenCL/m01100_a0.cl @@ -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); diff --git a/OpenCL/m01300_a0.cl b/OpenCL/m01300_a0.cl index 57802bd69..9a3220c58 100644 --- a/OpenCL/m01300_a0.cl +++ b/OpenCL/m01300_a0.cl @@ -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); diff --git a/OpenCL/m01400_a0.cl b/OpenCL/m01400_a0.cl index 0f7281e04..f03630b5b 100644 --- a/OpenCL/m01400_a0.cl +++ b/OpenCL/m01400_a0.cl @@ -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); diff --git a/OpenCL/m01410_a0.cl b/OpenCL/m01410_a0.cl index 1fe1eacbe..011626049 100644 --- a/OpenCL/m01410_a0.cl +++ b/OpenCL/m01410_a0.cl @@ -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); diff --git a/OpenCL/m01420_a0.cl b/OpenCL/m01420_a0.cl index 833950267..42249b287 100644 --- a/OpenCL/m01420_a0.cl +++ b/OpenCL/m01420_a0.cl @@ -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); diff --git a/OpenCL/m01430_a0.cl b/OpenCL/m01430_a0.cl index 57e79232c..18126888d 100644 --- a/OpenCL/m01430_a0.cl +++ b/OpenCL/m01430_a0.cl @@ -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); diff --git a/OpenCL/m01440_a0.cl b/OpenCL/m01440_a0.cl index a12d41519..5fb73cb3f 100644 --- a/OpenCL/m01440_a0.cl +++ b/OpenCL/m01440_a0.cl @@ -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); diff --git a/OpenCL/m01450_a0.cl b/OpenCL/m01450_a0.cl index d7b6f3e2e..c1e86674a 100644 --- a/OpenCL/m01450_a0.cl +++ b/OpenCL/m01450_a0.cl @@ -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); diff --git a/OpenCL/m01460_a0.cl b/OpenCL/m01460_a0.cl index 3f839ce50..e5ec645a8 100644 --- a/OpenCL/m01460_a0.cl +++ b/OpenCL/m01460_a0.cl @@ -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); diff --git a/OpenCL/m01700_a0.cl b/OpenCL/m01700_a0.cl index 948653460..df00d5daf 100644 --- a/OpenCL/m01700_a0.cl +++ b/OpenCL/m01700_a0.cl @@ -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); diff --git a/OpenCL/m01710_a0.cl b/OpenCL/m01710_a0.cl index 34f98e222..e6963d2ea 100644 --- a/OpenCL/m01710_a0.cl +++ b/OpenCL/m01710_a0.cl @@ -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); diff --git a/OpenCL/m01720_a0.cl b/OpenCL/m01720_a0.cl index 40830e5e7..dc4c0c60a 100644 --- a/OpenCL/m01720_a0.cl +++ b/OpenCL/m01720_a0.cl @@ -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); diff --git a/OpenCL/m01730_a0.cl b/OpenCL/m01730_a0.cl index b9b9f3d4a..02425162e 100644 --- a/OpenCL/m01730_a0.cl +++ b/OpenCL/m01730_a0.cl @@ -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); diff --git a/OpenCL/m01740_a0.cl b/OpenCL/m01740_a0.cl index 3d8b383f2..f98a85099 100644 --- a/OpenCL/m01740_a0.cl +++ b/OpenCL/m01740_a0.cl @@ -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); diff --git a/OpenCL/m01750_a0.cl b/OpenCL/m01750_a0.cl index 2d259c60c..8696e3d21 100644 --- a/OpenCL/m01750_a0.cl +++ b/OpenCL/m01750_a0.cl @@ -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); diff --git a/OpenCL/m01760_a0.cl b/OpenCL/m01760_a0.cl index 237e955ab..dcdb4707e 100644 --- a/OpenCL/m01760_a0.cl +++ b/OpenCL/m01760_a0.cl @@ -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); diff --git a/OpenCL/m02610_a0.cl b/OpenCL/m02610_a0.cl index 459b40ba4..5e60f2aec 100644 --- a/OpenCL/m02610_a0.cl +++ b/OpenCL/m02610_a0.cl @@ -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); diff --git a/OpenCL/m02810_a0.cl b/OpenCL/m02810_a0.cl index 4afbd8bf0..b77201544 100644 --- a/OpenCL/m02810_a0.cl +++ b/OpenCL/m02810_a0.cl @@ -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); diff --git a/OpenCL/m03710_a0.cl b/OpenCL/m03710_a0.cl index d6f1ed07a..b2b3630e4 100644 --- a/OpenCL/m03710_a0.cl +++ b/OpenCL/m03710_a0.cl @@ -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); diff --git a/OpenCL/m03800_a0.cl b/OpenCL/m03800_a0.cl index 75e4673ed..4c6aa9812 100644 --- a/OpenCL/m03800_a0.cl +++ b/OpenCL/m03800_a0.cl @@ -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); diff --git a/OpenCL/m03910_a0.cl b/OpenCL/m03910_a0.cl index 570587c2c..710cfdb3e 100644 --- a/OpenCL/m03910_a0.cl +++ b/OpenCL/m03910_a0.cl @@ -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); diff --git a/OpenCL/m04010_a0.cl b/OpenCL/m04010_a0.cl index 17b75c12f..640707e2b 100644 --- a/OpenCL/m04010_a0.cl +++ b/OpenCL/m04010_a0.cl @@ -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); diff --git a/OpenCL/m04110_a0.cl b/OpenCL/m04110_a0.cl index 413ed0e55..47476d3fb 100644 --- a/OpenCL/m04110_a0.cl +++ b/OpenCL/m04110_a0.cl @@ -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); diff --git a/OpenCL/m04310_a0.cl b/OpenCL/m04310_a0.cl index 04ea44233..47deba9d3 100644 --- a/OpenCL/m04310_a0.cl +++ b/OpenCL/m04310_a0.cl @@ -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); diff --git a/OpenCL/m04400_a0.cl b/OpenCL/m04400_a0.cl index 73457f6c5..47cd66e43 100644 --- a/OpenCL/m04400_a0.cl +++ b/OpenCL/m04400_a0.cl @@ -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); diff --git a/OpenCL/m04500_a0.cl b/OpenCL/m04500_a0.cl index b76522bfc..4bac012af 100644 --- a/OpenCL/m04500_a0.cl +++ b/OpenCL/m04500_a0.cl @@ -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); diff --git a/OpenCL/m04520_a0.cl b/OpenCL/m04520_a0.cl index b414de223..5fb591ec1 100644 --- a/OpenCL/m04520_a0.cl +++ b/OpenCL/m04520_a0.cl @@ -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); diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0.cl index bf2db7bb0..d90f09cdb 100644 --- a/OpenCL/m04700_a0.cl +++ b/OpenCL/m04700_a0.cl @@ -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); diff --git a/OpenCL/m04800_a0.cl b/OpenCL/m04800_a0.cl index 3e4a08686..b96ac9ebc 100644 --- a/OpenCL/m04800_a0.cl +++ b/OpenCL/m04800_a0.cl @@ -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); diff --git a/OpenCL/m04900_a0.cl b/OpenCL/m04900_a0.cl index ff71dc8b5..ac5c1c944 100644 --- a/OpenCL/m04900_a0.cl +++ b/OpenCL/m04900_a0.cl @@ -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); diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0.cl index c69ed3e78..0f0d3b4d2 100644 --- a/OpenCL/m05100_a0.cl +++ b/OpenCL/m05100_a0.cl @@ -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); diff --git a/OpenCL/m05300_a0.cl b/OpenCL/m05300_a0.cl index 53cd62281..f378cd53e 100644 --- a/OpenCL/m05300_a0.cl +++ b/OpenCL/m05300_a0.cl @@ -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); diff --git a/OpenCL/m05400_a0.cl b/OpenCL/m05400_a0.cl index 27d5be542..e55a42c07 100644 --- a/OpenCL/m05400_a0.cl +++ b/OpenCL/m05400_a0.cl @@ -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); diff --git a/OpenCL/m05500_a0.cl b/OpenCL/m05500_a0.cl index 7dccde687..c45bff4a6 100644 --- a/OpenCL/m05500_a0.cl +++ b/OpenCL/m05500_a0.cl @@ -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); diff --git a/OpenCL/m05600_a0.cl b/OpenCL/m05600_a0.cl index 8703ccd82..a3c72b430 100644 --- a/OpenCL/m05600_a0.cl +++ b/OpenCL/m05600_a0.cl @@ -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); diff --git a/OpenCL/m06000_a0.cl b/OpenCL/m06000_a0.cl index 203439658..a23547ec4 100644 --- a/OpenCL/m06000_a0.cl +++ b/OpenCL/m06000_a0.cl @@ -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); diff --git a/OpenCL/m06100_a0.cl b/OpenCL/m06100_a0.cl index ae45fa353..31b5d5ed2 100644 --- a/OpenCL/m06100_a0.cl +++ b/OpenCL/m06100_a0.cl @@ -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); diff --git a/OpenCL/m07000_a0.cl b/OpenCL/m07000_a0.cl index ccfc6d192..7c3aba5ad 100644 --- a/OpenCL/m07000_a0.cl +++ b/OpenCL/m07000_a0.cl @@ -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 diff --git a/OpenCL/m07300_a0.cl b/OpenCL/m07300_a0.cl index 5be49d782..a24566e0c 100644 --- a/OpenCL/m07300_a0.cl +++ b/OpenCL/m07300_a0.cl @@ -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); diff --git a/OpenCL/m07500_a0.cl b/OpenCL/m07500_a0.cl index e8e0878e2..45e4d9227 100644 --- a/OpenCL/m07500_a0.cl +++ b/OpenCL/m07500_a0.cl @@ -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); diff --git a/OpenCL/m08100_a0.cl b/OpenCL/m08100_a0.cl index 4c21e1d95..601a26e34 100644 --- a/OpenCL/m08100_a0.cl +++ b/OpenCL/m08100_a0.cl @@ -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; diff --git a/OpenCL/m08300_a0.cl b/OpenCL/m08300_a0.cl index 8296d0152..a88f30a51 100644 --- a/OpenCL/m08300_a0.cl +++ b/OpenCL/m08300_a0.cl @@ -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); diff --git a/OpenCL/m08400_a0.cl b/OpenCL/m08400_a0.cl index 28d218e5d..df6c976d7 100644 --- a/OpenCL/m08400_a0.cl +++ b/OpenCL/m08400_a0.cl @@ -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); diff --git a/OpenCL/m09900_a0.cl b/OpenCL/m09900_a0.cl index ab4f210d9..c2c99cc46 100644 --- a/OpenCL/m09900_a0.cl +++ b/OpenCL/m09900_a0.cl @@ -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; diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl index a49457b55..0861c9bf4 100644 --- a/OpenCL/m10800_a0.cl +++ b/OpenCL/m10800_a0.cl @@ -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); diff --git a/OpenCL/m11000_a0.cl b/OpenCL/m11000_a0.cl index 48759a64a..9c9682a6d 100644 --- a/OpenCL/m11000_a0.cl +++ b/OpenCL/m11000_a0.cl @@ -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); diff --git a/OpenCL/m11100_a0.cl b/OpenCL/m11100_a0.cl index 8dfd25225..d91b16c89 100644 --- a/OpenCL/m11100_a0.cl +++ b/OpenCL/m11100_a0.cl @@ -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]; diff --git a/OpenCL/m11200_a0.cl b/OpenCL/m11200_a0.cl index 5970d3d11..96f68ffd8 100644 --- a/OpenCL/m11200_a0.cl +++ b/OpenCL/m11200_a0.cl @@ -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); diff --git a/OpenCL/m11400_a0.cl b/OpenCL/m11400_a0.cl index 873bdc008..069228e56 100644 --- a/OpenCL/m11400_a0.cl +++ b/OpenCL/m11400_a0.cl @@ -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); diff --git a/OpenCL/m12600_a0.cl b/OpenCL/m12600_a0.cl index a45625a6c..bf050339e 100644 --- a/OpenCL/m12600_a0.cl +++ b/OpenCL/m12600_a0.cl @@ -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); diff --git a/OpenCL/m13100_a0.cl b/OpenCL/m13100_a0.cl index 1bb983243..e4bd50e1a 100644 --- a/OpenCL/m13100_a0.cl +++ b/OpenCL/m13100_a0.cl @@ -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); diff --git a/OpenCL/m13300_a0.cl b/OpenCL/m13300_a0.cl index 843c1588a..00105a20c 100644 --- a/OpenCL/m13300_a0.cl +++ b/OpenCL/m13300_a0.cl @@ -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); diff --git a/OpenCL/m13500_a0.cl b/OpenCL/m13500_a0.cl index d9acdec84..cd5f69795 100644 --- a/OpenCL/m13500_a0.cl +++ b/OpenCL/m13500_a0.cl @@ -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); diff --git a/OpenCL/m13800_a0.cl b/OpenCL/m13800_a0.cl index 7ee3d4a59..7f08f7461 100644 --- a/OpenCL/m13800_a0.cl +++ b/OpenCL/m13800_a0.cl @@ -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); diff --git a/OpenCL/m13900_a0.cl b/OpenCL/m13900_a0.cl index 910428649..8d9e9d835 100644 --- a/OpenCL/m13900_a0.cl +++ b/OpenCL/m13900_a0.cl @@ -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); diff --git a/OpenCL/m14400_a0.cl b/OpenCL/m14400_a0.cl index e4ec1a87a..92ec5b47a 100644 --- a/OpenCL/m14400_a0.cl +++ b/OpenCL/m14400_a0.cl @@ -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; diff --git a/OpenCL/m15000_a0.cl b/OpenCL/m15000_a0.cl index 65bf7cc18..298a20b42 100644 --- a/OpenCL/m15000_a0.cl +++ b/OpenCL/m15000_a0.cl @@ -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); diff --git a/OpenCL/m15500_a0.cl b/OpenCL/m15500_a0.cl index 1d736ae77..69caa3422 100644 --- a/OpenCL/m15500_a0.cl +++ b/OpenCL/m15500_a0.cl @@ -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); diff --git a/src/user_options.c b/src/user_options.c index 01c25f7e9..aaac6c4f1 100644 --- a/src/user_options.c +++ b/src/user_options.c @@ -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)