mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-22 08:08:10 +00:00
Added new hash mode -m 12900 = Android FDE (Samsung DEK)
This commit is contained in:
jsteube
parent
e137e0fde6
commit
332b3c35e5
552
OpenCL/m12900.cl
Normal file
552
OpenCL/m12900.cl
Normal file
@ -0,0 +1,552 @@
|
||||
/**
|
||||
* Author......: Jens Steube <jens.steube@gmail.com>
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define _ANDROIDFDE_SAMSUNG_
|
||||
|
||||
#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))) m12900_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))) m12900_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))) m12900_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 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
w0[0] = tmps[gid].out[0];
|
||||
w0[1] = tmps[gid].out[1];
|
||||
w0[2] = tmps[gid].out[2];
|
||||
w0[3] = tmps[gid].out[3];
|
||||
w1[0] = tmps[gid].out[4];
|
||||
w1[1] = tmps[gid].out[5];
|
||||
w1[2] = tmps[gid].out[6];
|
||||
w1[3] = tmps[gid].out[7];
|
||||
w2[0] = 0;
|
||||
w2[1] = 0;
|
||||
w2[2] = 0;
|
||||
w2[3] = 0;
|
||||
w3[0] = 0;
|
||||
w3[1] = 0;
|
||||
w3[2] = 0;
|
||||
w3[3] = 0;
|
||||
|
||||
u32 ipad[8];
|
||||
u32 opad[8];
|
||||
|
||||
hmac_sha256_pad (w0, w1, w2, w3, ipad, opad);
|
||||
|
||||
w0[0] = salt_bufs[salt_pos].salt_buf[ 4];
|
||||
w0[1] = salt_bufs[salt_pos].salt_buf[ 5];
|
||||
w0[2] = salt_bufs[salt_pos].salt_buf[ 6];
|
||||
w0[3] = salt_bufs[salt_pos].salt_buf[ 7];
|
||||
w1[0] = salt_bufs[salt_pos].salt_buf[ 8];
|
||||
w1[1] = salt_bufs[salt_pos].salt_buf[ 9];
|
||||
w1[2] = salt_bufs[salt_pos].salt_buf[10];
|
||||
w1[3] = salt_bufs[salt_pos].salt_buf[11];
|
||||
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;
|
||||
|
||||
u32 digest[8];
|
||||
|
||||
hmac_sha256_run (w0, w1, w2, w3, ipad, opad, digest);
|
||||
|
||||
const u32 r0 = digest[0];
|
||||
const u32 r1 = digest[1];
|
||||
const u32 r2 = digest[2];
|
||||
const u32 r3 = digest[3];
|
||||
|
||||
#define il_pos 0
|
||||
|
||||
#include COMPARE_M
|
||||
}
|
||||
@ -488,6 +488,8 @@ typedef struct
|
||||
u32 digest_buf[8];
|
||||
#elif defined _MS_DRSR_
|
||||
u32 digest_buf[8];
|
||||
#elif defined _ANDROIDFDE_SAMSUNG_
|
||||
u32 digest_buf[8];
|
||||
#endif
|
||||
|
||||
} digest_t;
|
||||
|
||||
@ -8,6 +8,11 @@ Type.: Feature
|
||||
File.: Kernel
|
||||
Desc.: Extended support from 14 to 255 functions calls per rule on GPU
|
||||
|
||||
Type.: Feature
|
||||
File.: Kernel
|
||||
Desc.: Added new hash mode -m 12900 = Android FDE (Samsung DEK)
|
||||
Issue: 86
|
||||
|
||||
Type.: Feature
|
||||
File.: Host
|
||||
Desc.: Added a Makefile target for native compilation
|
||||
|
||||
@ -204,6 +204,7 @@ AMD users require Catalyst 14.9 or later (recommended 15.9 or later)
|
||||
- Password Safe v3
|
||||
- eCryptfs
|
||||
- Android FDE <= 4.3
|
||||
- Android FDE (Samsung DEK)
|
||||
- TrueCrypt 5.0+
|
||||
|
||||
##
|
||||
|
||||
@ -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_
|
||||
#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_
|
||||
/**
|
||||
* 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_
|
||||
#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_
|
||||
|
||||
#define SHIFT_RIGHT_32(x,n) ((x) >> (n))
|
||||
|
||||
|
||||
@ -351,6 +351,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERNEL_ACCEL_12600 32
|
||||
#define KERNEL_ACCEL_12700 64
|
||||
#define KERNEL_ACCEL_12800 64
|
||||
#define KERNEL_ACCEL_12900 8
|
||||
|
||||
#define KERNEL_LOOPS_0 256
|
||||
#define KERNEL_LOOPS_10 256
|
||||
@ -528,6 +529,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERNEL_LOOPS_12600 32
|
||||
#define KERNEL_LOOPS_12700 10
|
||||
#define KERNEL_LOOPS_12800 100
|
||||
#define KERNEL_LOOPS_12900 64
|
||||
|
||||
/**
|
||||
* Strings
|
||||
@ -673,6 +675,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define HT_12600 "ColdFusion 10+"
|
||||
#define HT_12700 "Blockchain, My Wallet"
|
||||
#define HT_12800 "MS-AzureSync PBKDF2-HMAC-SHA256"
|
||||
#define HT_12900 "Android FDE (Samsung DEK)"
|
||||
|
||||
#define HT_00011 "Joomla < 2.5.18"
|
||||
#define HT_00012 "PostgreSQL"
|
||||
@ -1002,6 +1005,8 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define DISPLAY_LEN_MAX_12700 1 + 10 + 1 + 5 + 1 + 20000
|
||||
#define DISPLAY_LEN_MIN_12800 11 + 1 + 20 + 1 + 1 + 1 + 64
|
||||
#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_11 32 + 1 + 16
|
||||
#define DISPLAY_LEN_MAX_11 32 + 1 + 32
|
||||
@ -1258,6 +1263,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define KERN_TYPE_CF10 12600
|
||||
#define KERN_TYPE_MYWALLET 12700
|
||||
#define KERN_TYPE_MS_DRSR 12800
|
||||
#define KERN_TYPE_ANDROIDFDE_SAMSUNG 12900
|
||||
|
||||
/**
|
||||
* signatures
|
||||
@ -1375,6 +1381,7 @@ extern hc_thread_mutex_t mux_display;
|
||||
#define ROUNDS_RAR3 262144
|
||||
#define ROUNDS_MYWALLET 10
|
||||
#define ROUNDS_MS_DRSR 100
|
||||
#define ROUNDS_ANDROIDFDE_SAMSUNG 4096
|
||||
|
||||
/**
|
||||
* salt types
|
||||
@ -1857,6 +1864,7 @@ int rar3hp_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);
|
||||
int androidfde_samsung_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf);
|
||||
|
||||
void load_kernel (const char *kernel_file, int num_devices, size_t *kernel_lengths, const unsigned char **kernel_sources);
|
||||
void writeProgramBin (char *dst, unsigned char *binary, size_t binary_size);
|
||||
|
||||
@ -128,7 +128,7 @@ const uint RESTORE_MIN = 210;
|
||||
|
||||
#define MAX_DICTSTAT 10000
|
||||
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 128
|
||||
#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 129
|
||||
|
||||
#define global_free(attr) \
|
||||
{ \
|
||||
@ -254,6 +254,7 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
|
||||
6231,
|
||||
6241,
|
||||
8800,
|
||||
12900,
|
||||
12200,
|
||||
9700,
|
||||
9710,
|
||||
@ -679,6 +680,7 @@ const char *USAGE_BIG[] =
|
||||
" Y = 2 = XTS 1024 bit (Ciphers: AES or Serpent or Twofish or AES-Twofish or Serpent-AES or Twofish-Serpent)",
|
||||
" Y = 3 = XTS 1536 bit (Ciphers: All)",
|
||||
" 8800 = Android FDE < v4.3",
|
||||
" 12900 = Android FDE (Samsung DEK)",
|
||||
" 12200 = eCryptfs",
|
||||
"",
|
||||
"[[ Documents ]]",
|
||||
@ -5608,7 +5610,7 @@ int main (int argc, char **argv)
|
||||
return (-1);
|
||||
}
|
||||
|
||||
if (hash_mode_chgd && hash_mode > 12800) // just added to remove compiler warnings for hash_mode_chgd
|
||||
if (hash_mode_chgd && hash_mode > 12900) // just added to remove compiler warnings for hash_mode_chgd
|
||||
{
|
||||
log_error ("ERROR: Invalid hash-type specified");
|
||||
|
||||
@ -9775,6 +9777,21 @@ int main (int argc, char **argv)
|
||||
dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
case 12900: 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_ANDROIDFDE_SAMSUNG;
|
||||
dgst_size = DGST_SIZE_4_8;
|
||||
parse_func = androidfde_samsung_parse_hash;
|
||||
sort_by_digest = sort_by_digest_4_8;
|
||||
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);
|
||||
}
|
||||
|
||||
@ -11129,6 +11146,8 @@ int main (int argc, char **argv)
|
||||
break;
|
||||
case 12800: hashes_buf[0].salt->salt_iter = ROUNDS_MS_DRSR - 1;
|
||||
break;
|
||||
case 12900: hashes_buf[0].salt->salt_iter = ROUNDS_ANDROIDFDE_SAMSUNG - 1;
|
||||
break;
|
||||
}
|
||||
|
||||
// set special tuning for benchmark-mode 1
|
||||
@ -11311,6 +11330,9 @@ int main (int argc, char **argv)
|
||||
case 12800: kernel_loops = ROUNDS_MS_DRSR;
|
||||
kernel_accel = 512;
|
||||
break;
|
||||
case 12900: kernel_loops = ROUNDS_ANDROIDFDE_SAMSUNG;
|
||||
kernel_accel = 8;
|
||||
break;
|
||||
}
|
||||
|
||||
// some algorithm collide too fast, make that impossible
|
||||
@ -13089,6 +13111,7 @@ int main (int argc, char **argv)
|
||||
case 12500: size_tmps = kernel_blocks * sizeof (rar3_tmp_t); break;
|
||||
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;
|
||||
};
|
||||
|
||||
uint size_hooks = 4;
|
||||
|
||||
83
src/shared.c
83
src/shared.c
@ -5610,6 +5610,7 @@ char *strhashtype (const uint hash_mode)
|
||||
case 12600: return ((char *) HT_12600); break;
|
||||
case 12700: return ((char *) HT_12700); break;
|
||||
case 12800: return ((char *) HT_12800); break;
|
||||
case 12900: return ((char *) HT_12900); break;
|
||||
}
|
||||
|
||||
return ((char *) "Unknown");
|
||||
@ -8042,6 +8043,31 @@ void ascii_digest (char out_buf[4096], uint salt_pos, uint digest_pos)
|
||||
byte_swap_32 (digest_buf[7])
|
||||
);
|
||||
}
|
||||
else if (hash_mode == 12900)
|
||||
{
|
||||
snprintf (out_buf, len-1, "%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x",
|
||||
salt.salt_buf[ 4],
|
||||
salt.salt_buf[ 5],
|
||||
salt.salt_buf[ 6],
|
||||
salt.salt_buf[ 7],
|
||||
salt.salt_buf[ 8],
|
||||
salt.salt_buf[ 9],
|
||||
salt.salt_buf[10],
|
||||
salt.salt_buf[11],
|
||||
byte_swap_32 (digest_buf[0]),
|
||||
byte_swap_32 (digest_buf[1]),
|
||||
byte_swap_32 (digest_buf[2]),
|
||||
byte_swap_32 (digest_buf[3]),
|
||||
byte_swap_32 (digest_buf[4]),
|
||||
byte_swap_32 (digest_buf[5]),
|
||||
byte_swap_32 (digest_buf[6]),
|
||||
byte_swap_32 (digest_buf[7]),
|
||||
salt.salt_buf[ 0],
|
||||
salt.salt_buf[ 1],
|
||||
salt.salt_buf[ 2],
|
||||
salt.salt_buf[ 3]
|
||||
);
|
||||
}
|
||||
else
|
||||
{
|
||||
if (hash_type == HASH_TYPE_MD4)
|
||||
@ -8996,6 +9022,7 @@ uint set_kernel_accel (uint hash_mode)
|
||||
case 12600: return GET_ACCEL (12600);
|
||||
case 12700: return GET_ACCEL (12700);
|
||||
case 12800: return GET_ACCEL (12800);
|
||||
case 12900: return GET_ACCEL (12900);
|
||||
}
|
||||
|
||||
return 0;
|
||||
@ -9181,6 +9208,7 @@ uint set_kernel_loops (uint hash_mode)
|
||||
case 12600: return GET_LOOPS (12600);
|
||||
case 12700: return GET_LOOPS (12700);
|
||||
case 12800: return GET_LOOPS (12800);
|
||||
case 12900: return GET_LOOPS (12900);
|
||||
}
|
||||
|
||||
return 0;
|
||||
@ -18570,6 +18598,61 @@ int ms_drsr_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int androidfde_samsung_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_12900) || (input_len > DISPLAY_LEN_MAX_12900)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
uint32_t *digest = (uint32_t *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
/**
|
||||
* parse line
|
||||
*/
|
||||
|
||||
char *hash_pos = input_buf + 64;
|
||||
char *salt1_pos = input_buf + 128;
|
||||
char *salt2_pos = input_buf;
|
||||
|
||||
/**
|
||||
* salt
|
||||
*/
|
||||
|
||||
salt->salt_buf[ 0] = hex_to_uint (&salt1_pos[ 0]);
|
||||
salt->salt_buf[ 1] = hex_to_uint (&salt1_pos[ 8]);
|
||||
salt->salt_buf[ 2] = hex_to_uint (&salt1_pos[16]);
|
||||
salt->salt_buf[ 3] = hex_to_uint (&salt1_pos[24]);
|
||||
|
||||
salt->salt_buf[ 4] = hex_to_uint (&salt2_pos[ 0]);
|
||||
salt->salt_buf[ 5] = hex_to_uint (&salt2_pos[ 8]);
|
||||
salt->salt_buf[ 6] = hex_to_uint (&salt2_pos[16]);
|
||||
salt->salt_buf[ 7] = hex_to_uint (&salt2_pos[24]);
|
||||
|
||||
salt->salt_buf[ 8] = hex_to_uint (&salt2_pos[32]);
|
||||
salt->salt_buf[ 9] = hex_to_uint (&salt2_pos[40]);
|
||||
salt->salt_buf[10] = hex_to_uint (&salt2_pos[48]);
|
||||
salt->salt_buf[11] = hex_to_uint (&salt2_pos[56]);
|
||||
|
||||
salt->salt_len = 48;
|
||||
|
||||
salt->salt_iter = ROUNDS_ANDROIDFDE_SAMSUNG - 1;
|
||||
|
||||
/**
|
||||
* digest buf
|
||||
*/
|
||||
|
||||
digest[0] = hex_to_uint (&hash_pos[ 0]);
|
||||
digest[1] = hex_to_uint (&hash_pos[ 8]);
|
||||
digest[2] = hex_to_uint (&hash_pos[16]);
|
||||
digest[3] = hex_to_uint (&hash_pos[24]);
|
||||
digest[4] = hex_to_uint (&hash_pos[32]);
|
||||
digest[5] = hex_to_uint (&hash_pos[40]);
|
||||
digest[6] = hex_to_uint (&hash_pos[48]);
|
||||
digest[7] = hex_to_uint (&hash_pos[56]);
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
/**
|
||||
* parallel running threads
|
||||
*/
|
||||
|
||||
@ -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);
|
||||
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 %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);
|
||||
@ -2160,6 +2160,21 @@ sub verify
|
||||
|
||||
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
|
||||
}
|
||||
elsif ($mode == 12900)
|
||||
{
|
||||
($hash_in, $word) = split ":", $line;
|
||||
|
||||
next unless defined $hash_in;
|
||||
next unless defined $word;
|
||||
|
||||
next unless length $hash_in == 160;
|
||||
|
||||
$param = substr ($hash_in, 0, 64);
|
||||
$salt = substr ($hash_in, 128, 32);
|
||||
$iter = 4096;
|
||||
|
||||
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: hash mode is not supported\n";
|
||||
@ -2426,6 +2441,14 @@ sub verify
|
||||
|
||||
$hash_out = $hash_in;
|
||||
}
|
||||
elsif ($mode == 12900)
|
||||
{
|
||||
$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);
|
||||
@ -2838,6 +2861,10 @@ sub passthrough
|
||||
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 20));
|
||||
}
|
||||
elsif ($mode == 12900)
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
|
||||
}
|
||||
else
|
||||
{
|
||||
print "ERROR: Unsupported hash type\n";
|
||||
@ -3537,6 +3564,20 @@ sub single
|
||||
}
|
||||
}
|
||||
}
|
||||
elsif ($mode == 12900)
|
||||
{
|
||||
for (my $i = 1; $i < 32; $i++)
|
||||
{
|
||||
if ($len != 0)
|
||||
{
|
||||
rnd ($mode, $len, 32);
|
||||
}
|
||||
else
|
||||
{
|
||||
rnd ($mode, $i, 32);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -6444,6 +6485,39 @@ END_CODE
|
||||
|
||||
$tmp_hash = sprintf ("v1;PPH1_MD4,%s,%d,%s", $salt_buf, $iterations, unpack ("H*", $hash));
|
||||
}
|
||||
elsif ($mode == 12900)
|
||||
{
|
||||
my $iterations = 4096;
|
||||
|
||||
if (length ($iter))
|
||||
{
|
||||
$iterations = int ($iter);
|
||||
}
|
||||
|
||||
my $salt2 = $salt_buf . $salt_buf;
|
||||
|
||||
if (defined $additional_param)
|
||||
{
|
||||
$salt2 = $additional_param;
|
||||
}
|
||||
|
||||
my $pbkdf2 = Crypt::PBKDF2->new
|
||||
(
|
||||
hasher => Crypt::PBKDF2->hasher_from_algorithm ('HMACSHA2', 256),
|
||||
iterations => $iterations,
|
||||
output_len => 32
|
||||
);
|
||||
|
||||
my $salt_buf_bin = pack ("H*", $salt_buf);
|
||||
|
||||
my $hash = $pbkdf2->PBKDF2 ($salt_buf_bin, $word_buf);
|
||||
|
||||
my $salt2_bin = pack ("H*", $salt2);
|
||||
|
||||
my $hash_hmac = hmac_hex ($salt2_bin, $hash, \&sha256, 64);
|
||||
|
||||
$tmp_hash = sprintf ("%s%s%s", $salt2, $hash_hmac, $salt_buf);
|
||||
}
|
||||
|
||||
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"
|
||||
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"
|
||||
|
||||
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"
|
||||
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"
|
||||
|
||||
OPTS="--quiet --force --potfile-disable --runtime 200 --gpu-temp-disable -d 1 --weak-hash-threshold=0"
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user