mirror of https://github.com/hashcat/hashcat.git
- Added hash-mode 19000 QNX /etc/shadow (MD5) - Added hash-mode 19100 QNX /etc/shadow (SHA256) - Added hash-mode 19200 QNX /etc/shadow (SHA512) Implements #35pull/1940/head
jsteube
5 years ago
parent
58390e34b7
commit
87c24200da
@ -0,0 +1,93 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
//#define NEW_SIMD_CODE
|
||||
|
||||
#include "inc_vendor.cl"
|
||||
#include "inc_hash_constants.h"
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
typedef struct qnx_md5_tmp
|
||||
{
|
||||
md5_ctx_t md5_ctx;
|
||||
|
||||
} qnx_md5_tmp_t;
|
||||
|
||||
__kernel void m19000_init (KERN_ATTR_TMPS (qnx_md5_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* init
|
||||
*/
|
||||
|
||||
md5_ctx_t md5_ctx;
|
||||
|
||||
md5_init (&md5_ctx);
|
||||
|
||||
md5_update_global (&md5_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
|
||||
tmps[gid].md5_ctx = md5_ctx;
|
||||
}
|
||||
|
||||
__kernel void m19000_loop (KERN_ATTR_TMPS (qnx_md5_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
md5_ctx_t md5_ctx = tmps[gid].md5_ctx;
|
||||
|
||||
for (u32 i = 0; i < loop_cnt; i++)
|
||||
{
|
||||
md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
}
|
||||
|
||||
tmps[gid].md5_ctx = md5_ctx;
|
||||
}
|
||||
|
||||
__kernel void m19000_comp (KERN_ATTR_TMPS (qnx_md5_tmp_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
md5_ctx_t md5_ctx = tmps[gid].md5_ctx;
|
||||
|
||||
md5_final (&md5_ctx);
|
||||
|
||||
const u32 r0 = md5_ctx.h[0];
|
||||
const u32 r1 = md5_ctx.h[1];
|
||||
const u32 r2 = md5_ctx.h[2];
|
||||
const u32 r3 = md5_ctx.h[3];
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#include COMPARE_M
|
||||
}
|
||||
@ -0,0 +1,93 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
//#define NEW_SIMD_CODE
|
||||
|
||||
#include "inc_vendor.cl"
|
||||
#include "inc_hash_constants.h"
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha256.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
typedef struct qnx_sha256_tmp
|
||||
{
|
||||
sha256_ctx_t sha256_ctx;
|
||||
|
||||
} qnx_sha256_tmp_t;
|
||||
|
||||
__kernel void m19100_init (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* init
|
||||
*/
|
||||
|
||||
sha256_ctx_t sha256_ctx;
|
||||
|
||||
sha256_init (&sha256_ctx);
|
||||
|
||||
sha256_update_global_swap (&sha256_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha256_update_global_swap (&sha256_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
|
||||
tmps[gid].sha256_ctx = sha256_ctx;
|
||||
}
|
||||
|
||||
__kernel void m19100_loop (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
sha256_ctx_t sha256_ctx = tmps[gid].sha256_ctx;
|
||||
|
||||
for (u32 i = 0; i < loop_cnt; i++)
|
||||
{
|
||||
sha256_update_global_swap (&sha256_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
}
|
||||
|
||||
tmps[gid].sha256_ctx = sha256_ctx;
|
||||
}
|
||||
|
||||
__kernel void m19100_comp (KERN_ATTR_TMPS (qnx_sha256_tmp_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
sha256_ctx_t sha256_ctx = tmps[gid].sha256_ctx;
|
||||
|
||||
sha256_final (&sha256_ctx);
|
||||
|
||||
const u32 r0 = swap32_S (sha256_ctx.h[0]);
|
||||
const u32 r1 = swap32_S (sha256_ctx.h[1]);
|
||||
const u32 r2 = swap32_S (sha256_ctx.h[2]);
|
||||
const u32 r3 = swap32_S (sha256_ctx.h[3]);
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#include COMPARE_M
|
||||
}
|
||||
@ -0,0 +1,455 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
//#define NEW_SIMD_CODE
|
||||
|
||||
#include "inc_vendor.cl"
|
||||
#include "inc_hash_constants.h"
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_sha512.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
typedef struct qnx_sha512_tmp
|
||||
{
|
||||
sha512_ctx_t sha512_ctx;
|
||||
|
||||
u32 sav; // to trigger sha512 bug
|
||||
|
||||
} qnx_sha512_tmp_t;
|
||||
|
||||
DECLSPEC u32 sha512_update_128_qnxbug (sha512_ctx_t *ctx, u32 *w0, u32 *w1, u32 *w2, u32 *w3, u32 *w4, u32 *w5, u32 *w6, u32 *w7, const int len, u32 sav)
|
||||
{
|
||||
MAYBE_VOLATILE const int pos = ctx->len & 127;
|
||||
|
||||
ctx->len += len;
|
||||
|
||||
if ((pos + len) < 128)
|
||||
{
|
||||
switch_buffer_by_offset_8x4_be_S (w0, w1, w2, w3, w4, w5, w6, w7, pos);
|
||||
|
||||
ctx->w0[0] |= w0[0];
|
||||
ctx->w0[1] |= w0[1];
|
||||
ctx->w0[2] |= w0[2];
|
||||
ctx->w0[3] |= w0[3];
|
||||
ctx->w1[0] |= w1[0];
|
||||
ctx->w1[1] |= w1[1];
|
||||
ctx->w1[2] |= w1[2];
|
||||
ctx->w1[3] |= w1[3];
|
||||
ctx->w2[0] |= w2[0];
|
||||
ctx->w2[1] |= w2[1];
|
||||
ctx->w2[2] |= w2[2];
|
||||
ctx->w2[3] |= w2[3];
|
||||
ctx->w3[0] |= w3[0];
|
||||
ctx->w3[1] |= w3[1];
|
||||
ctx->w3[2] |= w3[2];
|
||||
ctx->w3[3] |= w3[3];
|
||||
ctx->w4[0] |= w4[0];
|
||||
ctx->w4[1] |= w4[1];
|
||||
ctx->w4[2] |= w4[2];
|
||||
ctx->w4[3] |= w4[3];
|
||||
ctx->w5[0] |= w5[0];
|
||||
ctx->w5[1] |= w5[1];
|
||||
ctx->w5[2] |= w5[2];
|
||||
ctx->w5[3] |= w5[3];
|
||||
ctx->w6[0] |= w6[0];
|
||||
ctx->w6[1] |= w6[1];
|
||||
ctx->w6[2] |= w6[2];
|
||||
ctx->w6[3] |= w6[3];
|
||||
ctx->w7[0] |= w7[0];
|
||||
ctx->w7[1] |= w7[1];
|
||||
ctx->w7[2] |= w7[2];
|
||||
ctx->w7[3] |= w7[3];
|
||||
}
|
||||
else
|
||||
{
|
||||
u32 c0[4] = { 0 };
|
||||
u32 c1[4] = { 0 };
|
||||
u32 c2[4] = { 0 };
|
||||
u32 c3[4] = { 0 };
|
||||
u32 c4[4] = { 0 };
|
||||
u32 c5[4] = { 0 };
|
||||
u32 c6[4] = { 0 };
|
||||
u32 c7[4] = { 0 };
|
||||
|
||||
switch_buffer_by_offset_8x4_carry_be_S (w0, w1, w2, w3, w4, w5, w6, w7, c0, c1, c2, c3, c4, c5, c6, c7, pos);
|
||||
|
||||
ctx->w0[0] |= w0[0];
|
||||
ctx->w0[1] |= w0[1];
|
||||
ctx->w0[2] |= w0[2];
|
||||
ctx->w0[3] |= w0[3];
|
||||
ctx->w1[0] |= w1[0];
|
||||
ctx->w1[1] |= w1[1];
|
||||
ctx->w1[2] |= w1[2];
|
||||
ctx->w1[3] |= w1[3];
|
||||
ctx->w2[0] |= w2[0];
|
||||
ctx->w2[1] |= w2[1];
|
||||
ctx->w2[2] |= w2[2];
|
||||
ctx->w2[3] |= w2[3];
|
||||
ctx->w3[0] |= w3[0];
|
||||
ctx->w3[1] |= w3[1];
|
||||
ctx->w3[2] |= w3[2];
|
||||
ctx->w3[3] |= w3[3];
|
||||
ctx->w4[0] |= w4[0];
|
||||
ctx->w4[1] |= w4[1];
|
||||
ctx->w4[2] |= w4[2];
|
||||
ctx->w4[3] |= w4[3];
|
||||
ctx->w5[0] |= w5[0];
|
||||
ctx->w5[1] |= w5[1];
|
||||
ctx->w5[2] |= w5[2];
|
||||
ctx->w5[3] |= w5[3];
|
||||
ctx->w6[0] |= w6[0];
|
||||
ctx->w6[1] |= w6[1];
|
||||
ctx->w6[2] |= w6[2];
|
||||
ctx->w6[3] |= w6[3];
|
||||
ctx->w7[0] |= w7[0];
|
||||
ctx->w7[1] |= w7[1];
|
||||
ctx->w7[2] |= w7[2];
|
||||
ctx->w7[3] |= w7[3];
|
||||
|
||||
sav = ctx->w7[1];
|
||||
|
||||
sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
|
||||
|
||||
ctx->w0[0] = c0[0];
|
||||
ctx->w0[1] = c0[1];
|
||||
ctx->w0[2] = c0[2];
|
||||
ctx->w0[3] = c0[3];
|
||||
ctx->w1[0] = c1[0];
|
||||
ctx->w1[1] = c1[1];
|
||||
ctx->w1[2] = c1[2];
|
||||
ctx->w1[3] = c1[3];
|
||||
ctx->w2[0] = c2[0];
|
||||
ctx->w2[1] = c2[1];
|
||||
ctx->w2[2] = c2[2];
|
||||
ctx->w2[3] = c2[3];
|
||||
ctx->w3[0] = c3[0];
|
||||
ctx->w3[1] = c3[1];
|
||||
ctx->w3[2] = c3[2];
|
||||
ctx->w3[3] = c3[3];
|
||||
ctx->w4[0] = c4[0];
|
||||
ctx->w4[1] = c4[1];
|
||||
ctx->w4[2] = c4[2];
|
||||
ctx->w4[3] = c4[3];
|
||||
ctx->w5[0] = c5[0];
|
||||
ctx->w5[1] = c5[1];
|
||||
ctx->w5[2] = c5[2];
|
||||
ctx->w5[3] = c5[3];
|
||||
ctx->w6[0] = c6[0];
|
||||
ctx->w6[1] = c6[1];
|
||||
ctx->w6[2] = c6[2];
|
||||
ctx->w6[3] = c6[3];
|
||||
ctx->w7[0] = c7[0];
|
||||
ctx->w7[1] = c7[1];
|
||||
ctx->w7[2] = c7[2];
|
||||
ctx->w7[3] = c7[3];
|
||||
}
|
||||
|
||||
return sav;
|
||||
}
|
||||
|
||||
DECLSPEC u32 sha512_update_global_swap_qnxbug (sha512_ctx_t *ctx, const __global u32 *w, const int len, u32 sav)
|
||||
{
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
u32 w4[4];
|
||||
u32 w5[4];
|
||||
u32 w6[4];
|
||||
u32 w7[4];
|
||||
|
||||
int pos1;
|
||||
int pos4;
|
||||
|
||||
for (pos1 = 0, pos4 = 0; pos1 < len - 128; pos1 += 128, pos4 += 32)
|
||||
{
|
||||
w0[0] = w[pos4 + 0];
|
||||
w0[1] = w[pos4 + 1];
|
||||
w0[2] = w[pos4 + 2];
|
||||
w0[3] = w[pos4 + 3];
|
||||
w1[0] = w[pos4 + 4];
|
||||
w1[1] = w[pos4 + 5];
|
||||
w1[2] = w[pos4 + 6];
|
||||
w1[3] = w[pos4 + 7];
|
||||
w2[0] = w[pos4 + 8];
|
||||
w2[1] = w[pos4 + 9];
|
||||
w2[2] = w[pos4 + 10];
|
||||
w2[3] = w[pos4 + 11];
|
||||
w3[0] = w[pos4 + 12];
|
||||
w3[1] = w[pos4 + 13];
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
w4[0] = w[pos4 + 16];
|
||||
w4[1] = w[pos4 + 17];
|
||||
w4[2] = w[pos4 + 18];
|
||||
w4[3] = w[pos4 + 19];
|
||||
w5[0] = w[pos4 + 20];
|
||||
w5[1] = w[pos4 + 21];
|
||||
w5[2] = w[pos4 + 22];
|
||||
w5[3] = w[pos4 + 23];
|
||||
w6[0] = w[pos4 + 24];
|
||||
w6[1] = w[pos4 + 25];
|
||||
w6[2] = w[pos4 + 26];
|
||||
w6[3] = w[pos4 + 27];
|
||||
w7[0] = w[pos4 + 28];
|
||||
w7[1] = w[pos4 + 29];
|
||||
w7[2] = w[pos4 + 30];
|
||||
w7[3] = w[pos4 + 31];
|
||||
|
||||
w0[0] = swap32_S (w0[0]);
|
||||
w0[1] = swap32_S (w0[1]);
|
||||
w0[2] = swap32_S (w0[2]);
|
||||
w0[3] = swap32_S (w0[3]);
|
||||
w1[0] = swap32_S (w1[0]);
|
||||
w1[1] = swap32_S (w1[1]);
|
||||
w1[2] = swap32_S (w1[2]);
|
||||
w1[3] = swap32_S (w1[3]);
|
||||
w2[0] = swap32_S (w2[0]);
|
||||
w2[1] = swap32_S (w2[1]);
|
||||
w2[2] = swap32_S (w2[2]);
|
||||
w2[3] = swap32_S (w2[3]);
|
||||
w3[0] = swap32_S (w3[0]);
|
||||
w3[1] = swap32_S (w3[1]);
|
||||
w3[2] = swap32_S (w3[2]);
|
||||
w3[3] = swap32_S (w3[3]);
|
||||
w4[0] = swap32_S (w4[0]);
|
||||
w4[1] = swap32_S (w4[1]);
|
||||
w4[2] = swap32_S (w4[2]);
|
||||
w4[3] = swap32_S (w4[3]);
|
||||
w5[0] = swap32_S (w5[0]);
|
||||
w5[1] = swap32_S (w5[1]);
|
||||
w5[2] = swap32_S (w5[2]);
|
||||
w5[3] = swap32_S (w5[3]);
|
||||
w6[0] = swap32_S (w6[0]);
|
||||
w6[1] = swap32_S (w6[1]);
|
||||
w6[2] = swap32_S (w6[2]);
|
||||
w6[3] = swap32_S (w6[3]);
|
||||
w7[0] = swap32_S (w7[0]);
|
||||
w7[1] = swap32_S (w7[1]);
|
||||
w7[2] = swap32_S (w7[2]);
|
||||
w7[3] = swap32_S (w7[3]);
|
||||
|
||||
sav = sha512_update_128_qnxbug (ctx, w0, w1, w2, w3, w4, w5, w6, w7, 128, sav);
|
||||
}
|
||||
|
||||
w0[0] = w[pos4 + 0];
|
||||
w0[1] = w[pos4 + 1];
|
||||
w0[2] = w[pos4 + 2];
|
||||
w0[3] = w[pos4 + 3];
|
||||
w1[0] = w[pos4 + 4];
|
||||
w1[1] = w[pos4 + 5];
|
||||
w1[2] = w[pos4 + 6];
|
||||
w1[3] = w[pos4 + 7];
|
||||
w2[0] = w[pos4 + 8];
|
||||
w2[1] = w[pos4 + 9];
|
||||
w2[2] = w[pos4 + 10];
|
||||
w2[3] = w[pos4 + 11];
|
||||
w3[0] = w[pos4 + 12];
|
||||
w3[1] = w[pos4 + 13];
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
w4[0] = w[pos4 + 16];
|
||||
w4[1] = w[pos4 + 17];
|
||||
w4[2] = w[pos4 + 18];
|
||||
w4[3] = w[pos4 + 19];
|
||||
w5[0] = w[pos4 + 20];
|
||||
w5[1] = w[pos4 + 21];
|
||||
w5[2] = w[pos4 + 22];
|
||||
w5[3] = w[pos4 + 23];
|
||||
w6[0] = w[pos4 + 24];
|
||||
w6[1] = w[pos4 + 25];
|
||||
w6[2] = w[pos4 + 26];
|
||||
w6[3] = w[pos4 + 27];
|
||||
w7[0] = w[pos4 + 28];
|
||||
w7[1] = w[pos4 + 29];
|
||||
w7[2] = w[pos4 + 30];
|
||||
w7[3] = w[pos4 + 31];
|
||||
|
||||
w0[0] = swap32_S (w0[0]);
|
||||
w0[1] = swap32_S (w0[1]);
|
||||
w0[2] = swap32_S (w0[2]);
|
||||
w0[3] = swap32_S (w0[3]);
|
||||
w1[0] = swap32_S (w1[0]);
|
||||
w1[1] = swap32_S (w1[1]);
|
||||
w1[2] = swap32_S (w1[2]);
|
||||
w1[3] = swap32_S (w1[3]);
|
||||
w2[0] = swap32_S (w2[0]);
|
||||
w2[1] = swap32_S (w2[1]);
|
||||
w2[2] = swap32_S (w2[2]);
|
||||
w2[3] = swap32_S (w2[3]);
|
||||
w3[0] = swap32_S (w3[0]);
|
||||
w3[1] = swap32_S (w3[1]);
|
||||
w3[2] = swap32_S (w3[2]);
|
||||
w3[3] = swap32_S (w3[3]);
|
||||
w4[0] = swap32_S (w4[0]);
|
||||
w4[1] = swap32_S (w4[1]);
|
||||
w4[2] = swap32_S (w4[2]);
|
||||
w4[3] = swap32_S (w4[3]);
|
||||
w5[0] = swap32_S (w5[0]);
|
||||
w5[1] = swap32_S (w5[1]);
|
||||
w5[2] = swap32_S (w5[2]);
|
||||
w5[3] = swap32_S (w5[3]);
|
||||
w6[0] = swap32_S (w6[0]);
|
||||
w6[1] = swap32_S (w6[1]);
|
||||
w6[2] = swap32_S (w6[2]);
|
||||
w6[3] = swap32_S (w6[3]);
|
||||
w7[0] = swap32_S (w7[0]);
|
||||
w7[1] = swap32_S (w7[1]);
|
||||
w7[2] = swap32_S (w7[2]);
|
||||
w7[3] = swap32_S (w7[3]);
|
||||
|
||||
sav = sha512_update_128_qnxbug (ctx, w0, w1, w2, w3, w4, w5, w6, w7, len - pos1, sav);
|
||||
|
||||
return sav;
|
||||
}
|
||||
|
||||
DECLSPEC void sha512_final_qnxbug (sha512_ctx_t *ctx, u32 sav)
|
||||
{
|
||||
MAYBE_VOLATILE const int pos = ctx->len & 127;
|
||||
|
||||
append_0x80_8x4_S (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, pos ^ 3);
|
||||
|
||||
if (pos >= 112)
|
||||
{
|
||||
sav = ctx->w7[1];
|
||||
|
||||
sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
|
||||
|
||||
ctx->w0[0] = 0;
|
||||
ctx->w0[1] = 0;
|
||||
ctx->w0[2] = 0;
|
||||
ctx->w0[3] = 0;
|
||||
ctx->w1[0] = 0;
|
||||
ctx->w1[1] = 0;
|
||||
ctx->w1[2] = 0;
|
||||
ctx->w1[3] = 0;
|
||||
ctx->w2[0] = 0;
|
||||
ctx->w2[1] = 0;
|
||||
ctx->w2[2] = 0;
|
||||
ctx->w2[3] = 0;
|
||||
ctx->w3[0] = 0;
|
||||
ctx->w3[1] = 0;
|
||||
ctx->w3[2] = 0;
|
||||
ctx->w3[3] = 0;
|
||||
ctx->w4[0] = 0;
|
||||
ctx->w4[1] = 0;
|
||||
ctx->w4[2] = 0;
|
||||
ctx->w4[3] = 0;
|
||||
ctx->w5[0] = 0;
|
||||
ctx->w5[1] = 0;
|
||||
ctx->w5[2] = 0;
|
||||
ctx->w5[3] = 0;
|
||||
ctx->w6[0] = 0;
|
||||
ctx->w6[1] = 0;
|
||||
ctx->w6[2] = 0;
|
||||
ctx->w6[3] = 0;
|
||||
ctx->w7[0] = 0;
|
||||
ctx->w7[1] = 0;
|
||||
ctx->w7[2] = 0;
|
||||
ctx->w7[3] = 0;
|
||||
}
|
||||
|
||||
ctx->w7[1] = sav;
|
||||
ctx->w7[2] = 0;
|
||||
ctx->w7[3] = ctx->len * 8;
|
||||
|
||||
sha512_transform (ctx->w0, ctx->w1, ctx->w2, ctx->w3, ctx->w4, ctx->w5, ctx->w6, ctx->w7, ctx->h);
|
||||
}
|
||||
|
||||
__kernel void m19200_init (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* init
|
||||
*/
|
||||
|
||||
sha512_ctx_t sha512_ctx;
|
||||
|
||||
sha512_init (&sha512_ctx);
|
||||
|
||||
sha512_update_global_swap (&sha512_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha512_update_global_swap (&sha512_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
|
||||
tmps[gid].sha512_ctx = sha512_ctx;
|
||||
tmps[gid].sav = 0;
|
||||
}
|
||||
|
||||
__kernel void m19200_loop (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
sha512_ctx_t sha512_ctx = tmps[gid].sha512_ctx;
|
||||
u32 sav = tmps[gid].sav;
|
||||
|
||||
for (u32 i = 0; i < loop_cnt; i++)
|
||||
{
|
||||
sav = sha512_update_global_swap_qnxbug (&sha512_ctx, pws[gid].i, pws[gid].pw_len, sav);
|
||||
}
|
||||
|
||||
tmps[gid].sha512_ctx = sha512_ctx;
|
||||
tmps[gid].sav = sav;
|
||||
}
|
||||
|
||||
__kernel void m19200_comp (KERN_ATTR_TMPS (qnx_sha512_tmp_t))
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 gid = get_global_id (0);
|
||||
const u64 lid = get_local_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
sha512_ctx_t sha512_ctx = tmps[gid].sha512_ctx;
|
||||
|
||||
sha512_final (&sha512_ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (swap64_S (sha512_ctx.h[0]));
|
||||
const u32 r1 = h32_from_64_S (swap64_S (sha512_ctx.h[0]));
|
||||
const u32 r2 = l32_from_64_S (swap64_S (sha512_ctx.h[1]));
|
||||
const u32 r3 = h32_from_64_S (swap64_S (sha512_ctx.h[1]));
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#include COMPARE_M
|
||||
|
||||
// we should also handle the buggy qnx sha512 implementation
|
||||
// see https://github.com/magnumripper/JohnTheRipper/blob/bleeding-jumbo/src/sha2.c#L578-L595
|
||||
|
||||
sha512_ctx_t sha512_ctx2 = tmps[gid].sha512_ctx;
|
||||
u32 sav = tmps[gid].sav;
|
||||
|
||||
if (sha512_ctx2.len >= 116)
|
||||
{
|
||||
sha512_final_qnxbug (&sha512_ctx2, sav);
|
||||
|
||||
const u32 r0 = l32_from_64_S (swap64_S (sha512_ctx2.h[0]));
|
||||
const u32 r1 = h32_from_64_S (swap64_S (sha512_ctx2.h[0]));
|
||||
const u32 r2 = l32_from_64_S (swap64_S (sha512_ctx2.h[1]));
|
||||
const u32 r3 = h32_from_64_S (swap64_S (sha512_ctx2.h[1]));
|
||||
|
||||
#include COMPARE_M
|
||||
}
|
||||
}
|
||||
@ -0,0 +1,257 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#include "common.h"
|
||||
#include "types.h"
|
||||
#include "modules.h"
|
||||
#include "bitops.h"
|
||||
#include "convert.h"
|
||||
#include "shared.h"
|
||||
#include "cpu_md5.h"
|
||||
#include "memory.h"
|
||||
|
||||
static const u32 ATTACK_EXEC = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
static const u32 DGST_POS0 = 0;
|
||||
static const u32 DGST_POS1 = 1;
|
||||
static const u32 DGST_POS2 = 2;
|
||||
static const u32 DGST_POS3 = 3;
|
||||
static const u32 DGST_SIZE = DGST_SIZE_8_8;
|
||||
static const u32 HASH_CATEGORY = HASH_CATEGORY_OS;
|
||||
static const char *HASH_NAME = "QNX /etc/shadow (MD5)";
|
||||
static const u64 KERN_TYPE = 19000;
|
||||
static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE;
|
||||
static const u64 OPTS_TYPE = OPTS_TYPE_PT_GENERATE_LE;
|
||||
static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED;
|
||||
static const char *ST_PASS = "hashcat";
|
||||
static const char *ST_HASH = "@m@75f6f129f9c9e77b6b1b78f791ed764a@8741857532330050";
|
||||
|
||||
u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; }
|
||||
u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; }
|
||||
u32 module_dgst_pos1 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1; }
|
||||
u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2; }
|
||||
u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3; }
|
||||
u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE; }
|
||||
u32 module_hash_category (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY; }
|
||||
const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME; }
|
||||
u64 module_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE; }
|
||||
u32 module_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE; }
|
||||
u64 module_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE; }
|
||||
u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE; }
|
||||
const char *module_st_hash (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH; }
|
||||
const char *module_st_pass (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS; }
|
||||
|
||||
typedef struct md5_ctx
|
||||
{
|
||||
u32 h[4];
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
int len;
|
||||
|
||||
} md5_ctx_t;
|
||||
|
||||
typedef struct qnx_md5_tmp
|
||||
{
|
||||
md5_ctx_t md5_ctx;
|
||||
|
||||
} qnx_md5_tmp_t;
|
||||
|
||||
static const int ROUNDS_QNX = 1000;
|
||||
|
||||
u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
|
||||
{
|
||||
const u64 tmp_size = (const u64) sizeof (qnx_md5_tmp_t);
|
||||
|
||||
return tmp_size;
|
||||
}
|
||||
|
||||
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
|
||||
{
|
||||
u32 *digest = (u32 *) digest_buf;
|
||||
|
||||
token_t token;
|
||||
|
||||
token.token_cnt = 4;
|
||||
|
||||
token.sep[0] = '@';
|
||||
token.len_min[0] = 0;
|
||||
token.len_max[0] = 0;
|
||||
token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[1] = '@';
|
||||
token.len_min[1] = 1;
|
||||
token.len_max[1] = 8;
|
||||
token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[2] = '@';
|
||||
token.len_min[2] = 32;
|
||||
token.len_max[2] = 128;
|
||||
token.attr[2] = TOKEN_ATTR_VERIFY_LENGTH
|
||||
| TOKEN_ATTR_VERIFY_HEX;
|
||||
|
||||
token.sep[3] = '@';
|
||||
token.len_min[3] = 8;
|
||||
token.len_max[3] = 16;
|
||||
token.attr[3] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
|
||||
|
||||
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
|
||||
|
||||
// check hash type
|
||||
|
||||
if (token.buf[1][0] != 'm') return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
// check iter
|
||||
|
||||
u32 iter = ROUNDS_QNX;
|
||||
|
||||
if (token.len[1] > 1)
|
||||
{
|
||||
if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
|
||||
}
|
||||
|
||||
// iter++; the additinal round is added in the init kernel
|
||||
|
||||
salt->salt_iter = iter;
|
||||
|
||||
// digest
|
||||
|
||||
if (token.len[2] != 32) return (PARSER_HASH_LENGTH);
|
||||
|
||||
digest[0] = hex_to_u32 ((const u8 *) token.buf[2] + 0);
|
||||
digest[1] = hex_to_u32 ((const u8 *) token.buf[2] + 8);
|
||||
digest[2] = hex_to_u32 ((const u8 *) token.buf[2] + 16);
|
||||
digest[3] = hex_to_u32 ((const u8 *) token.buf[2] + 24);
|
||||
|
||||
// salt
|
||||
|
||||
if ((token.len[3] == 8) || (token.len[3] == 16))
|
||||
{
|
||||
memcpy (salt->salt_buf, token.buf[3], token.len[3]);
|
||||
|
||||
salt->salt_len = token.len[3];
|
||||
}
|
||||
else
|
||||
{
|
||||
return (PARSER_SALT_LENGTH);
|
||||
}
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
|
||||
{
|
||||
const u32 *digest = (const u32 *) digest_buf;
|
||||
|
||||
int line_len = 0;
|
||||
|
||||
const int iter = salt->salt_iter;
|
||||
|
||||
if (iter == ROUNDS_QNX)
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 'm';
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
else
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 'm';
|
||||
line_buf[line_len++] = ',';
|
||||
|
||||
line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
|
||||
u32_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 8;
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
|
||||
memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
|
||||
|
||||
line_len += salt->salt_len;
|
||||
|
||||
line_buf[line_len] = 0;
|
||||
|
||||
return line_len;
|
||||
}
|
||||
|
||||
void module_init (module_ctx_t *module_ctx)
|
||||
{
|
||||
module_ctx->module_context_size = MODULE_CONTEXT_SIZE_CURRENT;
|
||||
module_ctx->module_interface_version = MODULE_INTERFACE_VERSION_CURRENT;
|
||||
|
||||
module_ctx->module_attack_exec = module_attack_exec;
|
||||
module_ctx->module_benchmark_esalt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_hook_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_mask = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_build_plain_postprocess = MODULE_DEFAULT;
|
||||
module_ctx->module_deep_comp_kernel = MODULE_DEFAULT;
|
||||
module_ctx->module_dgst_pos0 = module_dgst_pos0;
|
||||
module_ctx->module_dgst_pos1 = module_dgst_pos1;
|
||||
module_ctx->module_dgst_pos2 = module_dgst_pos2;
|
||||
module_ctx->module_dgst_pos3 = module_dgst_pos3;
|
||||
module_ctx->module_dgst_size = module_dgst_size;
|
||||
module_ctx->module_dictstat_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_save = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_outfile = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_zero_hash = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode = module_hash_decode;
|
||||
module_ctx->module_hash_encode_status = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_encode = module_hash_encode;
|
||||
module_ctx->module_hash_init_selftest = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_mode = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_category = module_hash_category;
|
||||
module_ctx->module_hash_name = module_hash_name;
|
||||
module_ctx->module_hlfmt_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_hook12 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook23 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_salt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_size = MODULE_DEFAULT;
|
||||
module_ctx->module_jit_build_options = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kern_type = module_kern_type;
|
||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||
module_ctx->module_opti_type = module_opti_type;
|
||||
module_ctx->module_opts_type = module_opts_type;
|
||||
module_ctx->module_outfile_check_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_outfile_check_nocomp = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_keep_all_hashes = MODULE_DEFAULT;
|
||||
module_ctx->module_pwdump_column = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_max = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_max = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_type = module_salt_type;
|
||||
module_ctx->module_separator = MODULE_DEFAULT;
|
||||
module_ctx->module_st_hash = module_st_hash;
|
||||
module_ctx->module_st_pass = module_st_pass;
|
||||
module_ctx->module_tmp_size = module_tmp_size;
|
||||
module_ctx->module_unstable_warning = MODULE_DEFAULT;
|
||||
module_ctx->module_warmup_disable = MODULE_DEFAULT;
|
||||
}
|
||||
@ -0,0 +1,265 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#include "common.h"
|
||||
#include "types.h"
|
||||
#include "modules.h"
|
||||
#include "bitops.h"
|
||||
#include "convert.h"
|
||||
#include "shared.h"
|
||||
#include "cpu_md5.h"
|
||||
#include "memory.h"
|
||||
|
||||
static const u32 ATTACK_EXEC = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
static const u32 DGST_POS0 = 0;
|
||||
static const u32 DGST_POS1 = 1;
|
||||
static const u32 DGST_POS2 = 2;
|
||||
static const u32 DGST_POS3 = 3;
|
||||
static const u32 DGST_SIZE = DGST_SIZE_8_8;
|
||||
static const u32 HASH_CATEGORY = HASH_CATEGORY_OS;
|
||||
static const char *HASH_NAME = "QNX /etc/shadow (SHA256)";
|
||||
static const u64 KERN_TYPE = 19100;
|
||||
static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE;
|
||||
static const u64 OPTS_TYPE = OPTS_TYPE_PT_GENERATE_LE;
|
||||
static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED;
|
||||
static const char *ST_PASS = "hashcat";
|
||||
static const char *ST_HASH = "@s@0b365cab7e17ee1e7e1a90078501cc1aa85888d6da34e2f5b04f5c614b882a93@5498317092471604";
|
||||
|
||||
u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; }
|
||||
u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; }
|
||||
u32 module_dgst_pos1 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1; }
|
||||
u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2; }
|
||||
u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3; }
|
||||
u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE; }
|
||||
u32 module_hash_category (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY; }
|
||||
const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME; }
|
||||
u64 module_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE; }
|
||||
u32 module_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE; }
|
||||
u64 module_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE; }
|
||||
u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE; }
|
||||
const char *module_st_hash (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH; }
|
||||
const char *module_st_pass (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS; }
|
||||
|
||||
typedef struct sha256_ctx
|
||||
{
|
||||
u32 h[8];
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
int len;
|
||||
|
||||
} sha256_ctx_t;
|
||||
|
||||
typedef struct qnx_sha256_tmp
|
||||
{
|
||||
sha256_ctx_t sha256_ctx;
|
||||
|
||||
} qnx_sha256_tmp_t;
|
||||
|
||||
static const int ROUNDS_QNX = 1000;
|
||||
|
||||
u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
|
||||
{
|
||||
const u64 tmp_size = (const u64) sizeof (qnx_sha256_tmp_t);
|
||||
|
||||
return tmp_size;
|
||||
}
|
||||
|
||||
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
|
||||
{
|
||||
u32 *digest = (u32 *) digest_buf;
|
||||
|
||||
token_t token;
|
||||
|
||||
token.token_cnt = 4;
|
||||
|
||||
token.sep[0] = '@';
|
||||
token.len_min[0] = 0;
|
||||
token.len_max[0] = 0;
|
||||
token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[1] = '@';
|
||||
token.len_min[1] = 1;
|
||||
token.len_max[1] = 8;
|
||||
token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[2] = '@';
|
||||
token.len_min[2] = 32;
|
||||
token.len_max[2] = 128;
|
||||
token.attr[2] = TOKEN_ATTR_VERIFY_LENGTH
|
||||
| TOKEN_ATTR_VERIFY_HEX;
|
||||
|
||||
token.sep[3] = '@';
|
||||
token.len_min[3] = 8;
|
||||
token.len_max[3] = 16;
|
||||
token.attr[3] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
|
||||
|
||||
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
|
||||
|
||||
// check hash type
|
||||
|
||||
if (token.buf[1][0] != 's') return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
// check iter
|
||||
|
||||
u32 iter = ROUNDS_QNX;
|
||||
|
||||
if (token.len[1] > 1)
|
||||
{
|
||||
if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
|
||||
}
|
||||
|
||||
// iter++; the additinal round is added in the init kernel
|
||||
|
||||
salt->salt_iter = iter;
|
||||
|
||||
// digest
|
||||
|
||||
if (token.len[2] != 64) return (PARSER_HASH_LENGTH);
|
||||
|
||||
digest[0] = hex_to_u32 ((const u8 *) token.buf[2] + 0);
|
||||
digest[1] = hex_to_u32 ((const u8 *) token.buf[2] + 8);
|
||||
digest[2] = hex_to_u32 ((const u8 *) token.buf[2] + 16);
|
||||
digest[3] = hex_to_u32 ((const u8 *) token.buf[2] + 24);
|
||||
digest[4] = hex_to_u32 ((const u8 *) token.buf[2] + 32);
|
||||
digest[5] = hex_to_u32 ((const u8 *) token.buf[2] + 40);
|
||||
digest[6] = hex_to_u32 ((const u8 *) token.buf[2] + 48);
|
||||
digest[7] = hex_to_u32 ((const u8 *) token.buf[2] + 56);
|
||||
|
||||
// salt
|
||||
|
||||
if ((token.len[3] == 8) || (token.len[3] == 16))
|
||||
{
|
||||
memcpy (salt->salt_buf, token.buf[3], token.len[3]);
|
||||
|
||||
salt->salt_len = token.len[3];
|
||||
}
|
||||
else
|
||||
{
|
||||
return (PARSER_SALT_LENGTH);
|
||||
}
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
|
||||
{
|
||||
const u32 *digest = (const u32 *) digest_buf;
|
||||
|
||||
int line_len = 0;
|
||||
|
||||
const int iter = salt->salt_iter;
|
||||
|
||||
if (iter == ROUNDS_QNX)
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 's';
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
else
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 's';
|
||||
line_buf[line_len++] = ',';
|
||||
|
||||
line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
|
||||
u32_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[4], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[5], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[6], (u8 *) line_buf + line_len); line_len += 8;
|
||||
u32_to_hex (digest[7], (u8 *) line_buf + line_len); line_len += 8;
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
|
||||
memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
|
||||
|
||||
line_len += salt->salt_len;
|
||||
|
||||
line_buf[line_len] = 0;
|
||||
|
||||
return line_len;
|
||||
}
|
||||
|
||||
void module_init (module_ctx_t *module_ctx)
|
||||
{
|
||||
module_ctx->module_context_size = MODULE_CONTEXT_SIZE_CURRENT;
|
||||
module_ctx->module_interface_version = MODULE_INTERFACE_VERSION_CURRENT;
|
||||
|
||||
module_ctx->module_attack_exec = module_attack_exec;
|
||||
module_ctx->module_benchmark_esalt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_hook_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_mask = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_build_plain_postprocess = MODULE_DEFAULT;
|
||||
module_ctx->module_deep_comp_kernel = MODULE_DEFAULT;
|
||||
module_ctx->module_dgst_pos0 = module_dgst_pos0;
|
||||
module_ctx->module_dgst_pos1 = module_dgst_pos1;
|
||||
module_ctx->module_dgst_pos2 = module_dgst_pos2;
|
||||
module_ctx->module_dgst_pos3 = module_dgst_pos3;
|
||||
module_ctx->module_dgst_size = module_dgst_size;
|
||||
module_ctx->module_dictstat_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_save = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_outfile = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_zero_hash = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode = module_hash_decode;
|
||||
module_ctx->module_hash_encode_status = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_encode = module_hash_encode;
|
||||
module_ctx->module_hash_init_selftest = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_mode = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_category = module_hash_category;
|
||||
module_ctx->module_hash_name = module_hash_name;
|
||||
module_ctx->module_hlfmt_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_hook12 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook23 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_salt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_size = MODULE_DEFAULT;
|
||||
module_ctx->module_jit_build_options = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kern_type = module_kern_type;
|
||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||
module_ctx->module_opti_type = module_opti_type;
|
||||
module_ctx->module_opts_type = module_opts_type;
|
||||
module_ctx->module_outfile_check_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_outfile_check_nocomp = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_keep_all_hashes = MODULE_DEFAULT;
|
||||
module_ctx->module_pwdump_column = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_max = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_max = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_type = module_salt_type;
|
||||
module_ctx->module_separator = MODULE_DEFAULT;
|
||||
module_ctx->module_st_hash = module_st_hash;
|
||||
module_ctx->module_st_pass = module_st_pass;
|
||||
module_ctx->module_tmp_size = module_tmp_size;
|
||||
module_ctx->module_unstable_warning = MODULE_DEFAULT;
|
||||
module_ctx->module_warmup_disable = MODULE_DEFAULT;
|
||||
}
|
||||
@ -0,0 +1,271 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#include "common.h"
|
||||
#include "types.h"
|
||||
#include "modules.h"
|
||||
#include "bitops.h"
|
||||
#include "convert.h"
|
||||
#include "shared.h"
|
||||
#include "cpu_md5.h"
|
||||
#include "memory.h"
|
||||
|
||||
static const u32 ATTACK_EXEC = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
static const u32 DGST_POS0 = 0;
|
||||
static const u32 DGST_POS1 = 1;
|
||||
static const u32 DGST_POS2 = 2;
|
||||
static const u32 DGST_POS3 = 3;
|
||||
static const u32 DGST_SIZE = DGST_SIZE_8_8;
|
||||
static const u32 HASH_CATEGORY = HASH_CATEGORY_OS;
|
||||
static const char *HASH_NAME = "QNX /etc/shadow (SHA512)";
|
||||
static const u64 KERN_TYPE = 19200;
|
||||
static const u32 OPTI_TYPE = OPTI_TYPE_ZERO_BYTE;
|
||||
static const u64 OPTS_TYPE = OPTS_TYPE_PT_GENERATE_LE;
|
||||
static const u32 SALT_TYPE = SALT_TYPE_EMBEDDED;
|
||||
static const char *ST_PASS = "hashcat";
|
||||
static const char *ST_HASH = "@S@715df9e94c097805dd1e13c6a40f331d02ce589765a2100ec7435e76b978d5efc364ce10870780622cee003c9951bd92ec1020c924b124cfff7e0fa1f73e3672@2257314490293159";
|
||||
|
||||
u32 module_attack_exec (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ATTACK_EXEC; }
|
||||
u32 module_dgst_pos0 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS0; }
|
||||
u32 module_dgst_pos1 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS1; }
|
||||
u32 module_dgst_pos2 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS2; }
|
||||
u32 module_dgst_pos3 (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_POS3; }
|
||||
u32 module_dgst_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return DGST_SIZE; }
|
||||
u32 module_hash_category (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_CATEGORY; }
|
||||
const char *module_hash_name (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return HASH_NAME; }
|
||||
u64 module_kern_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return KERN_TYPE; }
|
||||
u32 module_opti_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTI_TYPE; }
|
||||
u64 module_opts_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return OPTS_TYPE; }
|
||||
u32 module_salt_type (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return SALT_TYPE; }
|
||||
const char *module_st_hash (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_HASH; }
|
||||
const char *module_st_pass (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra) { return ST_PASS; }
|
||||
|
||||
typedef struct sha512_ctx
|
||||
{
|
||||
u64 h[8];
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
u32 w4[4];
|
||||
u32 w5[4];
|
||||
u32 w6[4];
|
||||
u32 w7[4];
|
||||
|
||||
int len;
|
||||
|
||||
} sha512_ctx_t;
|
||||
|
||||
typedef struct qnx_sha512_tmp
|
||||
{
|
||||
sha512_ctx_t sha512_ctx;
|
||||
|
||||
u32 sav; // to trigger sha512 bug
|
||||
|
||||
} qnx_sha512_tmp_t;
|
||||
|
||||
static const int ROUNDS_QNX = 1000;
|
||||
|
||||
u64 module_tmp_size (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const user_options_t *user_options, MAYBE_UNUSED const user_options_extra_t *user_options_extra)
|
||||
{
|
||||
const u64 tmp_size = (const u64) sizeof (qnx_sha512_tmp_t);
|
||||
|
||||
return tmp_size;
|
||||
}
|
||||
|
||||
int module_hash_decode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED void *digest_buf, MAYBE_UNUSED salt_t *salt, MAYBE_UNUSED void *esalt_buf, MAYBE_UNUSED void *hook_salt_buf, MAYBE_UNUSED hashinfo_t *hash_info, const char *line_buf, MAYBE_UNUSED const int line_len)
|
||||
{
|
||||
u64 *digest = (u64 *) digest_buf;
|
||||
|
||||
token_t token;
|
||||
|
||||
token.token_cnt = 4;
|
||||
|
||||
token.sep[0] = '@';
|
||||
token.len_min[0] = 0;
|
||||
token.len_max[0] = 0;
|
||||
token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[1] = '@';
|
||||
token.len_min[1] = 1;
|
||||
token.len_max[1] = 8;
|
||||
token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
token.sep[2] = '@';
|
||||
token.len_min[2] = 32;
|
||||
token.len_max[2] = 128;
|
||||
token.attr[2] = TOKEN_ATTR_VERIFY_LENGTH
|
||||
| TOKEN_ATTR_VERIFY_HEX;
|
||||
|
||||
token.sep[3] = '@';
|
||||
token.len_min[3] = 8;
|
||||
token.len_max[3] = 16;
|
||||
token.attr[3] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
const int rc_tokenizer = input_tokenizer ((const u8 *) line_buf, line_len, &token);
|
||||
|
||||
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
|
||||
|
||||
// check hash type
|
||||
|
||||
if (token.buf[1][0] != 'S') return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
// check iter
|
||||
|
||||
u32 iter = ROUNDS_QNX;
|
||||
|
||||
if (token.len[1] > 1)
|
||||
{
|
||||
if (token.buf[1][1] != ',') return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
iter = hc_strtoul ((const char *) token.buf[1] + 2, NULL, 10);
|
||||
}
|
||||
|
||||
// iter++; the additinal round is added in the init kernel
|
||||
|
||||
salt->salt_iter = iter;
|
||||
|
||||
// digest
|
||||
|
||||
if (token.len[2] != 128) return (PARSER_HASH_LENGTH);
|
||||
|
||||
digest[0] = hex_to_u64 ((const u8 *) token.buf[2] + 0);
|
||||
digest[1] = hex_to_u64 ((const u8 *) token.buf[2] + 16);
|
||||
digest[2] = hex_to_u64 ((const u8 *) token.buf[2] + 32);
|
||||
digest[3] = hex_to_u64 ((const u8 *) token.buf[2] + 48);
|
||||
digest[4] = hex_to_u64 ((const u8 *) token.buf[2] + 64);
|
||||
digest[5] = hex_to_u64 ((const u8 *) token.buf[2] + 80);
|
||||
digest[6] = hex_to_u64 ((const u8 *) token.buf[2] + 96);
|
||||
digest[7] = hex_to_u64 ((const u8 *) token.buf[2] + 112);
|
||||
|
||||
// salt
|
||||
|
||||
if ((token.len[3] == 8) || (token.len[3] == 16))
|
||||
{
|
||||
memcpy (salt->salt_buf, token.buf[3], token.len[3]);
|
||||
|
||||
salt->salt_len = token.len[3];
|
||||
}
|
||||
else
|
||||
{
|
||||
return (PARSER_SALT_LENGTH);
|
||||
}
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int module_hash_encode (MAYBE_UNUSED const hashconfig_t *hashconfig, MAYBE_UNUSED const void *digest_buf, MAYBE_UNUSED const salt_t *salt, MAYBE_UNUSED const void *esalt_buf, MAYBE_UNUSED const void *hook_salt_buf, MAYBE_UNUSED const hashinfo_t *hash_info, char *line_buf, MAYBE_UNUSED const int line_size)
|
||||
{
|
||||
const u64 *digest = (const u64 *) digest_buf;
|
||||
|
||||
int line_len = 0;
|
||||
|
||||
const int iter = salt->salt_iter;
|
||||
|
||||
if (iter == ROUNDS_QNX)
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 'S';
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
else
|
||||
{
|
||||
line_buf[line_len++] = '@';
|
||||
line_buf[line_len++] = 'S';
|
||||
line_buf[line_len++] = ',';
|
||||
|
||||
line_len += snprintf (line_buf + line_len, line_size - line_len, "%u", iter);
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
}
|
||||
|
||||
u64_to_hex (digest[0], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[1], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[2], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[3], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[4], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[5], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[6], (u8 *) line_buf + line_len); line_len += 16;
|
||||
u64_to_hex (digest[7], (u8 *) line_buf + line_len); line_len += 16;
|
||||
|
||||
line_buf[line_len++] = '@';
|
||||
|
||||
memcpy (line_buf + line_len, salt->salt_buf, salt->salt_len);
|
||||
|
||||
line_len += salt->salt_len;
|
||||
|
||||
line_buf[line_len] = 0;
|
||||
|
||||
return line_len;
|
||||
}
|
||||
|
||||
void module_init (module_ctx_t *module_ctx)
|
||||
{
|
||||
module_ctx->module_context_size = MODULE_CONTEXT_SIZE_CURRENT;
|
||||
module_ctx->module_interface_version = MODULE_INTERFACE_VERSION_CURRENT;
|
||||
|
||||
module_ctx->module_attack_exec = module_attack_exec;
|
||||
module_ctx->module_benchmark_esalt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_hook_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_mask = MODULE_DEFAULT;
|
||||
module_ctx->module_benchmark_salt = MODULE_DEFAULT;
|
||||
module_ctx->module_build_plain_postprocess = MODULE_DEFAULT;
|
||||
module_ctx->module_deep_comp_kernel = MODULE_DEFAULT;
|
||||
module_ctx->module_dgst_pos0 = module_dgst_pos0;
|
||||
module_ctx->module_dgst_pos1 = module_dgst_pos1;
|
||||
module_ctx->module_dgst_pos2 = module_dgst_pos2;
|
||||
module_ctx->module_dgst_pos3 = module_dgst_pos3;
|
||||
module_ctx->module_dgst_size = module_dgst_size;
|
||||
module_ctx->module_dictstat_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_esalt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_buffer_size = MODULE_DEFAULT;
|
||||
module_ctx->module_extra_tmp_size = MODULE_DEFAULT;
|
||||
module_ctx->module_forced_outfile_format = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_count = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_parse = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_binary_save = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_outfile = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode_zero_hash = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_decode = module_hash_decode;
|
||||
module_ctx->module_hash_encode_status = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_encode = module_hash_encode;
|
||||
module_ctx->module_hash_init_selftest = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_mode = MODULE_DEFAULT;
|
||||
module_ctx->module_hash_category = module_hash_category;
|
||||
module_ctx->module_hash_name = module_hash_name;
|
||||
module_ctx->module_hlfmt_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_hook12 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook23 = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_salt_size = MODULE_DEFAULT;
|
||||
module_ctx->module_hook_size = MODULE_DEFAULT;
|
||||
module_ctx->module_jit_build_options = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_accel_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_loops_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_max = MODULE_DEFAULT;
|
||||
module_ctx->module_kernel_threads_min = MODULE_DEFAULT;
|
||||
module_ctx->module_kern_type = module_kern_type;
|
||||
module_ctx->module_kern_type_dynamic = MODULE_DEFAULT;
|
||||
module_ctx->module_opti_type = module_opti_type;
|
||||
module_ctx->module_opts_type = module_opts_type;
|
||||
module_ctx->module_outfile_check_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_outfile_check_nocomp = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_disable = MODULE_DEFAULT;
|
||||
module_ctx->module_potfile_keep_all_hashes = MODULE_DEFAULT;
|
||||
module_ctx->module_pwdump_column = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_max = MODULE_DEFAULT;
|
||||
module_ctx->module_pw_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_max = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_min = MODULE_DEFAULT;
|
||||
module_ctx->module_salt_type = module_salt_type;
|
||||
module_ctx->module_separator = MODULE_DEFAULT;
|
||||
module_ctx->module_st_hash = module_st_hash;
|
||||
module_ctx->module_st_pass = module_st_pass;
|
||||
module_ctx->module_tmp_size = module_tmp_size;
|
||||
module_ctx->module_unstable_warning = MODULE_DEFAULT;
|
||||
module_ctx->module_warmup_disable = MODULE_DEFAULT;
|
||||
}
|
||||
@ -0,0 +1,65 @@
|
||||
#!/usr/bin/env perl
|
||||
|
||||
##
|
||||
## Author......: See docs/credits.txt
|
||||
## License.....: MIT
|
||||
##
|
||||
|
||||
use strict;
|
||||
use warnings;
|
||||
|
||||
use Digest::MD5 qw (md5_hex);
|
||||
|
||||
sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
|
||||
|
||||
sub module_generate_hash
|
||||
{
|
||||
my $word = shift;
|
||||
my $salt = shift;
|
||||
my $iter = shift // 1000;
|
||||
|
||||
my $data;
|
||||
|
||||
$data .= $salt;
|
||||
$data .= $word x $iter;
|
||||
$data .= $word;
|
||||
|
||||
my $digest = md5_hex ($data);
|
||||
|
||||
my $hash;
|
||||
|
||||
if ($iter == 1000)
|
||||
{
|
||||
$hash = sprintf ("\@m\@%s\@%s", $digest, $salt);
|
||||
}
|
||||
else
|
||||
{
|
||||
$hash = sprintf ("\@m,%u\@%s\@%s", $iter, $digest, $salt);
|
||||
}
|
||||
|
||||
return $hash;
|
||||
}
|
||||
|
||||
sub module_verify_hash
|
||||
{
|
||||
my $line = shift;
|
||||
|
||||
my ($hash, $word) = split (':', $line);
|
||||
|
||||
return unless defined $hash;
|
||||
return unless defined $word;
|
||||
|
||||
my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
|
||||
|
||||
my ($type, $iter) = split (/\,/, $tag);
|
||||
|
||||
return unless ($type eq "m");
|
||||
|
||||
my $word_packed = pack_if_HEX_notation ($word);
|
||||
|
||||
my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
|
||||
|
||||
return ($new_hash, $word);
|
||||
}
|
||||
|
||||
1;
|
||||
@ -0,0 +1,65 @@
|
||||
#!/usr/bin/env perl
|
||||
|
||||
##
|
||||
## Author......: See docs/credits.txt
|
||||
## License.....: MIT
|
||||
##
|
||||
|
||||
use strict;
|
||||
use warnings;
|
||||
|
||||
use Digest::SHA qw (sha256_hex);
|
||||
|
||||
sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
|
||||
|
||||
sub module_generate_hash
|
||||
{
|
||||
my $word = shift;
|
||||
my $salt = shift;
|
||||
my $iter = shift // 1000;
|
||||
|
||||
my $data;
|
||||
|
||||
$data .= $salt;
|
||||
$data .= $word x $iter;
|
||||
$data .= $word;
|
||||
|
||||
my $digest = sha256_hex ($data);
|
||||
|
||||
my $hash;
|
||||
|
||||
if ($iter == 1000)
|
||||
{
|
||||
$hash = sprintf ("\@s\@%s\@%s", $digest, $salt);
|
||||
}
|
||||
else
|
||||
{
|
||||
$hash = sprintf ("\@s,%u\@%s\@%s", $iter, $digest, $salt);
|
||||
}
|
||||
|
||||
return $hash;
|
||||
}
|
||||
|
||||
sub module_verify_hash
|
||||
{
|
||||
my $line = shift;
|
||||
|
||||
my ($hash, $word) = split (':', $line);
|
||||
|
||||
return unless defined $hash;
|
||||
return unless defined $word;
|
||||
|
||||
my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
|
||||
|
||||
my ($type, $iter) = split (/\,/, $tag);
|
||||
|
||||
return unless ($type eq "s");
|
||||
|
||||
my $word_packed = pack_if_HEX_notation ($word);
|
||||
|
||||
my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
|
||||
|
||||
return ($new_hash, $word);
|
||||
}
|
||||
|
||||
1;
|
||||
@ -0,0 +1,65 @@
|
||||
#!/usr/bin/env perl
|
||||
|
||||
##
|
||||
## Author......: See docs/credits.txt
|
||||
## License.....: MIT
|
||||
##
|
||||
|
||||
use strict;
|
||||
use warnings;
|
||||
|
||||
use Digest::SHA qw (sha512_hex);
|
||||
|
||||
sub module_constraints { [[0, 256], [16, 16], [-1, -1], [-1, -1], [-1, -1]] }
|
||||
|
||||
sub module_generate_hash
|
||||
{
|
||||
my $word = shift;
|
||||
my $salt = shift;
|
||||
my $iter = shift // 1000;
|
||||
|
||||
my $data;
|
||||
|
||||
$data .= $salt;
|
||||
$data .= $word x $iter;
|
||||
$data .= $word;
|
||||
|
||||
my $digest = sha512_hex ($data);
|
||||
|
||||
my $hash;
|
||||
|
||||
if ($iter == 1000)
|
||||
{
|
||||
$hash = sprintf ("\@S\@%s\@%s", $digest, $salt);
|
||||
}
|
||||
else
|
||||
{
|
||||
$hash = sprintf ("\@S,%u\@%s\@%s", $iter, $digest, $salt);
|
||||
}
|
||||
|
||||
return $hash;
|
||||
}
|
||||
|
||||
sub module_verify_hash
|
||||
{
|
||||
my $line = shift;
|
||||
|
||||
my ($hash, $word) = split (':', $line);
|
||||
|
||||
return unless defined $hash;
|
||||
return unless defined $word;
|
||||
|
||||
my (undef, $tag, $digest, $salt) = split (/\@/, $hash);
|
||||
|
||||
my ($type, $iter) = split (/\,/, $tag);
|
||||
|
||||
return unless ($type eq "S");
|
||||
|
||||
my $word_packed = pack_if_HEX_notation ($word);
|
||||
|
||||
my $new_hash = module_generate_hash ($word_packed, $salt, $iter);
|
||||
|
||||
return ($new_hash, $word);
|
||||
}
|
||||
|
||||
1;
|
||||
Loading…
Reference in new issue