Added hash-mode 15100 = Juniper/NetBSD sha1crypt

pull/1205/head
jsteube 7 years ago
parent 826de76bd6
commit d1b2fa0b31

@ -675,9 +675,10 @@ typedef struct
typedef struct
{
u32 salt_buf[16];
u32 salt_buf_pc[8];
u32 salt_buf_pc[16];
u32 salt_len;
u32 salt_len_pc;
u32 salt_iter;
u32 salt_iter2;
u32 salt_sign[2];
@ -1165,6 +1166,17 @@ typedef struct
} tc64_tmp_t;
typedef struct
{
// pbkdf1-sha1 is limited to 160 bits
u32 ipad[5];
u32 opad[5];
u32 out[5];
} pbkdf1_sha1_tmp_t;
typedef struct
{
u32 ipad[4];

@ -0,0 +1,630 @@
/**
* 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"
#define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl"
void sha1_transform_S (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5])
{
u32 A = digest[0];
u32 B = digest[1];
u32 C = digest[2];
u32 D = digest[3];
u32 E = digest[4];
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];
#undef K
#define K SHA1C00
SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, w0_t);
SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w1_t);
SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w2_t);
SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w3_t);
SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w4_t);
SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, w5_t);
SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w6_t);
SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w7_t);
SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w8_t);
SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w9_t);
SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, wa_t);
SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, wb_t);
SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, wc_t);
SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, wd_t);
SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, we_t);
SHA1_STEP_S (SHA1_F0o, A, B, C, D, E, wf_t);
w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F0o, E, A, B, C, D, w0_t);
w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F0o, D, E, A, B, C, w1_t);
w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F0o, C, D, E, A, B, w2_t);
w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F0o, B, C, D, E, A, w3_t);
#undef K
#define K SHA1C01
w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w4_t);
w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w5_t);
w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w6_t);
w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w7_t);
w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w8_t);
w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w9_t);
wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wa_t);
wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, wb_t);
wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, wc_t);
wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wd_t);
we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, we_t);
wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wf_t);
w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w0_t);
w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w1_t);
w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w2_t);
w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w3_t);
w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w4_t);
w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w5_t);
w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w6_t);
w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w7_t);
#undef K
#define K SHA1C02
w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w8_t);
w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w9_t);
wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, wa_t);
wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, wb_t);
wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, wc_t);
wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, wd_t);
we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, we_t);
wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, wf_t);
w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, w0_t);
w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, w1_t);
w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w2_t);
w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w3_t);
w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, w4_t);
w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, w5_t);
w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, w6_t);
w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F2o, A, B, C, D, E, w7_t);
w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F2o, E, A, B, C, D, w8_t);
w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F2o, D, E, A, B, C, w9_t);
wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F2o, C, D, E, A, B, wa_t);
wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F2o, B, C, D, E, A, wb_t);
#undef K
#define K SHA1C03
wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, wc_t);
wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wd_t);
we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, we_t);
wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, wf_t);
w0_t = rotl32_S ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w0_t);
w1_t = rotl32_S ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w1_t);
w2_t = rotl32_S ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w2_t);
w3_t = rotl32_S ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w3_t);
w4_t = rotl32_S ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w4_t);
w5_t = rotl32_S ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, w5_t);
w6_t = rotl32_S ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, w6_t);
w7_t = rotl32_S ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, w7_t);
w8_t = rotl32_S ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, w8_t);
w9_t = rotl32_S ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, w9_t);
wa_t = rotl32_S ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wa_t);
wb_t = rotl32_S ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP_S (SHA1_F1, A, B, C, D, E, wb_t);
wc_t = rotl32_S ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP_S (SHA1_F1, E, A, B, C, D, wc_t);
wd_t = rotl32_S ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP_S (SHA1_F1, D, E, A, B, C, wd_t);
we_t = rotl32_S ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP_S (SHA1_F1, C, D, E, A, B, we_t);
wf_t = rotl32_S ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP_S (SHA1_F1, B, C, D, E, A, wf_t);
digest[0] += A;
digest[1] += B;
digest[2] += C;
digest[3] += D;
digest[4] += E;
}
void hmac_sha1_pad_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[5], u32 opad[5])
{
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] = SHA1M_A;
ipad[1] = SHA1M_B;
ipad[2] = SHA1M_C;
ipad[3] = SHA1M_D;
ipad[4] = SHA1M_E;
sha1_transform_S (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] = SHA1M_A;
opad[1] = SHA1M_B;
opad[2] = SHA1M_C;
opad[3] = SHA1M_D;
opad[4] = SHA1M_E;
sha1_transform_S (w0, w1, w2, w3, opad);
}
void hmac_sha1_run_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[5], u32 opad[5], u32 digest[5])
{
digest[0] = ipad[0];
digest[1] = ipad[1];
digest[2] = ipad[2];
digest[3] = ipad[3];
digest[4] = ipad[4];
sha1_transform_S (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] = 0x80000000;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = (64 + 20) * 8;
digest[0] = opad[0];
digest[1] = opad[1];
digest[2] = opad[2];
digest[3] = opad[3];
digest[4] = opad[4];
sha1_transform_S (w0, w1, w2, w3, digest);
}
void sha1_transform_V (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{
u32x A = digest[0];
u32x B = digest[1];
u32x C = digest[2];
u32x D = digest[3];
u32x E = digest[4];
u32x w0_t = w0[0];
u32x w1_t = w0[1];
u32x w2_t = w0[2];
u32x w3_t = w0[3];
u32x w4_t = w1[0];
u32x w5_t = w1[1];
u32x w6_t = w1[2];
u32x w7_t = w1[3];
u32x w8_t = w2[0];
u32x w9_t = w2[1];
u32x wa_t = w2[2];
u32x wb_t = w2[3];
u32x wc_t = w3[0];
u32x wd_t = w3[1];
u32x we_t = w3[2];
u32x wf_t = w3[3];
#undef K
#define K SHA1C00
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
#undef K
#define K SHA1C01
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
#undef K
#define K SHA1C02
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
#undef K
#define K SHA1C03
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
digest[0] += A;
digest[1] += B;
digest[2] += C;
digest[3] += D;
digest[4] += E;
}
void hmac_sha1_pad_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5])
{
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] = SHA1M_A;
ipad[1] = SHA1M_B;
ipad[2] = SHA1M_C;
ipad[3] = SHA1M_D;
ipad[4] = SHA1M_E;
sha1_transform_V (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] = SHA1M_A;
opad[1] = SHA1M_B;
opad[2] = SHA1M_C;
opad[3] = SHA1M_D;
opad[4] = SHA1M_E;
sha1_transform_V (w0, w1, w2, w3, opad);
}
void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5])
{
digest[0] = ipad[0];
digest[1] = ipad[1];
digest[2] = ipad[2];
digest[3] = ipad[3];
digest[4] = ipad[4];
sha1_transform_V (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] = 0x80000000;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = (64 + 20) * 8;
digest[0] = opad[0];
digest[1] = opad[1];
digest[2] = opad[2];
digest[3] = opad[3];
digest[4] = opad[4];
sha1_transform_V (w0, w1, w2, w3, digest);
}
__kernel void m15100_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global pbkdf1_sha1_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
*/
const u32 gid = get_global_id (0);
if (gid >= gid_max) return;
u32 w0[4];
w0[0] = swap32_S (pws[gid].i[ 0]);
w0[1] = swap32_S (pws[gid].i[ 1]);
w0[2] = swap32_S (pws[gid].i[ 2]);
w0[3] = swap32_S (pws[gid].i[ 3]);
u32 w1[4];
w1[0] = swap32_S (pws[gid].i[ 4]);
w1[1] = swap32_S (pws[gid].i[ 5]);
w1[2] = swap32_S (pws[gid].i[ 6]);
w1[3] = swap32_S (pws[gid].i[ 7]);
u32 w2[4];
w2[0] = swap32_S (pws[gid].i[ 8]);
w2[1] = swap32_S (pws[gid].i[ 9]);
w2[2] = swap32_S (pws[gid].i[10]);
w2[3] = swap32_S (pws[gid].i[11]);
u32 w3[4];
w3[0] = swap32_S (pws[gid].i[12]);
w3[1] = swap32_S (pws[gid].i[13]);
w3[2] = swap32_S (pws[gid].i[14]);
w3[3] = swap32_S (pws[gid].i[15]);
/**
* salt
*/
const u32 salt_len = salt_bufs[salt_pos].salt_len_pc;
u32 salt_buf0[4];
u32 salt_buf1[4];
u32 salt_buf2[4];
u32 salt_buf3[4];
salt_buf0[0] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 0]);
salt_buf0[1] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 1]);
salt_buf0[2] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 2]);
salt_buf0[3] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 3]);
salt_buf1[0] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 4]);
salt_buf1[1] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 5]);
salt_buf1[2] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 6]);
salt_buf1[3] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 7]);
salt_buf2[0] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 8]);
salt_buf2[1] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[ 9]);
salt_buf2[2] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[10]);
salt_buf2[3] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[11]);
salt_buf3[0] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[12]);
salt_buf3[1] = swap32_S (salt_bufs[salt_pos].salt_buf_pc[13]);
salt_buf3[2] = 0;
salt_buf3[3] = (64 + salt_len) * 8;
u32 ipad[5];
u32 opad[5];
hmac_sha1_pad_S (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].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];
u32 out[5];
hmac_sha1_run_S (salt_buf0, salt_buf1, salt_buf2, salt_buf3, ipad, opad, out);
tmps[gid].out[0] = out[0];
tmps[gid].out[1] = out[1];
tmps[gid].out[2] = out[2];
tmps[gid].out[3] = out[3];
tmps[gid].out[4] = out[4];
}
__kernel void m15100_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global pbkdf1_sha1_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
const u32 gid = get_global_id (0);
if ((gid * VECT_SIZE) >= gid_max) return;
u32x ipad[5];
u32x opad[5];
ipad[0] = packv (tmps, ipad, gid, 0);
ipad[1] = packv (tmps, ipad, gid, 1);
ipad[2] = packv (tmps, ipad, gid, 2);
ipad[3] = packv (tmps, ipad, gid, 3);
ipad[4] = packv (tmps, ipad, gid, 4);
opad[0] = packv (tmps, opad, gid, 0);
opad[1] = packv (tmps, opad, gid, 1);
opad[2] = packv (tmps, opad, gid, 2);
opad[3] = packv (tmps, opad, gid, 3);
opad[4] = packv (tmps, opad, gid, 4);
u32x out[5];
out[0] = packv (tmps, out, gid, 0);
out[1] = packv (tmps, out, gid, 1);
out[2] = packv (tmps, out, gid, 2);
out[3] = packv (tmps, out, gid, 3);
out[4] = packv (tmps, out, gid, 4);
for (u32 j = 0; j < loop_cnt; j++)
{
u32x w0[4];
u32x w1[4];
u32x w2[4];
u32x w3[4];
w0[0] = out[0];
w0[1] = out[1];
w0[2] = out[2];
w0[3] = out[3];
w1[0] = out[4];
w1[1] = 0x80000000;
w1[2] = 0;
w1[3] = 0;
w2[0] = 0;
w2[1] = 0;
w2[2] = 0;
w2[3] = 0;
w3[0] = 0;
w3[1] = 0;
w3[2] = 0;
w3[3] = (64 + 20) * 8;
hmac_sha1_run_V (w0, w1, w2, w3, ipad, opad, out);
}
unpackv (tmps, out, gid, 0, out[0]);
unpackv (tmps, out, gid, 1, out[1]);
unpackv (tmps, out, gid, 2, out[2]);
unpackv (tmps, out, gid, 3, out[3]);
unpackv (tmps, out, gid, 4, out[4]);
}
__kernel void m15100_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global pbkdf1_sha1_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
* base
*/
const u32 gid = get_global_id (0);
if (gid >= gid_max) return;
const u32 lid = get_local_id (0);
const u32 r0 = tmps[gid].out[DGST_R0];
const u32 r1 = tmps[gid].out[DGST_R1];
const u32 r2 = tmps[gid].out[DGST_R2];
const u32 r3 = tmps[gid].out[DGST_R3];
#define il_pos 0
#include COMPARE_M
}

@ -9,6 +9,12 @@
- WPA cracking: Added support for a special bit (bit 8) of the message_pair that indicates if replay counters did match
- Added visual indicator to the status screen to indicate if the checkpoint quit feature is enabled
##
## Algorithms
##
- Added hash-mode 15100 = Juniper/NetBSD sha1crypt
##
## Improvements
##

@ -172,6 +172,7 @@ NVidia users require "NVIDIA Driver" (367.x or later)
- Cisco $9$
- Juniper IVE
- Juniper Netscreen/SSG (ScreenOS)
- Juniper/NetBSD sha1crypt
- Fortigate (FortiOS)
- Android PIN
- Windows 8+ phone PIN/Password

@ -176,7 +176,7 @@ _hashcat ()
{
local VERSION=3.40
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 200 300 400 500 501 900 1000 1100 1400 1410 1411 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 3910 4010 4110 4300 4400 4500 4520 4521 4522 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 7000 7100 7200 7300 7400 7500 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 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000"
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 200 300 400 500 501 900 1000 1100 1400 1410 1411 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 3910 4010 4110 4300 4400 4500 4520 4521 4522 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 7000 7100 7200 7300 7400 7500 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 12001 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14700 14800 14900 15000 15100"
local ATTACK_MODES="0 1 3 6 7"
local HCCAPX_MESSAGE_PAIR="0 1 2 3 4 5"
local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15"

@ -667,6 +667,17 @@ typedef struct saph_sha1_tmp
} saph_sha1_tmp_t;
typedef struct pbkdf1_sha1_tmp
{
// pbkdf1-sha1 is limited to 160 bits
u32 ipad[5];
u32 opad[5];
u32 out[5];
} pbkdf1_sha1_tmp_t;
typedef struct pbkdf2_md5_tmp
{
u32 ipad[4];
@ -1173,6 +1184,8 @@ typedef enum display_len
DISPLAY_LEN_MAX_14900 = 8 + 1 + 8,
DISPLAY_LEN_MIN_15000 = 128 + 1 + 64,
DISPLAY_LEN_MAX_15000 = 128 + 1 + 64,
DISPLAY_LEN_MIN_15100 = 6 + 3 + 1 + 8 + 1 + 28,
DISPLAY_LEN_MAX_15100 = 6 + 6 + 1 + 8 + 1 + 28,
DISPLAY_LEN_MIN_99999 = 1,
DISPLAY_LEN_MAX_99999 = 55,
@ -1498,6 +1511,7 @@ typedef enum kern_type
KERN_TYPE_ITUNES_BACKUP_10 = 14800,
KERN_TYPE_SKIP32 = 14900,
KERN_TYPE_FILEZILLA_SERVER = 15000,
KERN_TYPE_NETBSD_SHA1CRYPT = 15100,
KERN_TYPE_PLAINTEXT = 99999,
} kern_type_t;
@ -1566,6 +1580,7 @@ typedef enum rounds_count
ROUNDS_ITUNES101_BACKUP = 10000000, // wtf, i mean, really?
ROUNDS_ITUNES102_BACKUP = 10000,
ROUNDS_ATLASSIAN = 10000,
ROUNDS_NETBSD_SHA1CRYPT = 20000,
ROUNDS_STDOUT = 0
} rounds_count_t;
@ -1740,6 +1755,7 @@ int skip32_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu
int fortigate_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int sha256b64s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int filezilla_server_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int netbsd_sha1crypt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
int atlassian_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
/**

@ -651,26 +651,27 @@ typedef enum user_options_map
* structs
*/
typedef struct salt
typedef struct
{
u32 salt_buf[16];
u32 salt_buf_pc[8];
u32 salt_buf[16];
u32 salt_buf_pc[16];
u32 salt_len;
u32 salt_iter;
u32 salt_iter2;
u32 salt_sign[2];
u32 salt_len;
u32 salt_len_pc;
u32 salt_iter;
u32 salt_iter2;
u32 salt_sign[2];
u32 keccak_mdlen;
u32 keccak_mdlen;
u32 digests_cnt;
u32 digests_done;
u32 digests_cnt;
u32 digests_done;
u32 digests_offset;
u32 digests_offset;
u32 scrypt_N;
u32 scrypt_r;
u32 scrypt_p;
u32 scrypt_N;
u32 scrypt_r;
u32 scrypt_p;
} salt_t;

@ -238,6 +238,7 @@ static const char HT_14700[] = "iTunes Backup < 10.0";
static const char HT_14800[] = "iTunes Backup >= 10.0";
static const char HT_14900[] = "Skip32";
static const char HT_15000[] = "FileZilla Server >= 0.9.55";
static const char HT_15100[] = "Juniper/NetBSD sha1crypt";
static const char HT_99999[] = "Plaintext";
static const char HT_00011[] = "Joomla < 2.5.18";
@ -300,81 +301,81 @@ static const char HT_13762[] = "VeraCrypt PBKDF2-HMAC-SHA256 + XTS 1024 bit + bo
static const char HT_13763[] = "VeraCrypt PBKDF2-HMAC-SHA256 + XTS 1536 bit + boot-mode";
static const char HT_12001[] = "Atlassian (PBKDF2-HMAC-SHA1)";
static const char SIGNATURE_ANDROIDFDE[] = "$fde$";
static const char SIGNATURE_AXCRYPT[] = "$axcrypt$*1";
static const char SIGNATURE_AXCRYPT_SHA1[] = "$axcrypt_sha1";
static const char SIGNATURE_BCRYPT1[] = "$2a$";
static const char SIGNATURE_BCRYPT2[] = "$2b$";
static const char SIGNATURE_BCRYPT3[] = "$2x$";
static const char SIGNATURE_BCRYPT4[] = "$2y$";
static const char SIGNATURE_BITCOIN_WALLET[] = "$bitcoin$";
static const char SIGNATURE_BSDICRYPT[] = "_";
static const char SIGNATURE_CISCO8[] = "$8$";
static const char SIGNATURE_CISCO9[] = "$9$";
static const char SIGNATURE_CRAM_MD5[] = "$cram_md5$";
static const char SIGNATURE_DCC2[] = "$DCC2$";
static const char SIGNATURE_DJANGOPBKDF2[] = "pbkdf2_sha256$";
static const char SIGNATURE_DJANGOSHA1[] = "sha1$";
static const char SIGNATURE_DRUPAL7[] = "$S$";
static const char SIGNATURE_ECRYPTFS[] = "$ecryptfs$";
static const char SIGNATURE_EPISERVER4[] = "$episerver$*1*";
static const char SIGNATURE_EPISERVER[] = "$episerver$*0*";
static const char SIGNATURE_KEEPASS[] = "$keepass$";
static const char SIGNATURE_KRB5PA[] = "$krb5pa$23";
static const char SIGNATURE_KRB5TGS[] = "$krb5tgs$23";
static const char SIGNATURE_MD5AIX[] = "{smd5}";
static const char SIGNATURE_MD5APR1[] = "$apr1$";
static const char SIGNATURE_MD5CRYPT[] = "$1$";
static const char SIGNATURE_MEDIAWIKI_B[] = "$B$";
static const char SIGNATURE_MS_DRSR[] = "v1;PPH1_MD4";
static const char SIGNATURE_MSSQL[] = "0x0100";
static const char SIGNATURE_MSSQL2012[] = "0x0200";
static const char SIGNATURE_MYSQL_AUTH[] = "$mysqlna$";
static const char SIGNATURE_MYWALLET[] = "$blockchain$";
static const char SIGNATURE_NETSCALER[] = "1";
static const char SIGNATURE_OFFICE2007[] = "$office$";
static const char SIGNATURE_OFFICE2010[] = "$office$";
static const char SIGNATURE_OFFICE2013[] = "$office$";
static const char SIGNATURE_OLDOFFICE0[] = "$oldoffice$0";
static const char SIGNATURE_OLDOFFICE1[] = "$oldoffice$1";
static const char SIGNATURE_OLDOFFICE3[] = "$oldoffice$3";
static const char SIGNATURE_OLDOFFICE4[] = "$oldoffice$4";
static const char SIGNATURE_PBKDF2_MD5[] = "md5:";
static const char SIGNATURE_PBKDF2_SHA1[] = "sha1:";
static const char SIGNATURE_PBKDF2_SHA256[] = "sha256:";
static const char SIGNATURE_PBKDF2_SHA512[] = "sha512:";
static const char SIGNATURE_PDF[] = "$pdf$";
static const char SIGNATURE_PHPASS1[] = "$P$";
static const char SIGNATURE_PHPASS2[] = "$H$";
static const char SIGNATURE_PHPS[] = "$PHPS$";
static const char SIGNATURE_POSTGRESQL_AUTH[] = "$postgres$";
static const char SIGNATURE_PSAFE3[] = "PWS3";
static const char SIGNATURE_RACF[] = "$racf$";
static const char SIGNATURE_RAR3[] = "$RAR3$";
static const char SIGNATURE_RAR5[] = "$rar5$";
static const char SIGNATURE_SAPH_SHA1[] = "{x-issha, ";
static const char SIGNATURE_SCRYPT[] = "SCRYPT";
static const char SIGNATURE_SEVEN_ZIP[] = "$7z$";
static const char SIGNATURE_SHA1AIX[] = "{ssha1}";
static const char SIGNATURE_SHA1B64[] = "{SHA}";
static const char SIGNATURE_SHA256AIX[] = "{ssha256}";
static const char SIGNATURE_SHA256B64S[] = "{SSHA256}";
static const char SIGNATURE_SHA256CRYPT[] = "$5$";
static const char SIGNATURE_SHA512AIX[] = "{ssha512}";
static const char SIGNATURE_SHA512B64S[] = "{SSHA512}";
static const char SIGNATURE_SHA512CRYPT[] = "$6$";
static const char SIGNATURE_SHA512GRUB[] = "grub.pbkdf2.sha512.";
static const char SIGNATURE_SHA512OSX[] = "$ml$";
static const char SIGNATURE_SIP_AUTH[] = "$sip$*";
static const char SIGNATURE_SSHA1B64_lower[] = "{ssha}";
static const char SIGNATURE_SSHA1B64_upper[] = "{SSHA}";
static const char SIGNATURE_SYBASEASE[] = "0xc007";
//static const char SIGNATURE_TRUECRYPT[] = "TRUE";
static const char SIGNATURE_ZIP2_START[] = "$zip2$";
static const char SIGNATURE_ZIP2_STOP[] = "$/zip2$";
static const char SIGNATURE_ITUNES_BACKUP[] = "$itunes_backup$";
static const char SIGNATURE_FORTIGATE[] = "AK1";
static const char SIGNATURE_ATLASSIAN[] = "{PKCS5S2}";
static const char SIGNATURE_ANDROIDFDE[] = "$fde$";
static const char SIGNATURE_AXCRYPT[] = "$axcrypt$*1";
static const char SIGNATURE_AXCRYPT_SHA1[] = "$axcrypt_sha1";
static const char SIGNATURE_BCRYPT1[] = "$2a$";
static const char SIGNATURE_BCRYPT2[] = "$2b$";
static const char SIGNATURE_BCRYPT3[] = "$2x$";
static const char SIGNATURE_BCRYPT4[] = "$2y$";
static const char SIGNATURE_BITCOIN_WALLET[] = "$bitcoin$";
static const char SIGNATURE_BSDICRYPT[] = "_";
static const char SIGNATURE_CISCO8[] = "$8$";
static const char SIGNATURE_CISCO9[] = "$9$";
static const char SIGNATURE_CRAM_MD5[] = "$cram_md5$";
static const char SIGNATURE_DCC2[] = "$DCC2$";
static const char SIGNATURE_DJANGOPBKDF2[] = "pbkdf2_sha256$";
static const char SIGNATURE_DJANGOSHA1[] = "sha1$";
static const char SIGNATURE_DRUPAL7[] = "$S$";
static const char SIGNATURE_ECRYPTFS[] = "$ecryptfs$";
static const char SIGNATURE_EPISERVER4[] = "$episerver$*1*";
static const char SIGNATURE_EPISERVER[] = "$episerver$*0*";
static const char SIGNATURE_KEEPASS[] = "$keepass$";
static const char SIGNATURE_KRB5PA[] = "$krb5pa$23";
static const char SIGNATURE_KRB5TGS[] = "$krb5tgs$23";
static const char SIGNATURE_MD5AIX[] = "{smd5}";
static const char SIGNATURE_MD5APR1[] = "$apr1$";
static const char SIGNATURE_MD5CRYPT[] = "$1$";
static const char SIGNATURE_MEDIAWIKI_B[] = "$B$";
static const char SIGNATURE_MS_DRSR[] = "v1;PPH1_MD4";
static const char SIGNATURE_MSSQL[] = "0x0100";
static const char SIGNATURE_MSSQL2012[] = "0x0200";
static const char SIGNATURE_MYSQL_AUTH[] = "$mysqlna$";
static const char SIGNATURE_MYWALLET[] = "$blockchain$";
static const char SIGNATURE_NETSCALER[] = "1";
static const char SIGNATURE_OFFICE2007[] = "$office$";
static const char SIGNATURE_OFFICE2010[] = "$office$";
static const char SIGNATURE_OFFICE2013[] = "$office$";
static const char SIGNATURE_OLDOFFICE0[] = "$oldoffice$0";
static const char SIGNATURE_OLDOFFICE1[] = "$oldoffice$1";
static const char SIGNATURE_OLDOFFICE3[] = "$oldoffice$3";
static const char SIGNATURE_OLDOFFICE4[] = "$oldoffice$4";
static const char SIGNATURE_PBKDF2_MD5[] = "md5:";
static const char SIGNATURE_PBKDF2_SHA1[] = "sha1:";
static const char SIGNATURE_PBKDF2_SHA256[] = "sha256:";
static const char SIGNATURE_PBKDF2_SHA512[] = "sha512:";
static const char SIGNATURE_PDF[] = "$pdf$";
static const char SIGNATURE_PHPASS1[] = "$P$";
static const char SIGNATURE_PHPASS2[] = "$H$";
static const char SIGNATURE_PHPS[] = "$PHPS$";
static const char SIGNATURE_POSTGRESQL_AUTH[] = "$postgres$";
static const char SIGNATURE_PSAFE3[] = "PWS3";
static const char SIGNATURE_RACF[] = "$racf$";
static const char SIGNATURE_RAR3[] = "$RAR3$";
static const char SIGNATURE_RAR5[] = "$rar5$";
static const char SIGNATURE_SAPH_SHA1[] = "{x-issha, ";
static const char SIGNATURE_SCRYPT[] = "SCRYPT";
static const char SIGNATURE_SEVEN_ZIP[] = "$7z$";
static const char SIGNATURE_SHA1AIX[] = "{ssha1}";
static const char SIGNATURE_SHA1B64[] = "{SHA}";
static const char SIGNATURE_SHA256AIX[] = "{ssha256}";
static const char SIGNATURE_SHA256B64S[] = "{SSHA256}";
static const char SIGNATURE_SHA256CRYPT[] = "$5$";
static const char SIGNATURE_SHA512AIX[] = "{ssha512}";
static const char SIGNATURE_SHA512B64S[] = "{SSHA512}";
static const char SIGNATURE_SHA512CRYPT[] = "$6$";
static const char SIGNATURE_SHA512GRUB[] = "grub.pbkdf2.sha512.";
static const char SIGNATURE_SHA512OSX[] = "$ml$";
static const char SIGNATURE_SIP_AUTH[] = "$sip$*";
static const char SIGNATURE_SSHA1B64_lower[] = "{ssha}";
static const char SIGNATURE_SSHA1B64_upper[] = "{SSHA}";
static const char SIGNATURE_SYBASEASE[] = "0xc007";
static const char SIGNATURE_ZIP2_START[] = "$zip2$";
static const char SIGNATURE_ZIP2_STOP[] = "$/zip2$";
static const char SIGNATURE_ITUNES_BACKUP[] = "$itunes_backup$";
static const char SIGNATURE_FORTIGATE[] = "AK1";
static const char SIGNATURE_ATLASSIAN[] = "{PKCS5S2}";
static const char SIGNATURE_NETBSD_SHA1CRYPT[] = "$sha1$";
/**
* decoder / encoder
@ -1628,6 +1629,129 @@ static void sha512aix_encode (u8 digest[64], u8 buf[86])
buf[85] = int_to_itoa64 (l & 0x3f); //l >>= 6;
}
static void netbsd_sha1crypt_decode (u8 digest[20], u8 buf[28], u8 *additional_byte)
{
int l;
l = itoa64_to_int (buf[ 0]) << 0;
l |= itoa64_to_int (buf[ 1]) << 6;
l |= itoa64_to_int (buf[ 2]) << 12;
l |= itoa64_to_int (buf[ 3]) << 18;
digest[ 2] = (l >> 0) & 0xff;
digest[ 1] = (l >> 8) & 0xff;
digest[ 0] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[ 4]) << 0;
l |= itoa64_to_int (buf[ 5]) << 6;
l |= itoa64_to_int (buf[ 6]) << 12;
l |= itoa64_to_int (buf[ 7]) << 18;
digest[ 5] = (l >> 0) & 0xff;
digest[ 4] = (l >> 8) & 0xff;
digest[ 3] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[ 8]) << 0;
l |= itoa64_to_int (buf[ 9]) << 6;
l |= itoa64_to_int (buf[10]) << 12;
l |= itoa64_to_int (buf[11]) << 18;
digest[ 8] = (l >> 0) & 0xff;
digest[ 7] = (l >> 8) & 0xff;
digest[ 6] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[12]) << 0;
l |= itoa64_to_int (buf[13]) << 6;
l |= itoa64_to_int (buf[14]) << 12;
l |= itoa64_to_int (buf[15]) << 18;
digest[11] = (l >> 0) & 0xff;
digest[10] = (l >> 8) & 0xff;
digest[ 9] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[16]) << 0;
l |= itoa64_to_int (buf[17]) << 6;
l |= itoa64_to_int (buf[18]) << 12;
l |= itoa64_to_int (buf[19]) << 18;
digest[14] = (l >> 0) & 0xff;
digest[13] = (l >> 8) & 0xff;
digest[12] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[20]) << 0;
l |= itoa64_to_int (buf[21]) << 6;
l |= itoa64_to_int (buf[22]) << 12;
l |= itoa64_to_int (buf[23]) << 18;
digest[17] = (l >> 0) & 0xff;
digest[16] = (l >> 8) & 0xff;
digest[15] = (l >> 16) & 0xff;
l = itoa64_to_int (buf[24]) << 0;
l |= itoa64_to_int (buf[25]) << 6;
l |= itoa64_to_int (buf[26]) << 12;
l |= itoa64_to_int (buf[27]) << 18;
additional_byte[0] = (l >> 0) & 0xff;
digest[19] = (l >> 8) & 0xff;
digest[18] = (l >> 16) & 0xff;
}
static void netbsd_sha1crypt_encode (u8 digest[20], u8 additional_byte, u8 buf[30])
{
int l;
l = (digest[ 2] << 0) | (digest[ 1] << 8) | (digest[ 0] << 16);
buf[ 0] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 1] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 2] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 3] = int_to_itoa64 (l & 0x3f);
l = (digest[ 5] << 0) | (digest[ 4] << 8) | (digest[ 3] << 16);
buf[ 4] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 5] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 6] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 7] = int_to_itoa64 (l & 0x3f);
l = (digest[ 8] << 0) | (digest[ 7] << 8) | (digest[ 6] << 16);
buf[ 8] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[ 9] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[10] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[11] = int_to_itoa64 (l & 0x3f);
l = (digest[11] << 0) | (digest[10] << 8) | (digest[ 9] << 16);
buf[12] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[13] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[14] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[15] = int_to_itoa64 (l & 0x3f);
l = (digest[14] << 0) | (digest[13] << 8) | (digest[12] << 16);
buf[16] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[17] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[18] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[19] = int_to_itoa64 (l & 0x3f);
l = (digest[17] << 0) | (digest[16] << 8) | (digest[15] << 16);
buf[20] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[21] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[22] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[23] = int_to_itoa64 (l & 0x3f);
l = (additional_byte << 0) | (digest[19] << 8) | (digest[18] << 16);
buf[24] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[25] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[26] = int_to_itoa64 (l & 0x3f); l >>= 6;
buf[27] = int_to_itoa64 (l & 0x3f);
buf[28] = 0;
}
static void sha256crypt_decode (u8 digest[32], u8 buf[43])
{
int l;
@ -14211,6 +14335,89 @@ int filezilla_server_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf,
return (PARSER_OK);
}
int netbsd_sha1crypt_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
{
if ((input_len < DISPLAY_LEN_MIN_15100) || (input_len > DISPLAY_LEN_MAX_15100)) return (PARSER_GLOBAL_LENGTH);
if (memcmp (SIGNATURE_NETBSD_SHA1CRYPT, input_buf, 6)) return (PARSER_SIGNATURE_UNMATCHED);
u32 *digest = (u32 *) hash_buf->digest;
salt_t *salt = hash_buf->salt;
u8 *iter_pos = input_buf + 6;
/**
* parse line
*/
u8 *salt_pos = (u8 *) strchr ((const char *) iter_pos, '$');
if (salt_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
salt_pos++;
u8 *hash_pos = (u8 *) strchr ((const char *) salt_pos, '$');
if (hash_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
u32 salt_len = hash_pos - salt_pos;
hash_pos++;
u32 hash_len = input_len - (hash_pos - input_buf);
/**
* verify data
*/
u32 iter = atoi ((const char *) iter_pos);
if (iter < 99) return (PARSER_SALT_ITERATION); // (actually: CRYPT_SHA1_ITERATIONS should be 24680 or more)
if (salt_len != 8) return (PARSER_SALT_LENGTH);
if (hash_len != 28) return (PARSER_HASH_LENGTH);
/**
* store data
*/
// iterations:
salt->salt_iter = iter - 1;
// salt:
memcpy ((u8 *) salt->salt_buf, salt_pos, salt_len);
// salt length:
salt->salt_len = salt_len;
// digest:
netbsd_sha1crypt_decode ((u8 *) digest, (u8 *) hash_pos, (u8 *) salt->salt_sign);
digest[0] = byte_swap_32 (digest[0]);
digest[1] = byte_swap_32 (digest[1]);
digest[2] = byte_swap_32 (digest[2]);
digest[3] = byte_swap_32 (digest[3]);
digest[4] = byte_swap_32 (digest[4]);
// precompute salt
char *ptr = (char *) salt->salt_buf_pc;
const int salt_len_pc = snprintf (ptr, 64, "%s$sha1$%u", (char *) salt->salt_buf, iter);
ptr[salt_len_pc] = 0x80;
salt->salt_len_pc = salt_len_pc;
return (PARSER_OK);
}
int atlassian_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
{
if ((input_len < DISPLAY_LEN_MIN_12001) || (input_len > DISPLAY_LEN_MAX_12001)) return (PARSER_GLOBAL_LENGTH);
@ -14874,6 +15081,7 @@ char *strhashtype (const u32 hash_mode)
case 14800: return ((char *) HT_14800);
case 14900: return ((char *) HT_14900);
case 15000: return ((char *) HT_15000);
case 15100: return ((char *) HT_15100);
case 99999: return ((char *) HT_99999);
}
@ -17958,6 +18166,19 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
{
snprintf (out_buf, out_len - 1, "%08x:%08x", digest_buf[0], salt.salt_buf[0]);
}
else if (hash_mode == 15100)
{
// encode the digest:
netbsd_sha1crypt_encode ((unsigned char *) digest_buf, salt.salt_sign[0], (unsigned char *) ptr_plain);
// output:
snprintf (out_buf, out_len - 1, "$sha1$%i$%s$%s",
salt.salt_iter + 1,
(char *) salt.salt_buf,
ptr_plain);
}
else if (hash_mode == 99999)
{
char *ptr = (char *) digest_buf;
@ -22187,6 +22408,20 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
hashconfig->dgst_pos3 = 7;
break;
case 15100: hashconfig->hash_type = HASH_TYPE_SHA1;
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE;
hashconfig->kern_type = KERN_TYPE_NETBSD_SHA1CRYPT;
hashconfig->dgst_size = DGST_SIZE_4_5;
hashconfig->parse_func = netbsd_sha1crypt_parse_hash;
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE;
hashconfig->dgst_pos0 = 0;
hashconfig->dgst_pos1 = 1;
hashconfig->dgst_pos2 = 2;
hashconfig->dgst_pos3 = 3;
break;
case 99999: hashconfig->hash_type = HASH_TYPE_PLAINTEXT;
hashconfig->salt_type = SALT_TYPE_NONE;
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
@ -22423,6 +22658,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
case 14600: hashconfig->tmp_size = sizeof (luks_tmp_t); break;
case 14700: hashconfig->tmp_size = sizeof (pbkdf2_sha1_tmp_t); break;
case 14800: hashconfig->tmp_size = sizeof (pbkdf2_sha256_tmp_t); break;
case 15100: hashconfig->tmp_size = sizeof (pbkdf1_sha1_tmp_t); break;
};
// hook_size
@ -22763,6 +22999,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
break;
case 14900: salt->salt_len = 4;
break;
case 15100: salt->salt_len = 8;
break;
}
// special esalt handling
@ -23015,6 +23253,8 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
case 14800: salt->salt_iter = ROUNDS_ITUNES101_BACKUP - 1;
salt->salt_iter2 = ROUNDS_ITUNES102_BACKUP - 1;
break;
case 15100: salt->salt_iter = ROUNDS_NETBSD_SHA1CRYPT - 1;
break;
}
}

@ -46,7 +46,7 @@ my $hashcat = "./hashcat";
my $MAX_LEN = 55;
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 99999);
my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 125, 130, 131, 132, 133, 140, 141, 150, 160, 200, 300, 400, 500, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 99999);
my %is_unicode = map { $_ => 1 } qw(30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800);
my %less_fifteen = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
@ -2574,6 +2574,29 @@ sub verify
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
}
elsif ($mode == 15100)
{
($hash_in, $word) = split ":", $line;
next unless defined $hash_in;
next unless defined $word;
my @data = split ('\$', $hash_in);
next unless scalar @data == 5;
shift @data;
my $signature = shift @data;
next unless ($signature eq 'sha1');
$iter = shift @data;
$salt = shift @data;
$param = shift @data;
next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
}
else
{
print "ERROR: hash mode is not supported\n";
@ -2907,6 +2930,14 @@ sub verify
return unless (substr ($line, 0, $len) eq $hash_out);
}
elsif ($mode == 15100)
{
$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);
@ -3413,6 +3444,10 @@ sub passthrough
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 8));
}
elsif ($mode == 15100)
{
$tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 8));
}
else
{
print "ERROR: Unsupported hash type\n";
@ -4325,6 +4360,20 @@ sub single
rnd ($mode, 10, 8);
}
}
elsif ($mode == 15100)
{
for (my $i = 1; $i < 32; $i++)
{
if ($len != 0)
{
rnd ($mode, $len, 8);
}
else
{
rnd ($mode, $i, 8);
}
}
}
}
}
@ -7947,6 +7996,45 @@ END_CODE
$tmp_hash = sprintf ("%08x:%s", unpack ("N*", $hash), $salt_buf);
}
elsif ($mode == 15100)
{
my $iterations = 20000;
if (defined ($iter))
{
$iterations = $iter;
}
my $pbkdf1_salt_buf = sprintf ('%s$sha1$%u', $salt_buf, $iterations);
my $tmp = hmac ($pbkdf1_salt_buf, $word_buf, \&sha1, 64);
print unpack ("H*", $tmp), "\n";
for (my $r = 1; $r < $iterations; $r++)
{
$tmp = hmac ($tmp, $word_buf, \&sha1, 64);
}
my $hash_buf = "";
$hash_buf .= to64 ((int (ord (substr ($tmp, 0, 1))) << 16) | (int (ord (substr ($tmp, 1, 1))) << 8) | (int (ord (substr ($tmp, 2, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 3, 1))) << 16) | (int (ord (substr ($tmp, 4, 1))) << 8) | (int (ord (substr ($tmp, 5, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 6, 1))) << 16) | (int (ord (substr ($tmp, 7, 1))) << 8) | (int (ord (substr ($tmp, 8, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 9, 1))) << 16) | (int (ord (substr ($tmp, 10, 1))) << 8) | (int (ord (substr ($tmp, 11, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 12, 1))) << 16) | (int (ord (substr ($tmp, 13, 1))) << 8) | (int (ord (substr ($tmp, 14, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 15, 1))) << 16) | (int (ord (substr ($tmp, 16, 1))) << 8) | (int (ord (substr ($tmp, 17, 1)))), 4);
$hash_buf .= to64 ((int (ord (substr ($tmp, 18, 1))) << 16) | (int (ord (substr ($tmp, 19, 1))) << 8) | 0 , 4);
## super hackish, but we have no other choice, as this byte is kind of a random byte added to the digest when the hash was created
if (defined $additional_param)
{
$hash_buf = substr ($hash_buf, 0, 24) . substr ($additional_param, 24, 4);
}
$tmp_hash = sprintf ("\$sha1\$%d\$%s\$%s", $iterations, $salt_buf, $hash_buf);
}
elsif ($mode == 99999)
{
$tmp_hash = sprintf ("%s", $word_buf);

@ -9,7 +9,7 @@ TDIR="$( cd "$( dirname "${BASH_SOURCE[0]}" )" && pwd )"
# missing hash types: 5200,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 125 130 131 132 133 140 141 150 160 200 300 400 500 900 1000 1100 1300 1400 1410 1411 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 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 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 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 99999"
HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 125 130 131 132 133 140 141 150 160 200 300 400 500 900 1000 1100 1300 1400 1410 1411 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 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 5000 5100 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 7000 7100 7200 7300 7400 7500 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12001 12100 12200 12300 12400 12600 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 99999"
#ATTACK_MODES="0 1 3 6 7"
ATTACK_MODES="0 1 3 7"

Loading…
Cancel
Save