mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-22 14:48:12 +00:00
Added new hash mode -m 13000 = RAR5
This commit is contained in:
jsteube
parent
ea9ca63083
commit
f0a84a2410
528
OpenCL/m13000.cl
Normal file
528
OpenCL/m13000.cl
Normal file
@ -0,0 +1,528 @@
|
||||
/**
|
||||
* Author......: Jens Steube <jens.steube@gmail.com>
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _RAR5_
|
||||
|
||||
#include "include/constants.h"
|
||||
#include "include/kernel_vendor.h"
|
||||
|
||||
#define DGST_R0 0
|
||||
#define DGST_R1 1
|
||||
#define DGST_R2 2
|
||||
#define DGST_R3 3
|
||||
|
||||
#include "include/kernel_functions.c"
|
||||
|
||||
#include "OpenCL/types_ocl.c"
|
||||
#include "OpenCL/common.c"
|
||||
|
||||
#define COMPARE_S "OpenCL/check_single_comp4.c"
|
||||
#define COMPARE_M "OpenCL/check_multi_comp4.c"
|
||||
|
||||
__constant u32 k_sha256[64] =
|
||||
{
|
||||
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
|
||||
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
|
||||
SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
|
||||
SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
|
||||
SHA256C10, SHA256C11, SHA256C12, SHA256C13,
|
||||
SHA256C14, SHA256C15, SHA256C16, SHA256C17,
|
||||
SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
|
||||
SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
|
||||
SHA256C20, SHA256C21, SHA256C22, SHA256C23,
|
||||
SHA256C24, SHA256C25, SHA256C26, SHA256C27,
|
||||
SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
|
||||
SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
|
||||
SHA256C30, SHA256C31, SHA256C32, SHA256C33,
|
||||
SHA256C34, SHA256C35, SHA256C36, SHA256C37,
|
||||
SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
|
||||
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
|
||||
};
|
||||
|
||||
static void sha256_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[8])
|
||||
{
|
||||
u32 a = digest[0];
|
||||
u32 b = digest[1];
|
||||
u32 c = digest[2];
|
||||
u32 d = digest[3];
|
||||
u32 e = digest[4];
|
||||
u32 f = digest[5];
|
||||
u32 g = digest[6];
|
||||
u32 h = digest[7];
|
||||
|
||||
u32 w0_t = w0[0];
|
||||
u32 w1_t = w0[1];
|
||||
u32 w2_t = w0[2];
|
||||
u32 w3_t = w0[3];
|
||||
u32 w4_t = w1[0];
|
||||
u32 w5_t = w1[1];
|
||||
u32 w6_t = w1[2];
|
||||
u32 w7_t = w1[3];
|
||||
u32 w8_t = w2[0];
|
||||
u32 w9_t = w2[1];
|
||||
u32 wa_t = w2[2];
|
||||
u32 wb_t = w2[3];
|
||||
u32 wc_t = w3[0];
|
||||
u32 wd_t = w3[1];
|
||||
u32 we_t = w3[2];
|
||||
u32 wf_t = w3[3];
|
||||
|
||||
#define ROUND_EXPAND() \
|
||||
{ \
|
||||
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
|
||||
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
|
||||
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
|
||||
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
|
||||
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
|
||||
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
|
||||
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
|
||||
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
|
||||
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
|
||||
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
|
||||
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
|
||||
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
|
||||
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
|
||||
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
|
||||
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
|
||||
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
|
||||
}
|
||||
|
||||
#define ROUND_STEP(i) \
|
||||
{ \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
|
||||
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
|
||||
}
|
||||
|
||||
ROUND_STEP (0);
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 16; i < 64; i += 16)
|
||||
{
|
||||
ROUND_EXPAND (); ROUND_STEP (i);
|
||||
}
|
||||
|
||||
digest[0] += a;
|
||||
digest[1] += b;
|
||||
digest[2] += c;
|
||||
digest[3] += d;
|
||||
digest[4] += e;
|
||||
digest[5] += f;
|
||||
digest[6] += g;
|
||||
digest[7] += h;
|
||||
}
|
||||
|
||||
static void hmac_sha256_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u32 opad[8])
|
||||
{
|
||||
w0[0] = w0[0] ^ 0x36363636;
|
||||
w0[1] = w0[1] ^ 0x36363636;
|
||||
w0[2] = w0[2] ^ 0x36363636;
|
||||
w0[3] = w0[3] ^ 0x36363636;
|
||||
w1[0] = w1[0] ^ 0x36363636;
|
||||
w1[1] = w1[1] ^ 0x36363636;
|
||||
w1[2] = w1[2] ^ 0x36363636;
|
||||
w1[3] = w1[3] ^ 0x36363636;
|
||||
w2[0] = w2[0] ^ 0x36363636;
|
||||
w2[1] = w2[1] ^ 0x36363636;
|
||||
w2[2] = w2[2] ^ 0x36363636;
|
||||
w2[3] = w2[3] ^ 0x36363636;
|
||||
w3[0] = w3[0] ^ 0x36363636;
|
||||
w3[1] = w3[1] ^ 0x36363636;
|
||||
w3[2] = w3[2] ^ 0x36363636;
|
||||
w3[3] = w3[3] ^ 0x36363636;
|
||||
|
||||
ipad[0] = SHA256M_A;
|
||||
ipad[1] = SHA256M_B;
|
||||
ipad[2] = SHA256M_C;
|
||||
ipad[3] = SHA256M_D;
|
||||
ipad[4] = SHA256M_E;
|
||||
ipad[5] = SHA256M_F;
|
||||
ipad[6] = SHA256M_G;
|
||||
ipad[7] = SHA256M_H;
|
||||
|
||||
sha256_transform (w0, w1, w2, w3, ipad);
|
||||
|
||||
w0[0] = w0[0] ^ 0x6a6a6a6a;
|
||||
w0[1] = w0[1] ^ 0x6a6a6a6a;
|
||||
w0[2] = w0[2] ^ 0x6a6a6a6a;
|
||||
w0[3] = w0[3] ^ 0x6a6a6a6a;
|
||||
w1[0] = w1[0] ^ 0x6a6a6a6a;
|
||||
w1[1] = w1[1] ^ 0x6a6a6a6a;
|
||||
w1[2] = w1[2] ^ 0x6a6a6a6a;
|
||||
w1[3] = w1[3] ^ 0x6a6a6a6a;
|
||||
w2[0] = w2[0] ^ 0x6a6a6a6a;
|
||||
w2[1] = w2[1] ^ 0x6a6a6a6a;
|
||||
w2[2] = w2[2] ^ 0x6a6a6a6a;
|
||||
w2[3] = w2[3] ^ 0x6a6a6a6a;
|
||||
w3[0] = w3[0] ^ 0x6a6a6a6a;
|
||||
w3[1] = w3[1] ^ 0x6a6a6a6a;
|
||||
w3[2] = w3[2] ^ 0x6a6a6a6a;
|
||||
w3[3] = w3[3] ^ 0x6a6a6a6a;
|
||||
|
||||
opad[0] = SHA256M_A;
|
||||
opad[1] = SHA256M_B;
|
||||
opad[2] = SHA256M_C;
|
||||
opad[3] = SHA256M_D;
|
||||
opad[4] = SHA256M_E;
|
||||
opad[5] = SHA256M_F;
|
||||
opad[6] = SHA256M_G;
|
||||
opad[7] = SHA256M_H;
|
||||
|
||||
sha256_transform (w0, w1, w2, w3, opad);
|
||||
}
|
||||
|
||||
static void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u32 opad[8], u32 digest[8])
|
||||
{
|
||||
digest[0] = ipad[0];
|
||||
digest[1] = ipad[1];
|
||||
digest[2] = ipad[2];
|
||||
digest[3] = ipad[3];
|
||||
digest[4] = ipad[4];
|
||||
digest[5] = ipad[5];
|
||||
digest[6] = ipad[6];
|
||||
digest[7] = ipad[7];
|
||||
|
||||
sha256_transform (w0, w1, w2, w3, digest);
|
||||
|
||||
w0[0] = digest[0];
|
||||
w0[1] = digest[1];
|
||||
w0[2] = digest[2];
|
||||
w0[3] = digest[3];
|
||||
w1[0] = digest[4];
|
||||
w1[1] = digest[5];
|
||||
w1[2] = digest[6];
|
||||
w1[3] = digest[7];
|
||||
w2[0] = 0x80000000;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = (64 + 32) * 8;
|
||||
|
||||
digest[0] = opad[0];
|
||||
digest[1] = opad[1];
|
||||
digest[2] = opad[2];
|
||||
digest[3] = opad[3];
|
||||
digest[4] = opad[4];
|
||||
digest[5] = opad[5];
|
||||
digest[6] = opad[6];
|
||||
digest[7] = opad[7];
|
||||
|
||||
sha256_transform (w0, w1, w2, w3, digest);
|
||||
}
|
||||
|
||||
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
u32 w0[4];
|
||||
|
||||
w0[0] = pws[gid].i[ 0];
|
||||
w0[1] = pws[gid].i[ 1];
|
||||
w0[2] = pws[gid].i[ 2];
|
||||
w0[3] = pws[gid].i[ 3];
|
||||
|
||||
u32 w1[4];
|
||||
|
||||
w1[0] = pws[gid].i[ 4];
|
||||
w1[1] = pws[gid].i[ 5];
|
||||
w1[2] = pws[gid].i[ 6];
|
||||
w1[3] = pws[gid].i[ 7];
|
||||
|
||||
u32 w2[4];
|
||||
|
||||
w2[0] = pws[gid].i[ 8];
|
||||
w2[1] = pws[gid].i[ 9];
|
||||
w2[2] = pws[gid].i[10];
|
||||
w2[3] = pws[gid].i[11];
|
||||
|
||||
u32 w3[4];
|
||||
|
||||
w3[0] = pws[gid].i[12];
|
||||
w3[1] = pws[gid].i[13];
|
||||
w3[2] = pws[gid].i[14];
|
||||
w3[3] = pws[gid].i[15];
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
w0[0] = swap32 (w0[0]);
|
||||
w0[1] = swap32 (w0[1]);
|
||||
w0[2] = swap32 (w0[2]);
|
||||
w0[3] = swap32 (w0[3]);
|
||||
w1[0] = swap32 (w1[0]);
|
||||
w1[1] = swap32 (w1[1]);
|
||||
w1[2] = swap32 (w1[2]);
|
||||
w1[3] = swap32 (w1[3]);
|
||||
w2[0] = swap32 (w2[0]);
|
||||
w2[1] = swap32 (w2[1]);
|
||||
w2[2] = swap32 (w2[2]);
|
||||
w2[3] = swap32 (w2[3]);
|
||||
w3[0] = swap32 (w3[0]);
|
||||
w3[1] = swap32 (w3[1]);
|
||||
w3[2] = swap32 (w3[2]);
|
||||
w3[3] = swap32 (w3[3]);
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
u32 salt_len = 16;
|
||||
|
||||
u32 salt_buf0[4];
|
||||
|
||||
salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
|
||||
salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
|
||||
salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
|
||||
salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
|
||||
|
||||
u32 salt_buf1[4];
|
||||
|
||||
salt_buf1[0] = 0x00000001;
|
||||
salt_buf1[1] = 0x80000000;
|
||||
salt_buf1[2] = 0;
|
||||
salt_buf1[3] = 0;
|
||||
|
||||
u32 salt_buf2[4];
|
||||
|
||||
salt_buf2[0] = 0;
|
||||
salt_buf2[1] = 0;
|
||||
salt_buf2[2] = 0;
|
||||
salt_buf2[3] = 0;
|
||||
|
||||
u32 salt_buf3[4];
|
||||
|
||||
salt_buf3[0] = 0;
|
||||
salt_buf3[1] = 0;
|
||||
salt_buf3[2] = 0;
|
||||
salt_buf3[3] = (64 + salt_len + 4) * 8;
|
||||
|
||||
u32 ipad[8];
|
||||
u32 opad[8];
|
||||
|
||||
hmac_sha256_pad (w0, w1, w2, w3, ipad, opad);
|
||||
|
||||
tmps[gid].ipad[0] = ipad[0];
|
||||
tmps[gid].ipad[1] = ipad[1];
|
||||
tmps[gid].ipad[2] = ipad[2];
|
||||
tmps[gid].ipad[3] = ipad[3];
|
||||
tmps[gid].ipad[4] = ipad[4];
|
||||
tmps[gid].ipad[5] = ipad[5];
|
||||
tmps[gid].ipad[6] = ipad[6];
|
||||
tmps[gid].ipad[7] = ipad[7];
|
||||
|
||||
tmps[gid].opad[0] = opad[0];
|
||||
tmps[gid].opad[1] = opad[1];
|
||||
tmps[gid].opad[2] = opad[2];
|
||||
tmps[gid].opad[3] = opad[3];
|
||||
tmps[gid].opad[4] = opad[4];
|
||||
tmps[gid].opad[5] = opad[5];
|
||||
tmps[gid].opad[6] = opad[6];
|
||||
tmps[gid].opad[7] = opad[7];
|
||||
|
||||
for (u32 i = 0, j = 1; i < 8; i += 8, j += 1)
|
||||
{
|
||||
u32 dgst[8];
|
||||
|
||||
hmac_sha256_run (salt_buf0, salt_buf1, salt_buf2, salt_buf3, ipad, opad, dgst);
|
||||
|
||||
tmps[gid].dgst[i + 0] = dgst[0];
|
||||
tmps[gid].dgst[i + 1] = dgst[1];
|
||||
tmps[gid].dgst[i + 2] = dgst[2];
|
||||
tmps[gid].dgst[i + 3] = dgst[3];
|
||||
tmps[gid].dgst[i + 4] = dgst[4];
|
||||
tmps[gid].dgst[i + 5] = dgst[5];
|
||||
tmps[gid].dgst[i + 6] = dgst[6];
|
||||
tmps[gid].dgst[i + 7] = dgst[7];
|
||||
|
||||
tmps[gid].out[i + 0] = dgst[0];
|
||||
tmps[gid].out[i + 1] = dgst[1];
|
||||
tmps[gid].out[i + 2] = dgst[2];
|
||||
tmps[gid].out[i + 3] = dgst[3];
|
||||
tmps[gid].out[i + 4] = dgst[4];
|
||||
tmps[gid].out[i + 5] = dgst[5];
|
||||
tmps[gid].out[i + 6] = dgst[6];
|
||||
tmps[gid].out[i + 7] = dgst[7];
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
const u32 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 ipad[8];
|
||||
|
||||
ipad[0] = tmps[gid].ipad[0];
|
||||
ipad[1] = tmps[gid].ipad[1];
|
||||
ipad[2] = tmps[gid].ipad[2];
|
||||
ipad[3] = tmps[gid].ipad[3];
|
||||
ipad[4] = tmps[gid].ipad[4];
|
||||
ipad[5] = tmps[gid].ipad[5];
|
||||
ipad[6] = tmps[gid].ipad[6];
|
||||
ipad[7] = tmps[gid].ipad[7];
|
||||
|
||||
u32 opad[8];
|
||||
|
||||
opad[0] = tmps[gid].opad[0];
|
||||
opad[1] = tmps[gid].opad[1];
|
||||
opad[2] = tmps[gid].opad[2];
|
||||
opad[3] = tmps[gid].opad[3];
|
||||
opad[4] = tmps[gid].opad[4];
|
||||
opad[5] = tmps[gid].opad[5];
|
||||
opad[6] = tmps[gid].opad[6];
|
||||
opad[7] = tmps[gid].opad[7];
|
||||
|
||||
for (u32 i = 0; i < 8; i += 8)
|
||||
{
|
||||
u32 dgst[8];
|
||||
|
||||
dgst[0] = tmps[gid].dgst[i + 0];
|
||||
dgst[1] = tmps[gid].dgst[i + 1];
|
||||
dgst[2] = tmps[gid].dgst[i + 2];
|
||||
dgst[3] = tmps[gid].dgst[i + 3];
|
||||
dgst[4] = tmps[gid].dgst[i + 4];
|
||||
dgst[5] = tmps[gid].dgst[i + 5];
|
||||
dgst[6] = tmps[gid].dgst[i + 6];
|
||||
dgst[7] = tmps[gid].dgst[i + 7];
|
||||
|
||||
u32 out[8];
|
||||
|
||||
out[0] = tmps[gid].out[i + 0];
|
||||
out[1] = tmps[gid].out[i + 1];
|
||||
out[2] = tmps[gid].out[i + 2];
|
||||
out[3] = tmps[gid].out[i + 3];
|
||||
out[4] = tmps[gid].out[i + 4];
|
||||
out[5] = tmps[gid].out[i + 5];
|
||||
out[6] = tmps[gid].out[i + 6];
|
||||
out[7] = tmps[gid].out[i + 7];
|
||||
|
||||
for (u32 j = 0; j < loop_cnt; j++)
|
||||
{
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
w0[0] = dgst[0];
|
||||
w0[1] = dgst[1];
|
||||
w0[2] = dgst[2];
|
||||
w0[3] = dgst[3];
|
||||
w1[0] = dgst[4];
|
||||
w1[1] = dgst[5];
|
||||
w1[2] = dgst[6];
|
||||
w1[3] = dgst[7];
|
||||
w2[0] = 0x80000000;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = (64 + 32) * 8;
|
||||
|
||||
hmac_sha256_run (w0, w1, w2, w3, ipad, opad, dgst);
|
||||
|
||||
out[0] ^= dgst[0];
|
||||
out[1] ^= dgst[1];
|
||||
out[2] ^= dgst[2];
|
||||
out[3] ^= dgst[3];
|
||||
out[4] ^= dgst[4];
|
||||
out[5] ^= dgst[5];
|
||||
out[6] ^= dgst[6];
|
||||
out[7] ^= dgst[7];
|
||||
}
|
||||
|
||||
tmps[gid].dgst[i + 0] = dgst[0];
|
||||
tmps[gid].dgst[i + 1] = dgst[1];
|
||||
tmps[gid].dgst[i + 2] = dgst[2];
|
||||
tmps[gid].dgst[i + 3] = dgst[3];
|
||||
tmps[gid].dgst[i + 4] = dgst[4];
|
||||
tmps[gid].dgst[i + 5] = dgst[5];
|
||||
tmps[gid].dgst[i + 6] = dgst[6];
|
||||
tmps[gid].dgst[i + 7] = dgst[7];
|
||||
|
||||
tmps[gid].out[i + 0] = out[0];
|
||||
tmps[gid].out[i + 1] = out[1];
|
||||
tmps[gid].out[i + 2] = out[2];
|
||||
tmps[gid].out[i + 3] = out[3];
|
||||
tmps[gid].out[i + 4] = out[4];
|
||||
tmps[gid].out[i + 5] = out[5];
|
||||
tmps[gid].out[i + 6] = out[6];
|
||||
tmps[gid].out[i + 7] = out[7];
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
|
||||
{
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
u32 outl4[4];
|
||||
|
||||
outl4[0] = tmps[gid].out[0];
|
||||
outl4[1] = tmps[gid].out[2];
|
||||
outl4[2] = tmps[gid].out[4];
|
||||
outl4[3] = tmps[gid].out[6];
|
||||
|
||||
u32 outr4[4];
|
||||
|
||||
outr4[0] = tmps[gid].out[1];
|
||||
outr4[1] = tmps[gid].out[3];
|
||||
outr4[2] = tmps[gid].out[5];
|
||||
outr4[3] = tmps[gid].out[7];
|
||||
|
||||
u32 outl = outl4[0]
|
||||
^ outl4[1]
|
||||
^ outl4[2]
|
||||
^ outl4[3];
|
||||
|
||||
u32 outr = outr4[0]
|
||||
^ outr4[1]
|
||||
^ outr4[2]
|
||||
^ outr4[3];
|
||||
|
||||
const u32 r0 = outl;
|
||||
const u32 r1 = outr;
|
||||
const u32 r2 = 0;
|
||||
const u32 r3 = 0;
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#include COMPARE_M
|
||||
}
|
||||
@ -490,6 +490,8 @@ typedef struct
|
||||
u32 digest_buf[8];
|
||||
#elif defined _ANDROIDFDE_SAMSUNG_
|
||||
u32 digest_buf[8];
|
||||
#elif defined _RAR5_
|
||||
u32 digest_buf[4];
|
||||
#endif
|
||||
|
||||
} digest_t;
|
||||
|
||||
@ -13,6 +13,11 @@ File.: Kernel
|
||||
Desc.: Added new hash mode -m 12900 = Android FDE (Samsung DEK)
|
||||
Issue: 86
|
||||
|
||||
Type.: Feature
|
||||
File.: Kernel
|
||||
Desc.: Added new hash mode -m 13000 = RAR5
|
||||
Issue: 68
|
||||
|
||||
Type.: Feature
|
||||
File.: Host
|
||||
Desc.: Added a Makefile target for native compilation
|
||||
|
||||
@ -183,6 +183,7 @@ AMD users require Catalyst 14.9 or later (recommended 15.9 or later)
|
||||
- Skype
|
||||
- 7-Zip
|
||||
- RAR3-hp
|
||||
- RAR5
|
||||
- PDF 1.1 - 1.3 (Acrobat 2 - 4)
|
||||
- PDF 1.4 - 1.6 (Acrobat 5 - 8)
|
||||
- PDF 1.7 Level 3 (Acrobat 9)
|
||||
|
||||
@ -171,7 +171,7 @@ _oclHashcat ()
|
||||
{
|
||||
local VERSION=2.10
|
||||
|
||||
local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 190 200 300 400 500 501 900 1000 1100 1400 1410 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12100 12200 12300 12400 12500 12600 12700 12800 12900"
|
||||
local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 190 200 300 400 500 501 900 1000 1100 1400 1410 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000"
|
||||
local ATTACK_MODES="0 1 3 6 7"
|
||||
local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15"
|
||||
local BENCHMARK_MODE="0 1"
|
||||
|
||||
@ -71,7 +71,7 @@
|
||||
#define SHA1C03 0xca62c1d6u
|
||||
#endif
|
||||
|
||||
#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_
|
||||
#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_ || defined _RAR5_
|
||||
/**
|
||||
* SHA256 Constants
|
||||
*/
|
||||
|
||||
@ -193,7 +193,7 @@
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_
|
||||
#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_ || defined _RAR5_
|
||||
|
||||
#define SHIFT_RIGHT_32(x,n) ((x) >> (n))
|
||||
|
||||
|
||||
@ -352,6 +352,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERNEL_ACCEL_12700 64
|
||||
#define KERNEL_ACCEL_12800 64
|
||||
#define KERNEL_ACCEL_12900 8
|
||||
#define KERNEL_ACCEL_13000 8
|
||||
|
||||
#define KERNEL_LOOPS_0 256
|
||||
#define KERNEL_LOOPS_10 256
|
||||
@ -530,6 +531,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERNEL_LOOPS_12700 10
|
||||
#define KERNEL_LOOPS_12800 100
|
||||
#define KERNEL_LOOPS_12900 64
|
||||
#define KERNEL_LOOPS_13000 64
|
||||
|
||||
/**
|
||||
* Strings
|
||||
@ -676,6 +678,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define HT_12700 "Blockchain, My Wallet"
|
||||
#define HT_12800 "MS-AzureSync PBKDF2-HMAC-SHA256"
|
||||
#define HT_12900 "Android FDE (Samsung DEK)"
|
||||
#define HT_13000 "RAR5"
|
||||
|
||||
#define HT_00011 "Joomla < 2.5.18"
|
||||
#define HT_00012 "PostgreSQL"
|
||||
@ -1007,6 +1010,8 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define DISPLAY_LEN_MAX_12800 11 + 1 + 20 + 1 + 5 + 1 + 64
|
||||
#define DISPLAY_LEN_MIN_12900 64 + 64 + 32
|
||||
#define DISPLAY_LEN_MAX_12900 64 + 64 + 32
|
||||
#define DISPLAY_LEN_MIN_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
|
||||
|
||||
#define DISPLAY_LEN_MIN_11 32 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_11 32 + 1 + 32
|
||||
@ -1264,6 +1269,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERN_TYPE_MYWALLET 12700
|
||||
#define KERN_TYPE_MS_DRSR 12800
|
||||
#define KERN_TYPE_ANDROIDFDE_SAMSUNG 12900
|
||||
#define KERN_TYPE_RAR5 13000
|
||||
|
||||
/**
|
||||
* signatures
|
||||
@ -1332,6 +1338,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define SIGNATURE_RAR3 "$RAR3$"
|
||||
#define SIGNATURE_MYWALLET "$blockchain$"
|
||||
#define SIGNATURE_MS_DRSR "v1;PPH1_MD4"
|
||||
#define SIGNATURE_RAR5 "$rar5$"
|
||||
|
||||
/**
|
||||
* Default iteration numbers
|
||||
@ -1382,6 +1389,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define ROUNDS_MYWALLET 10
|
||||
#define ROUNDS_MS_DRSR 100
|
||||
#define ROUNDS_ANDROIDFDE_SAMSUNG 4096
|
||||
#define ROUNDS_RAR5 (1 << 15)
|
||||
|
||||
/**
|
||||
* salt types
|
||||
@ -1861,6 +1869,7 @@ int pbkdf2_sha512_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
int ecryptfs_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int bsdicrypt_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int rar3hp_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int rar5_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int cf10_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int mywallet_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
int ms_drsr_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
|
||||
@ -128,7 +128,7 @@ const uint RESTORE_MIN = 210;
|
||||
|
||||
#define MAX_DICTSTAT 10000
|
||||
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 129
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 130
|
||||
|
||||
#define global_free(attr) \
|
||||
{ \
|
||||
@ -249,6 +249,7 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
|
||||
133,
|
||||
11600,
|
||||
12500,
|
||||
13000,
|
||||
6211,
|
||||
6221,
|
||||
6231,
|
||||
@ -668,6 +669,7 @@ const char *USAGE_BIG[] =
|
||||
"",
|
||||
" 11600 = 7-Zip",
|
||||
" 12500 = RAR3-hp",
|
||||
" 13000 = RAR5",
|
||||
"",
|
||||
"[[ Full-Disk encryptions (FDE) ]]",
|
||||
"",
|
||||
@ -5610,7 +5612,7 @@ int main (int argc, char **argv)
|
||||
return (-1);
|
||||
}
|
||||
|
||||
if (hash_mode_chgd && hash_mode > 12900) // just added to remove compiler warnings for hash_mode_chgd
|
||||
if (hash_mode_chgd && hash_mode > 13000) // just added to remove compiler warnings for hash_mode_chgd
|
||||
{
|
||||
log_error ("ERROR: Invalid hash-type specified");
|
||||
|
||||
@ -9792,6 +9794,21 @@ int main (int argc, char **argv)
|
||||
dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
case 13000: hash_type = HASH_TYPE_PBKDF2_SHA256;
|
||||
salt_type = SALT_TYPE_EMBEDDED;
|
||||
attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
opts_type = OPTS_TYPE_PT_GENERATE_LE;
|
||||
kern_type = KERN_TYPE_RAR5;
|
||||
dgst_size = DGST_SIZE_4_4;
|
||||
parse_func = rar5_parse_hash;
|
||||
sort_by_digest = sort_by_digest_4_4;
|
||||
opti_type = OPTI_TYPE_ZERO_BYTE;
|
||||
dgst_pos0 = 0;
|
||||
dgst_pos1 = 1;
|
||||
dgst_pos2 = 2;
|
||||
dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
default: usage_mini_print (PROGNAME); return (-1);
|
||||
}
|
||||
|
||||
@ -11148,6 +11165,8 @@ int main (int argc, char **argv)
|
||||
break;
|
||||
case 12900: hashes_buf[0].salt->salt_iter = ROUNDS_ANDROIDFDE_SAMSUNG - 1;
|
||||
break;
|
||||
case 13000: hashes_buf[0].salt->salt_iter = ROUNDS_RAR5 - 1;
|
||||
break;
|
||||
}
|
||||
|
||||
// set special tuning for benchmark-mode 1
|
||||
@ -11333,6 +11352,9 @@ int main (int argc, char **argv)
|
||||
case 12900: kernel_loops = ROUNDS_ANDROIDFDE_SAMSUNG;
|
||||
kernel_accel = 8;
|
||||
break;
|
||||
case 13000: kernel_loops = ROUNDS_RAR5;
|
||||
kernel_accel = 8;
|
||||
break;
|
||||
}
|
||||
|
||||
// some algorithm collide too fast, make that impossible
|
||||
@ -13112,6 +13134,7 @@ int main (int argc, char **argv)
|
||||
case 12700: size_tmps = kernel_blocks * sizeof (mywallet_tmp_t); break;
|
||||
case 12800: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
|
||||
case 12900: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
|
||||
case 13000: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
|
||||
};
|
||||
|
||||
uint size_hooks = 4;
|
||||
|
||||
129
src/shared.c
129
src/shared.c
@ -5611,6 +5611,7 @@ char *strhashtype (const uint hash_mode)
|
||||
case 12700: return ((char *) HT_12700); break;
|
||||
case 12800: return ((char *) HT_12800); break;
|
||||
case 12900: return ((char *) HT_12900); break;
|
||||
case 13000: return ((char *) HT_13000); break;
|
||||
}
|
||||
|
||||
return ((char *) "Unknown");
|
||||
@ -8068,6 +8069,22 @@ void ascii_digest (char out_buf[4096], uint salt_pos, uint digest_pos)
|
||||
salt.salt_buf[ 3]
|
||||
);
|
||||
}
|
||||
else if (hash_mode == 13000)
|
||||
{
|
||||
snprintf (out_buf, len-1, "$rar5$16$%08x%08x%08x%08x$%u$%08x%08x%08x%08x$8$%08x%08x",
|
||||
salt.salt_buf[0],
|
||||
salt.salt_buf[1],
|
||||
salt.salt_buf[2],
|
||||
salt.salt_buf[3],
|
||||
salt.salt_sign[0],
|
||||
salt.salt_buf[4],
|
||||
salt.salt_buf[5],
|
||||
salt.salt_buf[6],
|
||||
salt.salt_buf[7],
|
||||
byte_swap_32 (digest_buf[0]),
|
||||
byte_swap_32 (digest_buf[1])
|
||||
);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (hash_type == HASH_TYPE_MD4)
|
||||
@ -9023,6 +9040,7 @@ uint set_kernel_accel (uint hash_mode)
|
||||
case 12700: return GET_ACCEL (12700);
|
||||
case 12800: return GET_ACCEL (12800);
|
||||
case 12900: return GET_ACCEL (12900);
|
||||
case 13000: return GET_ACCEL (13000);
|
||||
}
|
||||
|
||||
return 0;
|
||||
@ -9209,6 +9227,7 @@ uint set_kernel_loops (uint hash_mode)
|
||||
case 12700: return GET_LOOPS (12700);
|
||||
case 12800: return GET_LOOPS (12800);
|
||||
case 12900: return GET_LOOPS (12900);
|
||||
case 13000: return GET_LOOPS (13000);
|
||||
}
|
||||
|
||||
return 0;
|
||||
@ -18365,6 +18384,116 @@ int rar3hp_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int rar5_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_13000) || (input_len > DISPLAY_LEN_MAX_13000)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
if (memcmp (SIGNATURE_RAR5, input_buf, 1 + 4 + 1)) return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
uint32_t *digest = (uint32_t *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
/**
|
||||
* parse line
|
||||
*/
|
||||
|
||||
char *param0_pos = input_buf + 1 + 4 + 1;
|
||||
|
||||
char *param1_pos = strchr (param0_pos, '$');
|
||||
|
||||
if (param1_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint32_t param0_len = param1_pos - param0_pos;
|
||||
|
||||
param1_pos++;
|
||||
|
||||
char *param2_pos = strchr (param1_pos, '$');
|
||||
|
||||
if (param2_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint32_t param1_len = param2_pos - param1_pos;
|
||||
|
||||
param2_pos++;
|
||||
|
||||
char *param3_pos = strchr (param2_pos, '$');
|
||||
|
||||
if (param3_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint32_t param2_len = param3_pos - param2_pos;
|
||||
|
||||
param3_pos++;
|
||||
|
||||
char *param4_pos = strchr (param3_pos, '$');
|
||||
|
||||
if (param4_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint32_t param3_len = param4_pos - param3_pos;
|
||||
|
||||
param4_pos++;
|
||||
|
||||
char *param5_pos = strchr (param4_pos, '$');
|
||||
|
||||
if (param5_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
uint32_t param4_len = param5_pos - param4_pos;
|
||||
|
||||
param5_pos++;
|
||||
|
||||
uint32_t param5_len = input_len - 1 - 4 - 1 - param0_len - 1 - param1_len - 1 - param2_len - 1 - param3_len - 1 - param4_len - 1;
|
||||
|
||||
char *salt_buf = param1_pos;
|
||||
char *iv = param3_pos;
|
||||
char *pswcheck = param5_pos;
|
||||
|
||||
const uint salt_len = atoi (param0_pos);
|
||||
const uint iterations = atoi (param2_pos);
|
||||
const uint pswcheck_len = atoi (param4_pos);
|
||||
|
||||
/**
|
||||
* verify some data
|
||||
*/
|
||||
|
||||
if (param1_len != 32) return (PARSER_SALT_VALUE);
|
||||
if (param3_len != 32) return (PARSER_SALT_VALUE);
|
||||
if (param5_len != 16) return (PARSER_SALT_VALUE);
|
||||
|
||||
if (salt_len != 16) return (PARSER_SALT_VALUE);
|
||||
if (iterations == 0) return (PARSER_SALT_VALUE);
|
||||
if (pswcheck_len != 8) return (PARSER_SALT_VALUE);
|
||||
|
||||
/**
|
||||
* store data
|
||||
*/
|
||||
|
||||
salt->salt_buf[0] = hex_to_uint (&salt_buf[ 0]);
|
||||
salt->salt_buf[1] = hex_to_uint (&salt_buf[ 8]);
|
||||
salt->salt_buf[2] = hex_to_uint (&salt_buf[16]);
|
||||
salt->salt_buf[3] = hex_to_uint (&salt_buf[24]);
|
||||
|
||||
salt->salt_buf[4] = hex_to_uint (&iv[ 0]);
|
||||
salt->salt_buf[5] = hex_to_uint (&iv[ 8]);
|
||||
salt->salt_buf[6] = hex_to_uint (&iv[16]);
|
||||
salt->salt_buf[7] = hex_to_uint (&iv[24]);
|
||||
|
||||
salt->salt_len = 16 + 16;
|
||||
|
||||
salt->salt_sign[0] = iterations;
|
||||
|
||||
salt->salt_iter = ((1 << iterations) + 32) - 1;
|
||||
|
||||
/**
|
||||
* digest buf
|
||||
*/
|
||||
|
||||
digest[0] = hex_to_uint (&pswcheck[ 0]);
|
||||
digest[1] = hex_to_uint (&pswcheck[ 8]);
|
||||
digest[2] = 0;
|
||||
digest[3] = 0;
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int cf10_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_12600) || (input_len > DISPLAY_LEN_MAX_12600)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
@ -44,7 +44,7 @@ my $hashcat = "./oclHashcat";
|
||||
|
||||
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, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900);
|
||||
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 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, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 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, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000);
|
||||
|
||||
my %is_unicode = map { $_ => 1 } qw(30 40 130 131 132 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800);
|
||||
my %less_fifteen = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
|
||||
@ -2175,6 +2175,40 @@ sub verify
|
||||
|
||||
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
|
||||
}
|
||||
elsif ($mode == 13000)
|
||||
{
|
||||
my $hash_line;
|
||||
|
||||
($hash_line, $word) = split ":", $line;
|
||||
|
||||
next unless defined $hash_line;
|
||||
next unless defined $word;
|
||||
|
||||
my @data = split ('\$', $hash_line);
|
||||
|
||||
next unless scalar @data == 8;
|
||||
|
||||
shift @data;
|
||||
|
||||
my $signature = shift @data;
|
||||
my $salt_len = shift @data;
|
||||
my $salt_buf = shift @data;
|
||||
my $iterations = shift @data;
|
||||
my $iv = shift @data;
|
||||
my $pswcheck_len = shift @data;
|
||||
my $pswcheck = shift @data;
|
||||
|
||||
next unless ($signature eq "rar5");
|
||||
next unless ($salt_len == 16);
|
||||
next unless ($pswcheck_len == 8);
|
||||
|
||||
$salt = $salt_buf;
|
||||
$iter = $iterations;
|
||||
$hash_in = $pswcheck;
|
||||
$param = $iv;
|
||||
|
||||
next unless (exists ($db->{$hash_line}) and (! defined ($db->{$hash_line})));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: hash mode is not supported\n";
|
||||
@ -2449,6 +2483,14 @@ sub verify
|
||||
|
||||
return unless (substr ($line, 0, $len) eq $hash_out);
|
||||
}
|
||||
elsif ($mode == 13000)
|
||||
{
|
||||
$hash_out = gen_hash ($mode, $word, $salt, $iter, $param);
|
||||
|
||||
$len = length $hash_out;
|
||||
|
||||
return unless (substr ($line, 0, $len) eq $hash_out);
|
||||
}
|
||||
else
|
||||
{
|
||||
$hash_out = gen_hash ($mode, $word, $salt, $iter);
|
||||
@ -2865,6 +2907,10 @@ sub passthrough
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
|
||||
}
|
||||
elsif ($mode == 13000)
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: Unsupported hash type\n";
|
||||
@ -3578,6 +3624,20 @@ sub single
|
||||
}
|
||||
}
|
||||
}
|
||||
elsif ($mode == 13000)
|
||||
{
|
||||
for (my $i = 1; $i < 32; $i++)
|
||||
{
|
||||
if ($len != 0)
|
||||
{
|
||||
rnd ($mode, $len, 32);
|
||||
}
|
||||
else
|
||||
{
|
||||
rnd ($mode, $i, 32);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -6518,6 +6578,42 @@ END_CODE
|
||||
|
||||
$tmp_hash = sprintf ("%s%s%s", $salt2, $hash_hmac, $salt_buf);
|
||||
}
|
||||
elsif ($mode == 13000)
|
||||
{
|
||||
my $iterations = 15;
|
||||
|
||||
if (length ($iter))
|
||||
{
|
||||
$iterations = int ($iter);
|
||||
}
|
||||
|
||||
my $iv = "0" x 32;
|
||||
|
||||
if (defined $additional_param)
|
||||
{
|
||||
$iv = $additional_param;
|
||||
}
|
||||
|
||||
my $pbkdf2 = Crypt::PBKDF2->new
|
||||
(
|
||||
hasher => Crypt::PBKDF2->hasher_from_algorithm ('HMACSHA2', 256),
|
||||
iterations => (1 << $iterations) + 32,
|
||||
output_len => 32
|
||||
);
|
||||
|
||||
my $salt_buf_bin = pack ("H*", $salt_buf);
|
||||
|
||||
my $hash = $pbkdf2->PBKDF2 ($salt_buf_bin, $word_buf);
|
||||
|
||||
printf ("%s\n", unpack ("H*", $hash));
|
||||
|
||||
my $hash_final = substr ($hash, 0, 8)
|
||||
^ substr ($hash, 8, 8)
|
||||
^ substr ($hash, 16, 8)
|
||||
^ substr ($hash, 24, 8);
|
||||
|
||||
$tmp_hash = sprintf ('$rar5$16$%s$%d$%s$8$%s', $salt_buf, $iterations, $iv, unpack ("H*", $hash_final));
|
||||
}
|
||||
|
||||
return ($tmp_hash);
|
||||
}
|
||||
|
||||
@ -7,7 +7,7 @@
|
||||
|
||||
# missing hash types: 5200,6211,6221,6231,6241,6251,6261,6271,6281
|
||||
|
||||
HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900"
|
||||
HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 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 12100 12200 12300 12400 12600 12800 12900 13000"
|
||||
|
||||
ATTACK_MODES="0 1 3 6 7"
|
||||
|
||||
@ -17,7 +17,7 @@ HASHFILE_ONLY="2500"
|
||||
|
||||
NEVER_CRACK="11600"
|
||||
|
||||
SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6221 6231 6241 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12100 12200 12300 12400 12500 12800 12900"
|
||||
SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6221 6231 6241 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12100 12200 12300 12400 12500 12800 12900 13000"
|
||||
|
||||
OPTS="--quiet --force --potfile-disable --runtime 200 --gpu-temp-disable -d 1 --weak-hash-threshold=0"
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user