Rename function and variables containing 'unicode' to 'utf16le' because that's what's meant actually

pull/1198/merge
jsteube 7 years ago
parent 44110370d1
commit b8ad89c529

@ -201,7 +201,7 @@ inline void truncate_block (u32x w[4], const u32 len)
}
}
inline void make_unicode_be (const u32x in[4], u32x out1[4], u32x out2[4])
inline void make_utf16be (const u32x in[4], u32x out1[4], u32x out2[4])
{
#ifdef IS_NV
out2[3] = __byte_perm (in[3], 0, 0x3727);
@ -226,7 +226,7 @@ inline void make_unicode_be (const u32x in[4], u32x out1[4], u32x out2[4])
#endif
}
inline void make_unicode (const u32x in[4], u32x out1[4], u32x out2[4])
inline void make_utf16le (const u32x in[4], u32x out1[4], u32x out2[4])
{
#ifdef IS_NV
out2[3] = __byte_perm (in[3], 0, 0x7372);
@ -251,7 +251,7 @@ inline void make_unicode (const u32x in[4], u32x out1[4], u32x out2[4])
#endif
}
inline void undo_unicode_be (const u32x in1[4], const u32x in2[4], u32x out[4])
inline void undo_utf16be (const u32x in1[4], const u32x in2[4], u32x out[4])
{
#ifdef IS_NV
out[0] = __byte_perm (in1[0], in1[1], 0x4602);
@ -272,7 +272,7 @@ inline void undo_unicode_be (const u32x in1[4], const u32x in2[4], u32x out[4])
#endif
}
inline void undo_unicode (const u32x in1[4], const u32x in2[4], u32x out[4])
inline void undo_utf16le (const u32x in1[4], const u32x in2[4], u32x out[4])
{
#ifdef IS_NV
out[0] = __byte_perm (in1[0], in1[1], 0x6420);
@ -6112,7 +6112,7 @@ inline void truncate_block_S (u32 w[4], const u32 len)
}
}
inline void make_unicode_be_S (const u32 in[4], u32 out1[4], u32 out2[4])
inline void make_utf16be_S (const u32 in[4], u32 out1[4], u32 out2[4])
{
#ifdef IS_NV
out2[3] = __byte_perm_S (in[3], 0, 0x3727);
@ -6137,7 +6137,7 @@ inline void make_unicode_be_S (const u32 in[4], u32 out1[4], u32 out2[4])
#endif
}
inline void make_unicode_S (const u32 in[4], u32 out1[4], u32 out2[4])
inline void make_utf16le_S (const u32 in[4], u32 out1[4], u32 out2[4])
{
#ifdef IS_NV
out2[3] = __byte_perm_S (in[3], 0, 0x7372);
@ -6162,7 +6162,7 @@ inline void make_unicode_S (const u32 in[4], u32 out1[4], u32 out2[4])
#endif
}
inline void undo_unicode_be_S (const u32 in1[4], const u32 in2[4], u32 out[4])
inline void undo_utf16be_S (const u32 in1[4], const u32 in2[4], u32 out[4])
{
#ifdef IS_NV
out[0] = __byte_perm_S (in1[0], in1[1], 0x4602);
@ -6183,7 +6183,7 @@ inline void undo_unicode_be_S (const u32 in1[4], const u32 in2[4], u32 out[4])
#endif
}
inline void undo_unicode_S (const u32 in1[4], const u32 in2[4], u32 out[4])
inline void undo_utf16le_S (const u32 in1[4], const u32 in2[4], u32 out[4])
{
#ifdef IS_NV
out[0] = __byte_perm_S (in1[0], in1[1], 0x6420);

@ -14,8 +14,6 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00030_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -87,8 +85,8 @@ __kernel void m00030_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -311,8 +309,8 @@ __kernel void m00030_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -12,8 +12,6 @@
#include "inc_common.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00030_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -145,8 +143,8 @@ __kernel void m00030_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -429,8 +427,8 @@ __kernel void m00030_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -85,8 +85,8 @@ __kernel void m00040_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -289,8 +289,8 @@ __kernel void m00040_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -12,8 +12,6 @@
#include "inc_common.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00040_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -145,8 +143,8 @@ __kernel void m00040_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -407,8 +405,8 @@ __kernel void m00040_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -14,8 +14,6 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00130_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -87,8 +85,8 @@ __kernel void m00130_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -363,8 +361,8 @@ __kernel void m00130_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -12,8 +12,6 @@
#include "inc_common.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00130_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -145,8 +143,8 @@ __kernel void m00130_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -481,8 +479,8 @@ __kernel void m00130_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -14,8 +14,6 @@
#include "inc_rp.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00140_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -87,8 +85,8 @@ __kernel void m00140_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -343,8 +341,8 @@ __kernel void m00140_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -12,8 +12,6 @@
#include "inc_common.cl"
#include "inc_simd.cl"
// no unicode yet
__kernel void m00140_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -145,8 +143,8 @@ __kernel void m00140_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -459,8 +457,8 @@ __kernel void m00140_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -59,8 +59,8 @@ __kernel void m01000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;
@ -190,8 +190,8 @@ __kernel void m01000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;

@ -115,8 +115,8 @@ __kernel void m01000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;
@ -308,8 +308,8 @@ __kernel void m01000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;

@ -81,8 +81,8 @@ __kernel void m01100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 2 * 8;
w3[3] = 0;
@ -312,8 +312,8 @@ __kernel void m01100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 2 * 8;
w3[3] = 0;

@ -137,8 +137,8 @@ __kernel void m01100_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 2 * 8;
w3[3] = 0;
@ -430,8 +430,8 @@ __kernel void m01100_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 2 * 8;
w3[3] = 0;

@ -99,8 +99,8 @@ __kernel void m01430_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -362,8 +362,8 @@ __kernel void m01430_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -157,8 +157,8 @@ __kernel void m01430_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -480,8 +480,8 @@ __kernel void m01430_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -99,8 +99,8 @@ __kernel void m01440_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -342,8 +342,8 @@ __kernel void m01440_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -157,8 +157,8 @@ __kernel void m01440_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -458,8 +458,8 @@ __kernel void m01440_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -208,8 +208,8 @@ __kernel void m01730_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -399,8 +399,8 @@ __kernel void m01730_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -266,8 +266,8 @@ __kernel void m01730_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -517,8 +517,8 @@ __kernel void m01730_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -208,8 +208,8 @@ __kernel void m01740_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -379,8 +379,8 @@ __kernel void m01740_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -266,8 +266,8 @@ __kernel void m01740_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -495,8 +495,8 @@ __kernel void m01740_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = wordl3[2] | wordr3[2];
w3[3] = wordl3[3] | wordr3[3];
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -529,7 +529,7 @@ __kernel void m02100_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_1x4_S (w0, pw_len);
make_unicode_S (w0, w0, w1);
make_utf16le_S (w0, w0, w1);
w3[2] = pw_len * 2 * 8;

@ -581,8 +581,8 @@ __kernel void m05500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = out_len * 8 * 2;
w3_t[3] = 0;
@ -801,8 +801,8 @@ __kernel void m05500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = out_len * 8 * 2;
w3_t[3] = 0;

@ -632,8 +632,8 @@ __kernel void m05500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;
@ -910,8 +910,8 @@ __kernel void m05500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -354,8 +354,8 @@ __kernel void m05600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = out_len * 8 * 2;
w3_t[3] = 0;
@ -605,8 +605,8 @@ __kernel void m05600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = out_len * 8 * 2;
w3_t[3] = 0;

@ -410,8 +410,8 @@ __kernel void m05600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;
@ -719,8 +719,8 @@ __kernel void m05600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -447,8 +447,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -445,8 +445,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -445,8 +445,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -328,8 +328,8 @@ __kernel void m08000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
u32x w_t[16];
@ -511,8 +511,8 @@ __kernel void m08000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
u32x w_t[16];

@ -380,8 +380,8 @@ __kernel void m08000_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
u32x w_t[16];
@ -617,8 +617,8 @@ __kernel void m08000_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
u32x w2_t[4];
u32x w3_t[4];
make_unicode (w0, w0_t, w1_t);
make_unicode (w1, w2_t, w3_t);
make_utf16le (w0, w0_t, w1_t);
make_utf16le (w1, w2_t, w3_t);
u32x w_t[16];

@ -1353,8 +1353,8 @@ __kernel void m09400_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
/**
* salt

@ -1083,8 +1083,8 @@ __kernel void m09500_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
/**
* salt

@ -1123,8 +1123,8 @@ __kernel void m09600_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_4x4 (w0, w1, w2, w3, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
/**
* salt

@ -661,8 +661,8 @@ __kernel void m09700_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;
@ -862,8 +862,8 @@ __kernel void m09700_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;

@ -717,8 +717,8 @@ __kernel void m09700_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;
@ -976,8 +976,8 @@ __kernel void m09700_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;

@ -516,8 +516,8 @@ __kernel void m09720_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;
@ -634,8 +634,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = out_len * 8 * 2;
w3[3] = 0;

@ -564,8 +564,8 @@ __kernel void m09720_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;
@ -732,8 +732,8 @@ __kernel void m09720_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* md5
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 8 * 2;
w3[3] = 0;

@ -347,8 +347,8 @@ __kernel void m09800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (out_len * 2) + 16;
@ -559,8 +559,8 @@ __kernel void m09800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (out_len * 2) + 16;

@ -395,8 +395,8 @@ __kernel void m09800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (pw_len * 2) + 16;
@ -657,8 +657,8 @@ __kernel void m09800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (pw_len * 2) + 16;

@ -202,8 +202,8 @@ __kernel void m09820_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (out_len * 2) + 16;
@ -348,8 +348,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (out_len * 2) + 16;

@ -250,8 +250,8 @@ __kernel void m09820_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (pw_len * 2) + 16;
@ -446,8 +446,8 @@ __kernel void m09820_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
* sha1
*/
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_salt_len = (pw_len * 2) + 16;

@ -675,8 +675,8 @@ __kernel void m11600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
u32 pw_len = pws[gid].pw_len;
make_unicode (&pw[ 4], &pw[ 8], &pw[12]);
make_unicode (&pw[ 0], &pw[ 0], &pw[ 4]);
make_utf16le (&pw[ 4], &pw[ 8], &pw[12]);
make_utf16le (&pw[ 0], &pw[ 0], &pw[ 4]);
pw_len *= 2;

@ -391,8 +391,8 @@ __kernel void m12800_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_2x4 (w0, w1, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w3[2] = pw_len * 2 * 8;
@ -422,8 +422,8 @@ __kernel void m12800_init (__global pw_t *pws, __global const kernel_rule_t *rul
w1[3] = uint_to_hex_lower8 ((digest_md4[3] >> 16) & 255) << 0
| uint_to_hex_lower8 ((digest_md4[3] >> 24) & 255) << 16;
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
w0[0] = swap32 (w0[0]);
w0[1] = swap32 (w0[1]);

@ -617,8 +617,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -615,8 +615,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -615,8 +615,8 @@ void kerb_prepare (const u32 w0[4], const u32 w1[4], const u32 pw_len, const u32
append_0x80_2x4 (w0_t, w1_t, pw_len);
make_unicode (w1_t, w2_t, w3_t);
make_unicode (w0_t, w0_t, w1_t);
make_utf16le (w1_t, w2_t, w3_t);
make_utf16le (w0_t, w0_t, w1_t);
w3_t[2] = pw_len * 8 * 2;
w3_t[3] = 0;

@ -100,8 +100,8 @@ __kernel void m13500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;
@ -537,8 +537,8 @@ __kernel void m13500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, out_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x out_len2 = out_len * 2;

@ -150,8 +150,8 @@ __kernel void m13500_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;
@ -639,8 +639,8 @@ __kernel void m13500_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
append_0x80_2x4_VV (w0, w1, pw_len);
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
const u32x pw_len2 = pw_len * 2;

@ -556,8 +556,8 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len2 = out_len * 2;
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
u32x w[16];
@ -764,8 +764,8 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x out_len2 = out_len * 2;
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
u32x w[16];

@ -610,8 +610,8 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x pw_len2 = pw_len * 2;
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
u32x w[16];
@ -874,8 +874,8 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
const u32x pw_len2 = pw_len * 2;
make_unicode (w1, w2, w3);
make_unicode (w0, w0, w1);
make_utf16le (w1, w2, w3);
make_utf16le (w0, w0, w1);
u32x w[16];

@ -1563,8 +1563,8 @@ __kernel void m15300_init (__global pw_t *pws, __global const kernel_rule_t *rul
append_0x80_2x4_S (w0, w1, pw_len);
make_unicode_S (w1, w2, w3);
make_unicode_S (w0, w0, w1);
make_utf16le_S (w1, w2, w3);
make_utf16le_S (w0, w0, w1);
/**
* main

@ -347,8 +347,8 @@ typedef enum opti_type
typedef enum opts_type
{
OPTS_TYPE_PT_UNICODE_LE = (1ULL << 0),
OPTS_TYPE_PT_UNICODE_BE = (1ULL << 1),
OPTS_TYPE_PT_UTF16LE = (1ULL << 0),
OPTS_TYPE_PT_UTF16BE = (1ULL << 1),
OPTS_TYPE_PT_UPPER = (1ULL << 2),
OPTS_TYPE_PT_LOWER = (1ULL << 3),
OPTS_TYPE_PT_ADD01 = (1ULL << 4),
@ -361,26 +361,27 @@ typedef enum opts_type
OPTS_TYPE_PT_NEVERCRACK = (1ULL << 11), // if we want all possible results
OPTS_TYPE_PT_BITSLICE = (1ULL << 12),
OPTS_TYPE_PT_ALWAYS_ASCII = (1ULL << 13),
OPTS_TYPE_ST_UNICODE = (1ULL << 14),
OPTS_TYPE_ST_UPPER = (1ULL << 15),
OPTS_TYPE_ST_LOWER = (1ULL << 16),
OPTS_TYPE_ST_ADD01 = (1ULL << 17),
OPTS_TYPE_ST_ADD02 = (1ULL << 18),
OPTS_TYPE_ST_ADD80 = (1ULL << 19),
OPTS_TYPE_ST_ADDBITS14 = (1ULL << 20),
OPTS_TYPE_ST_ADDBITS15 = (1ULL << 21),
OPTS_TYPE_ST_GENERATE_LE = (1ULL << 22),
OPTS_TYPE_ST_GENERATE_BE = (1ULL << 23),
OPTS_TYPE_ST_HEX = (1ULL << 24),
OPTS_TYPE_ST_BASE64 = (1ULL << 25),
OPTS_TYPE_ST_HASH_MD5 = (1ULL << 26),
OPTS_TYPE_HASH_COPY = (1ULL << 27),
OPTS_TYPE_HASH_SPLIT = (1ULL << 28),
OPTS_TYPE_HOOK12 = (1ULL << 29),
OPTS_TYPE_HOOK23 = (1ULL << 30),
OPTS_TYPE_INIT2 = (1ULL << 31),
OPTS_TYPE_LOOP2 = (1ULL << 32),
OPTS_TYPE_BINARY_HASHFILE = (1ULL << 33),
OPTS_TYPE_ST_UTF16LE = (1ULL << 14),
OPTS_TYPE_ST_UTF16BE = (1ULL << 15),
OPTS_TYPE_ST_UPPER = (1ULL << 16),
OPTS_TYPE_ST_LOWER = (1ULL << 17),
OPTS_TYPE_ST_ADD01 = (1ULL << 18),
OPTS_TYPE_ST_ADD02 = (1ULL << 19),
OPTS_TYPE_ST_ADD80 = (1ULL << 20),
OPTS_TYPE_ST_ADDBITS14 = (1ULL << 21),
OPTS_TYPE_ST_ADDBITS15 = (1ULL << 22),
OPTS_TYPE_ST_GENERATE_LE = (1ULL << 23),
OPTS_TYPE_ST_GENERATE_BE = (1ULL << 24),
OPTS_TYPE_ST_HEX = (1ULL << 25),
OPTS_TYPE_ST_BASE64 = (1ULL << 26),
OPTS_TYPE_ST_HASH_MD5 = (1ULL << 27),
OPTS_TYPE_HASH_COPY = (1ULL << 28),
OPTS_TYPE_HASH_SPLIT = (1ULL << 29),
OPTS_TYPE_HOOK12 = (1ULL << 30),
OPTS_TYPE_HOOK23 = (1ULL << 31),
OPTS_TYPE_INIT2 = (1ULL << 32),
OPTS_TYPE_LOOP2 = (1ULL << 33),
OPTS_TYPE_BINARY_HASHFILE = (1ULL << 34),
} opts_type_t;

@ -79,15 +79,15 @@ static const char PA_255[] = "Unknown error";
static const char HT_00000[] = "MD5";
static const char HT_00010[] = "md5($pass.$salt)";
static const char HT_00020[] = "md5($salt.$pass)";
static const char HT_00030[] = "md5(unicode($pass).$salt)";
static const char HT_00040[] = "md5($salt.unicode($pass))";
static const char HT_00030[] = "md5(utf16le($pass).$salt)";
static const char HT_00040[] = "md5($salt.utf16le($pass))";
static const char HT_00050[] = "HMAC-MD5 (key = $pass)";
static const char HT_00060[] = "HMAC-MD5 (key = $salt)";
static const char HT_00100[] = "SHA1";
static const char HT_00110[] = "sha1($pass.$salt)";
static const char HT_00120[] = "sha1($salt.$pass)";
static const char HT_00130[] = "sha1(unicode($pass).$salt)";
static const char HT_00140[] = "sha1($salt.unicode($pass))";
static const char HT_00130[] = "sha1(utf16le($pass).$salt)";
static const char HT_00140[] = "sha1($salt.utf16le($pass))";
static const char HT_00150[] = "HMAC-SHA1 (key = $pass)";
static const char HT_00160[] = "HMAC-SHA1 (key = $salt)";
static const char HT_00200[] = "MySQL323";
@ -103,8 +103,8 @@ static const char HT_01300[] = "SHA-224";
static const char HT_01400[] = "SHA-256";
static const char HT_01410[] = "sha256($pass.$salt)";
static const char HT_01420[] = "sha256($salt.$pass)";
static const char HT_01430[] = "sha256(unicode($pass).$salt)";
static const char HT_01440[] = "sha256($salt.unicode($pass))";
static const char HT_01430[] = "sha256(utf16le($pass).$salt)";
static const char HT_01440[] = "sha256($salt.utf16le($pass))";
static const char HT_01450[] = "HMAC-SHA256 (key = $pass)";
static const char HT_01460[] = "HMAC-SHA256 (key = $salt)";
static const char HT_01500[] = "descrypt, DES (Unix), Traditional DES";
@ -112,8 +112,8 @@ static const char HT_01600[] = "Apache $apr1$ MD5, md5apr1, MD5 (APR)";
static const char HT_01700[] = "SHA-512";
static const char HT_01710[] = "sha512($pass.$salt)";
static const char HT_01720[] = "sha512($salt.$pass)";
static const char HT_01730[] = "sha512(unicode($pass).$salt)";
static const char HT_01740[] = "sha512($salt.unicode($pass))";
static const char HT_01730[] = "sha512(utf16le($pass).$salt)";
static const char HT_01740[] = "sha512($salt.utf16le($pass))";
static const char HT_01750[] = "HMAC-SHA512 (key = $pass)";
static const char HT_01760[] = "HMAC-SHA512 (key = $salt)";
static const char HT_01800[] = "sha512crypt $6$, SHA512 (Unix)";
@ -2210,7 +2210,7 @@ static u32 parse_and_store_salt (u8 *out, u8 *in, u32 salt_len, MAYBE_UNUSED con
salt_len = base64_decode (base64_to_int, (const u8 *) in, salt_len, (u8 *) tmp);
}
if (hashconfig->opts_type & OPTS_TYPE_ST_UNICODE)
if (hashconfig->opts_type & OPTS_TYPE_ST_UTF16LE)
{
if (salt_len < 20)
{
@ -2966,20 +2966,20 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
u32 SID_len = cipher_algo_pos - 1 - SID_pos;
/* maximum size of SID supported */
u8 *SID_unicode = (u8 *) hcmalloc (32 * 4);
memset (SID_unicode, 0, 32 * 4);
u8 *SID_utf16le = (u8 *) hcmalloc (32 * 4);
memset (SID_utf16le, 0, 32 * 4);
for (u32 i = 0; i < SID_len; i += 1)
{
SID_unicode[i * 2] = SID_pos[i];
SID_utf16le[i * 2] = SID_pos[i];
}
SID_unicode[(SID_len + 1) * 2] = 0x80;
SID_utf16le[(SID_len + 1) * 2] = 0x80;
/* Specific to DPAPI: needs trailing '\0' while computing hash */
dpapimk->SID_len = (SID_len + 1) * 2;
memcpy ((u8 *) dpapimk->SID, SID_unicode, 32 * 4);
memcpy ((u8 *) dpapimk->SID, SID_utf16le, 32 * 4);
for (u32 i = 0; i < 32; i++)
{
@ -2998,7 +2998,7 @@ int dpapimk_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
salt->salt_len = 16;
hcfree(SID_unicode);
hcfree(SID_utf16le);
return (PARSER_OK);
}
@ -15877,7 +15877,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
u32 salt_len = salt.salt_len;
if (opts_type & OPTS_TYPE_ST_UNICODE)
if (opts_type & OPTS_TYPE_ST_UTF16LE)
{
for (u32 i = 0, j = 0; i < salt_len; i += 1, j += 2)
{
@ -19284,7 +19284,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_INTERN;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS14;
hashconfig->kern_type = KERN_TYPE_MD5_PWUSLT;
@ -19310,7 +19310,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS14
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_MD5_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = md5s_parse_hash;
@ -19589,7 +19589,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_INTERN;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15;
hashconfig->kern_type = KERN_TYPE_SHA1_PWUSLT;
@ -19612,7 +19612,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_PT_UPPER
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15
@ -19637,7 +19637,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15
| OPTS_TYPE_ST_HEX;
@ -19661,7 +19661,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15;
hashconfig->kern_type = KERN_TYPE_SHA1_PWUSLT;
@ -19686,7 +19686,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS15
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_SHA1_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_4_5;
hashconfig->parse_func = sha1s_parse_hash;
@ -19709,7 +19709,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS15
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_BASE64;
hashconfig->kern_type = KERN_TYPE_SHA1_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_4_5;
@ -19885,7 +19885,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS14
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_MD4_PWU;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = md4_parse_hash;
@ -19909,9 +19909,9 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS14
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_UNICODE
| OPTS_TYPE_ST_UTF16LE
| OPTS_TYPE_ST_LOWER;
hashconfig->kern_type = KERN_TYPE_MD44_PWUSLT;
hashconfig->dgst_size = DGST_SIZE_4_4;
@ -20063,7 +20063,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_INTERN;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15;
hashconfig->kern_type = KERN_TYPE_SHA256_PWUSLT;
@ -20088,7 +20088,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS15
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_SHA256_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_4_8;
hashconfig->parse_func = sha256s_parse_hash;
@ -20111,7 +20111,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS15
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_BASE64;
hashconfig->kern_type = KERN_TYPE_SHA256_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_4_8;
@ -20312,7 +20312,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_INTERN;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15;
hashconfig->kern_type = KERN_TYPE_SHA512_PWSLTU;
@ -20336,7 +20336,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_ADDBITS15
| OPTS_TYPE_ST_HEX;
@ -20363,7 +20363,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS15
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_SHA512_SLTPWU;
hashconfig->dgst_size = DGST_SIZE_8_8;
hashconfig->parse_func = sha512s_parse_hash;
@ -20450,7 +20450,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_ST_LOWER
| OPTS_TYPE_ST_UNICODE;
| OPTS_TYPE_ST_UTF16LE;
hashconfig->kern_type = KERN_TYPE_DCC2;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = dcc2_parse_hash;
@ -21043,7 +21043,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS14
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_HEX;
hashconfig->kern_type = KERN_TYPE_NETNTLMv1;
hashconfig->dgst_size = DGST_SIZE_4_4;
@ -21062,7 +21062,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_ADDBITS14
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_NETNTLMv2;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = netntlmv2_parse_hash;
@ -21572,7 +21572,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_ST_ADD80
| OPTS_TYPE_ST_HEX;
hashconfig->kern_type = KERN_TYPE_SYBASEASE;
@ -21840,7 +21840,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_OLDOFFICE01;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = oldoffice01_parse_hash;
@ -21875,7 +21875,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_PT_NEVERCRACK;
hashconfig->kern_type = KERN_TYPE_OLDOFFICE01CM2;
hashconfig->dgst_size = DGST_SIZE_4_4;
@ -21894,7 +21894,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_OLDOFFICE34;
hashconfig->dgst_size = DGST_SIZE_4_4;
hashconfig->parse_func = oldoffice34_parse_hash;
@ -21928,7 +21928,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_ADD80
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_PT_NEVERCRACK;
hashconfig->kern_type = KERN_TYPE_OLDOFFICE34CM2;
hashconfig->dgst_size = DGST_SIZE_4_4;
@ -22576,7 +22576,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE
| OPTS_TYPE_PT_UTF16LE
| OPTS_TYPE_PT_ADD80;
hashconfig->kern_type = KERN_TYPE_PSTOKEN;
hashconfig->dgst_size = DGST_SIZE_4_5;
@ -22883,7 +22883,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_BE
| OPTS_TYPE_PT_UNICODE_LE;
| OPTS_TYPE_PT_UTF16LE;
hashconfig->kern_type = KERN_TYPE_WIN8PHONE;
hashconfig->dgst_size = DGST_SIZE_4_8;
hashconfig->parse_func = win8phone_parse_hash;
@ -23426,7 +23426,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->pw_max = PW_MAX;
if ((hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_BE))
if ((hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE))
{
hashconfig->pw_max = PW_MAX / 2;
}

@ -47,7 +47,7 @@ static void mp_css_split_cnt (hashcat_ctx_t *hashcat_ctx, const u32 css_cnt_orig
}
else
{
if ((hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_BE))
if ((hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE))
{
if (css_cnt_orig == 8 || css_cnt_orig == 10)
{
@ -122,50 +122,50 @@ static int mp_css_append_salt (hashcat_ctx_t *hashcat_ctx, salt_t *salt_buf)
return 0;
}
static int mp_css_unicode_expand_le (hashcat_ctx_t *hashcat_ctx)
static int mp_css_utf16le_expand (hashcat_ctx_t *hashcat_ctx)
{
mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx;
u32 css_cnt_unicode = mask_ctx->css_cnt * 2;
u32 css_cnt_utf16le = mask_ctx->css_cnt * 2;
cs_t *css_buf_unicode = (cs_t *) hccalloc (css_cnt_unicode, sizeof (cs_t));
cs_t *css_buf_utf16le = (cs_t *) hccalloc (css_cnt_utf16le, sizeof (cs_t));
for (u32 i = 0, j = 0; i < mask_ctx->css_cnt; i += 1, j += 2)
{
memcpy (&css_buf_unicode[j + 0], &mask_ctx->css_buf[i], sizeof (cs_t));
memcpy (&css_buf_utf16le[j + 0], &mask_ctx->css_buf[i], sizeof (cs_t));
css_buf_unicode[j + 1].cs_buf[0] = 0;
css_buf_unicode[j + 1].cs_len = 1;
css_buf_utf16le[j + 1].cs_buf[0] = 0;
css_buf_utf16le[j + 1].cs_len = 1;
}
hcfree (mask_ctx->css_buf);
mask_ctx->css_buf = css_buf_unicode;
mask_ctx->css_cnt = css_cnt_unicode;
mask_ctx->css_buf = css_buf_utf16le;
mask_ctx->css_cnt = css_cnt_utf16le;
return 0;
}
static int mp_css_unicode_expand_be (hashcat_ctx_t *hashcat_ctx)
static int mp_css_utf16be_expand (hashcat_ctx_t *hashcat_ctx)
{
mask_ctx_t *mask_ctx = hashcat_ctx->mask_ctx;
u32 css_cnt_unicode = mask_ctx->css_cnt * 2;
u32 css_cnt_utf16be = mask_ctx->css_cnt * 2;
cs_t *css_buf_unicode = (cs_t *) hccalloc (css_cnt_unicode, sizeof (cs_t));
cs_t *css_buf_utf16be = (cs_t *) hccalloc (css_cnt_utf16be, sizeof (cs_t));
for (u32 i = 0, j = 0; i < mask_ctx->css_cnt; i += 1, j += 2)
{
css_buf_unicode[j + 0].cs_buf[0] = 0;
css_buf_unicode[j + 0].cs_len = 1;
css_buf_utf16be[j + 0].cs_buf[0] = 0;
css_buf_utf16be[j + 0].cs_len = 1;
memcpy (&css_buf_unicode[j + 1], &mask_ctx->css_buf[i], sizeof (cs_t));
memcpy (&css_buf_utf16be[j + 1], &mask_ctx->css_buf[i], sizeof (cs_t));
}
hcfree (mask_ctx->css_buf);
mask_ctx->css_buf = css_buf_unicode;
mask_ctx->css_cnt = css_cnt_unicode;
mask_ctx->css_buf = css_buf_utf16be;
mask_ctx->css_cnt = css_cnt_utf16be;
return 0;
}
@ -1189,15 +1189,15 @@ int mask_ctx_update_loop (hashcat_ctx_t *hashcat_ctx)
return 0;
}
if (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_LE)
if (hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE)
{
const int rc = mp_css_unicode_expand_le (hashcat_ctx);
const int rc = mp_css_utf16le_expand (hashcat_ctx);
if (rc == -1) return -1;
}
else if (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_BE)
else if (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE)
{
const int rc = mp_css_unicode_expand_be (hashcat_ctx);
const int rc = mp_css_utf16be_expand (hashcat_ctx);
if (rc == -1) return -1;
}

@ -182,7 +182,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
}
}
if (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_LE)
if (hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE)
{
for (int i = 0, j = 0; i < plain_len; i += 2, j += 1)
{
@ -191,7 +191,7 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl
plain_len = plain_len / 2;
}
else if (hashconfig->opts_type & OPTS_TYPE_PT_UNICODE_BE)
else if (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE)
{
for (int i = 1, j = 0; i < plain_len; i += 2, j += 1)
{

@ -134,8 +134,8 @@ static const char *USAGE_BIG[] =
" 11800 | GOST R 34.11-2012 (Streebog) 512-bit | Raw Hash",
" 10 | md5($pass.$salt) | Raw Hash, Salted and/or Iterated",
" 20 | md5($salt.$pass) | Raw Hash, Salted and/or Iterated",
" 30 | md5(unicode($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 40 | md5($salt.unicode($pass)) | Raw Hash, Salted and/or Iterated",
" 30 | md5(utf16le($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 40 | md5($salt.utf16le($pass)) | Raw Hash, Salted and/or Iterated",
" 3800 | md5($salt.$pass.$salt) | Raw Hash, Salted and/or Iterated",
" 3710 | md5($salt.md5($pass)) | Raw Hash, Salted and/or Iterated",
" 4010 | md5($salt.md5($salt.$pass)) | Raw Hash, Salted and/or Iterated",
@ -146,8 +146,8 @@ static const char *USAGE_BIG[] =
" 4400 | md5(sha1($pass)) | Raw Hash, Salted and/or Iterated",
" 110 | sha1($pass.$salt) | Raw Hash, Salted and/or Iterated",
" 120 | sha1($salt.$pass) | Raw Hash, Salted and/or Iterated",
" 130 | sha1(unicode($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 140 | sha1($salt.unicode($pass)) | Raw Hash, Salted and/or Iterated",
" 130 | sha1(utf16le($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 140 | sha1($salt.utf16le($pass)) | Raw Hash, Salted and/or Iterated",
" 4500 | sha1(sha1($pass)) | Raw Hash, Salted and/or Iterated",
" 4520 | sha1($salt.sha1($pass)) | Raw Hash, Salted and/or Iterated",
" 4700 | sha1(md5($pass)) | Raw Hash, Salted and/or Iterated",
@ -155,12 +155,12 @@ static const char *USAGE_BIG[] =
" 14400 | sha1(CX) | Raw Hash, Salted and/or Iterated",
" 1410 | sha256($pass.$salt) | Raw Hash, Salted and/or Iterated",
" 1420 | sha256($salt.$pass) | Raw Hash, Salted and/or Iterated",
" 1430 | sha256(unicode($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 1440 | sha256($salt.unicode($pass)) | Raw Hash, Salted and/or Iterated",
" 1430 | sha256(utf16le($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 1440 | sha256($salt.utf16le($pass)) | Raw Hash, Salted and/or Iterated",
" 1710 | sha512($pass.$salt) | Raw Hash, Salted and/or Iterated",
" 1720 | sha512($salt.$pass) | Raw Hash, Salted and/or Iterated",
" 1730 | sha512(unicode($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 1740 | sha512($salt.unicode($pass)) | Raw Hash, Salted and/or Iterated",
" 1730 | sha512(utf16le($pass).$salt) | Raw Hash, Salted and/or Iterated",
" 1740 | sha512($salt.utf16le($pass)) | Raw Hash, Salted and/or Iterated",
" 50 | HMAC-MD5 (key = $pass) | Raw Hash, Authenticated",
" 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated",
" 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated",

@ -50,7 +50,7 @@ my $MAX_LEN = 55;
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15600, 15700, 99999);
my %is_unicode = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800);
my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800);
my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
my %allow_long_salt = map { $_ => 1 } qw (2500 4520 4521 5500 5600 7100 7200 7300 9400 9500 9600 9700 9800 10400 10500 10600 10700 1100 11000 11200 11300 11400 11600 12600 13500 13800 15000);
@ -8044,11 +8044,11 @@ END_CODE
}
elsif ($mode == 13800)
{
my $word_buf_unicode = encode ("UTF-16LE", $word_buf);
my $word_buf_utf16le = encode ("UTF-16LE", $word_buf);
my $salt_buf_bin = pack ("H*", $salt_buf);
$hash_buf = sha256_hex ($word_buf_unicode . $salt_buf_bin);
$hash_buf = sha256_hex ($word_buf_utf16le . $salt_buf_bin);
$tmp_hash = sprintf ("%s:%s", $hash_buf, $salt_buf);
}
@ -8692,7 +8692,7 @@ sub rnd
$max = 15 if ($mode == 2410);
if ($is_unicode{$mode})
if ($is_utf16le{$mode})
{
if (! $allow_long_salt{$mode})
{

Loading…
Cancel
Save