mirror of
https://github.com/hashcat/hashcat
synced 2024-11-20 23:27:31 +01:00
Add HMAC-Streebog-512 (pure kernels)
Implement HMAC based on GOST 34.11-2012 Streebog-512 as well as a test
case for it. Both the PyGOST + hmac python module and the VeraCrypt HMAC
for Streebog-512 were used as references. The kernels expect the digests
to be in big-endian order according to the RFC examples for Streebog.
Fix two bugs from commit 224315dd62.
- Add hash-mode 11850: HMAC-Streebog-512 (key = $pass), big-endian
- Add test case for hash-mode 11850
- Bugfix for a3-pure Streebog kernels (modes 11700 and 11800)
- Rename a few Streebog constants in interface.h
This commit is contained in:
R. Yushaev
parent
3b2c3f419d
commit
a8eb611b1c
@ -991,14 +991,14 @@ DECLSPEC void streebog256_final (streebog256_ctx_t *ctx)
|
||||
|
||||
u64 m[8];
|
||||
|
||||
m[0] = hl32_to_64 (ctx->w3[2], ctx->w3[3]);
|
||||
m[1] = hl32_to_64 (ctx->w3[0], ctx->w3[1]);
|
||||
m[2] = hl32_to_64 (ctx->w2[2], ctx->w2[3]);
|
||||
m[3] = hl32_to_64 (ctx->w2[0], ctx->w2[1]);
|
||||
m[4] = hl32_to_64 (ctx->w1[2], ctx->w1[3]);
|
||||
m[5] = hl32_to_64 (ctx->w1[0], ctx->w1[1]);
|
||||
m[6] = hl32_to_64 (ctx->w0[2], ctx->w0[3]);
|
||||
m[7] = hl32_to_64 (ctx->w0[0], ctx->w0[1]);
|
||||
m[0] = hl32_to_64_S (ctx->w3[2], ctx->w3[3]);
|
||||
m[1] = hl32_to_64_S (ctx->w3[0], ctx->w3[1]);
|
||||
m[2] = hl32_to_64_S (ctx->w2[2], ctx->w2[3]);
|
||||
m[3] = hl32_to_64_S (ctx->w2[0], ctx->w2[1]);
|
||||
m[4] = hl32_to_64_S (ctx->w1[2], ctx->w1[3]);
|
||||
m[5] = hl32_to_64_S (ctx->w1[0], ctx->w1[1]);
|
||||
m[6] = hl32_to_64_S (ctx->w0[2], ctx->w0[3]);
|
||||
m[7] = hl32_to_64_S (ctx->w0[0], ctx->w0[1]);
|
||||
|
||||
streebog256_g (ctx->h, ctx->n, m);
|
||||
|
||||
|
||||
@ -869,6 +869,63 @@ DECLSPEC void streebog512_update_64 (streebog512_ctx_t *ctx, u32 *w0, u32 *w1, u
|
||||
}
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_update (streebog512_ctx_t *ctx, const u32 *w, int len)
|
||||
{
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
int off = 0;
|
||||
|
||||
while (len > 63)
|
||||
{
|
||||
w0[0] = w[off + 0];
|
||||
w0[1] = w[off + 1];
|
||||
w0[2] = w[off + 2];
|
||||
w0[3] = w[off + 3];
|
||||
w1[0] = w[off + 4];
|
||||
w1[1] = w[off + 5];
|
||||
w1[2] = w[off + 6];
|
||||
w1[3] = w[off + 7];
|
||||
w2[0] = w[off + 8];
|
||||
w2[1] = w[off + 9];
|
||||
w2[2] = w[off + 10];
|
||||
w2[3] = w[off + 11];
|
||||
w3[0] = w[off + 12];
|
||||
w3[1] = w[off + 13];
|
||||
w3[2] = w[off + 14];
|
||||
w3[3] = w[off + 15];
|
||||
|
||||
off += 16;
|
||||
len -= 64;
|
||||
|
||||
streebog512_update_64 (ctx, w0, w1, w2, w3, 64);
|
||||
}
|
||||
|
||||
if (len > 0)
|
||||
{
|
||||
w0[0] = w[off + 0];
|
||||
w0[1] = w[off + 1];
|
||||
w0[2] = w[off + 2];
|
||||
w0[3] = w[off + 3];
|
||||
w1[0] = w[off + 4];
|
||||
w1[1] = w[off + 5];
|
||||
w1[2] = w[off + 6];
|
||||
w1[3] = w[off + 7];
|
||||
w2[0] = w[off + 8];
|
||||
w2[1] = w[off + 9];
|
||||
w2[2] = w[off + 10];
|
||||
w2[3] = w[off + 11];
|
||||
w3[0] = w[off + 12];
|
||||
w3[1] = w[off + 13];
|
||||
w3[2] = w[off + 14];
|
||||
w3[3] = w[off + 15];
|
||||
|
||||
streebog512_update_64 (ctx, w0, w1, w2, w3, len);
|
||||
}
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_update_swap (streebog512_ctx_t *ctx, const u32 *w, int len)
|
||||
{
|
||||
u32 w0[4];
|
||||
@ -991,14 +1048,14 @@ DECLSPEC void streebog512_final (streebog512_ctx_t *ctx)
|
||||
|
||||
u64 m[8];
|
||||
|
||||
m[0] = hl32_to_64 (ctx->w3[2], ctx->w3[3]);
|
||||
m[1] = hl32_to_64 (ctx->w3[0], ctx->w3[1]);
|
||||
m[2] = hl32_to_64 (ctx->w2[2], ctx->w2[3]);
|
||||
m[3] = hl32_to_64 (ctx->w2[0], ctx->w2[1]);
|
||||
m[4] = hl32_to_64 (ctx->w1[2], ctx->w1[3]);
|
||||
m[5] = hl32_to_64 (ctx->w1[0], ctx->w1[1]);
|
||||
m[6] = hl32_to_64 (ctx->w0[2], ctx->w0[3]);
|
||||
m[7] = hl32_to_64 (ctx->w0[0], ctx->w0[1]);
|
||||
m[0] = hl32_to_64_S (ctx->w3[2], ctx->w3[3]);
|
||||
m[1] = hl32_to_64_S (ctx->w3[0], ctx->w3[1]);
|
||||
m[2] = hl32_to_64_S (ctx->w2[2], ctx->w2[3]);
|
||||
m[3] = hl32_to_64_S (ctx->w2[0], ctx->w2[1]);
|
||||
m[4] = hl32_to_64_S (ctx->w1[2], ctx->w1[3]);
|
||||
m[5] = hl32_to_64_S (ctx->w1[0], ctx->w1[1]);
|
||||
m[6] = hl32_to_64_S (ctx->w0[2], ctx->w0[3]);
|
||||
m[7] = hl32_to_64_S (ctx->w0[0], ctx->w0[1]);
|
||||
|
||||
streebog512_g (ctx->h, ctx->n, m);
|
||||
|
||||
@ -1016,6 +1073,217 @@ DECLSPEC void streebog512_final (streebog512_ctx_t *ctx)
|
||||
streebog512_g (ctx->h, nullbuf, ctx->s);
|
||||
}
|
||||
|
||||
typedef struct streebog512_hmac_ctx
|
||||
{
|
||||
streebog512_ctx_t ipad;
|
||||
streebog512_ctx_t opad;
|
||||
|
||||
} streebog512_hmac_ctx_t;
|
||||
|
||||
DECLSPEC void streebog512_hmac_init_64 (streebog512_hmac_ctx_t *ctx, const u32 *w0, const u32 *w1, const u32 *w2, const u32 *w3)
|
||||
{
|
||||
u32 t0[4];
|
||||
u32 t1[4];
|
||||
u32 t2[4];
|
||||
u32 t3[4];
|
||||
|
||||
// ipad
|
||||
|
||||
t0[0] = w0[0] ^ 0x36363636;
|
||||
t0[1] = w0[1] ^ 0x36363636;
|
||||
t0[2] = w0[2] ^ 0x36363636;
|
||||
t0[3] = w0[3] ^ 0x36363636;
|
||||
t1[0] = w1[0] ^ 0x36363636;
|
||||
t1[1] = w1[1] ^ 0x36363636;
|
||||
t1[2] = w1[2] ^ 0x36363636;
|
||||
t1[3] = w1[3] ^ 0x36363636;
|
||||
t2[0] = w2[0] ^ 0x36363636;
|
||||
t2[1] = w2[1] ^ 0x36363636;
|
||||
t2[2] = w2[2] ^ 0x36363636;
|
||||
t2[3] = w2[3] ^ 0x36363636;
|
||||
t3[0] = w3[0] ^ 0x36363636;
|
||||
t3[1] = w3[1] ^ 0x36363636;
|
||||
t3[2] = w3[2] ^ 0x36363636;
|
||||
t3[3] = w3[3] ^ 0x36363636;
|
||||
|
||||
streebog512_init (&ctx->ipad);
|
||||
|
||||
streebog512_update_64 (&ctx->ipad, t0, t1, t2, t3, 64);
|
||||
|
||||
// opad
|
||||
|
||||
t0[0] = w0[0] ^ 0x5c5c5c5c;
|
||||
t0[1] = w0[1] ^ 0x5c5c5c5c;
|
||||
t0[2] = w0[2] ^ 0x5c5c5c5c;
|
||||
t0[3] = w0[3] ^ 0x5c5c5c5c;
|
||||
t1[0] = w1[0] ^ 0x5c5c5c5c;
|
||||
t1[1] = w1[1] ^ 0x5c5c5c5c;
|
||||
t1[2] = w1[2] ^ 0x5c5c5c5c;
|
||||
t1[3] = w1[3] ^ 0x5c5c5c5c;
|
||||
t2[0] = w2[0] ^ 0x5c5c5c5c;
|
||||
t2[1] = w2[1] ^ 0x5c5c5c5c;
|
||||
t2[2] = w2[2] ^ 0x5c5c5c5c;
|
||||
t2[3] = w2[3] ^ 0x5c5c5c5c;
|
||||
t3[0] = w3[0] ^ 0x5c5c5c5c;
|
||||
t3[1] = w3[1] ^ 0x5c5c5c5c;
|
||||
t3[2] = w3[2] ^ 0x5c5c5c5c;
|
||||
t3[3] = w3[3] ^ 0x5c5c5c5c;
|
||||
|
||||
streebog512_init (&ctx->opad);
|
||||
|
||||
streebog512_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_init (streebog512_hmac_ctx_t *ctx, const u32 *w, const int len)
|
||||
{
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
if (len > 64)
|
||||
{
|
||||
streebog512_ctx_t tmp;
|
||||
|
||||
streebog512_init (&tmp);
|
||||
|
||||
streebog512_update (&tmp, w, len);
|
||||
|
||||
streebog512_final (&tmp);
|
||||
|
||||
w0[0] = h32_from_64_S (tmp.h[0]);
|
||||
w0[1] = l32_from_64_S (tmp.h[0]);
|
||||
w0[2] = h32_from_64_S (tmp.h[1]);
|
||||
w0[3] = l32_from_64_S (tmp.h[1]);
|
||||
w1[0] = h32_from_64_S (tmp.h[2]);
|
||||
w1[1] = l32_from_64_S (tmp.h[2]);
|
||||
w1[2] = h32_from_64_S (tmp.h[3]);
|
||||
w1[3] = l32_from_64_S (tmp.h[3]);
|
||||
w2[0] = h32_from_64_S (tmp.h[4]);
|
||||
w2[1] = l32_from_64_S (tmp.h[4]);
|
||||
w2[2] = h32_from_64_S (tmp.h[5]);
|
||||
w2[3] = l32_from_64_S (tmp.h[5]);
|
||||
w3[0] = h32_from_64_S (tmp.h[6]);
|
||||
w3[1] = l32_from_64_S (tmp.h[6]);
|
||||
w3[2] = h32_from_64_S (tmp.h[7]);
|
||||
w3[3] = l32_from_64_S (tmp.h[7]);
|
||||
}
|
||||
else
|
||||
{
|
||||
w0[0] = w[ 0];
|
||||
w0[1] = w[ 1];
|
||||
w0[2] = w[ 2];
|
||||
w0[3] = w[ 3];
|
||||
w1[0] = w[ 4];
|
||||
w1[1] = w[ 5];
|
||||
w1[2] = w[ 6];
|
||||
w1[3] = w[ 7];
|
||||
w2[0] = w[ 8];
|
||||
w2[1] = w[ 9];
|
||||
w2[2] = w[10];
|
||||
w2[3] = w[11];
|
||||
w3[0] = w[12];
|
||||
w3[1] = w[13];
|
||||
w3[2] = w[14];
|
||||
w3[3] = w[15];
|
||||
}
|
||||
|
||||
streebog512_hmac_init_64 (ctx, w0, w1, w2, w3);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_init_swap (streebog512_hmac_ctx_t *ctx, const u32 *w, const int len)
|
||||
{
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
u32 w2[4];
|
||||
u32 w3[4];
|
||||
|
||||
if (len > 64)
|
||||
{
|
||||
streebog512_ctx_t tmp;
|
||||
|
||||
streebog512_init (&tmp);
|
||||
|
||||
streebog512_update_swap (&tmp, w, len);
|
||||
|
||||
streebog512_final (&tmp);
|
||||
|
||||
w0[0] = h32_from_64_S (tmp.h[0]);
|
||||
w0[1] = l32_from_64_S (tmp.h[0]);
|
||||
w0[2] = h32_from_64_S (tmp.h[1]);
|
||||
w0[3] = l32_from_64_S (tmp.h[1]);
|
||||
w1[0] = h32_from_64_S (tmp.h[2]);
|
||||
w1[1] = l32_from_64_S (tmp.h[2]);
|
||||
w1[2] = h32_from_64_S (tmp.h[3]);
|
||||
w1[3] = l32_from_64_S (tmp.h[3]);
|
||||
w2[0] = h32_from_64_S (tmp.h[4]);
|
||||
w2[1] = l32_from_64_S (tmp.h[4]);
|
||||
w2[2] = h32_from_64_S (tmp.h[5]);
|
||||
w2[3] = l32_from_64_S (tmp.h[5]);
|
||||
w3[0] = h32_from_64_S (tmp.h[6]);
|
||||
w3[1] = l32_from_64_S (tmp.h[6]);
|
||||
w3[2] = h32_from_64_S (tmp.h[7]);
|
||||
w3[3] = l32_from_64_S (tmp.h[7]);
|
||||
}
|
||||
else
|
||||
{
|
||||
w0[0] = swap32_S (w[ 0]);
|
||||
w0[1] = swap32_S (w[ 1]);
|
||||
w0[2] = swap32_S (w[ 2]);
|
||||
w0[3] = swap32_S (w[ 3]);
|
||||
w1[0] = swap32_S (w[ 4]);
|
||||
w1[1] = swap32_S (w[ 5]);
|
||||
w1[2] = swap32_S (w[ 6]);
|
||||
w1[3] = swap32_S (w[ 7]);
|
||||
w2[0] = swap32_S (w[ 8]);
|
||||
w2[1] = swap32_S (w[ 9]);
|
||||
w2[2] = swap32_S (w[10]);
|
||||
w2[3] = swap32_S (w[11]);
|
||||
w3[0] = swap32_S (w[12]);
|
||||
w3[1] = swap32_S (w[13]);
|
||||
w3[2] = swap32_S (w[14]);
|
||||
w3[3] = swap32_S (w[15]);
|
||||
}
|
||||
|
||||
streebog512_hmac_init_64 (ctx, w0, w1, w2, w3);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_update (streebog512_hmac_ctx_t *ctx, const u32 *w, const int len)
|
||||
{
|
||||
streebog512_update (&ctx->ipad, w, len);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_final (streebog512_hmac_ctx_t *ctx)
|
||||
{
|
||||
streebog512_final (&ctx->ipad);
|
||||
|
||||
u32 t0[4];
|
||||
u32 t1[4];
|
||||
u32 t2[4];
|
||||
u32 t3[4];
|
||||
|
||||
t0[0] = h32_from_64_S (ctx->ipad.h[7]);
|
||||
t0[1] = l32_from_64_S (ctx->ipad.h[7]);
|
||||
t0[2] = h32_from_64_S (ctx->ipad.h[6]);
|
||||
t0[3] = l32_from_64_S (ctx->ipad.h[6]);
|
||||
t1[0] = h32_from_64_S (ctx->ipad.h[5]);
|
||||
t1[1] = l32_from_64_S (ctx->ipad.h[5]);
|
||||
t1[2] = h32_from_64_S (ctx->ipad.h[4]);
|
||||
t1[3] = l32_from_64_S (ctx->ipad.h[4]);
|
||||
t2[0] = h32_from_64_S (ctx->ipad.h[3]);
|
||||
t2[1] = l32_from_64_S (ctx->ipad.h[3]);
|
||||
t2[2] = h32_from_64_S (ctx->ipad.h[2]);
|
||||
t2[3] = l32_from_64_S (ctx->ipad.h[2]);
|
||||
t3[0] = h32_from_64_S (ctx->ipad.h[1]);
|
||||
t3[1] = l32_from_64_S (ctx->ipad.h[1]);
|
||||
t3[2] = h32_from_64_S (ctx->ipad.h[0]);
|
||||
t3[3] = l32_from_64_S (ctx->ipad.h[0]);
|
||||
|
||||
streebog512_update_64 (&ctx->opad, t0, t1, t2, t3, 64);
|
||||
|
||||
streebog512_final (&ctx->opad);
|
||||
}
|
||||
|
||||
typedef struct streebog512_ctx_vector
|
||||
{
|
||||
u64x h[8];
|
||||
@ -1270,6 +1538,63 @@ DECLSPEC void streebog512_update_vector_64 (streebog512_ctx_vector_t *ctx, u32x
|
||||
}
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_update_vector (streebog512_ctx_vector_t *ctx, const u32x *w, int len)
|
||||
{
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
u32x w2[4];
|
||||
u32x w3[4];
|
||||
|
||||
int off = 0;
|
||||
|
||||
while (len > 63)
|
||||
{
|
||||
w0[0] = w[off + 0];
|
||||
w0[1] = w[off + 1];
|
||||
w0[2] = w[off + 2];
|
||||
w0[3] = w[off + 3];
|
||||
w1[0] = w[off + 4];
|
||||
w1[1] = w[off + 5];
|
||||
w1[2] = w[off + 6];
|
||||
w1[3] = w[off + 7];
|
||||
w2[0] = w[off + 8];
|
||||
w2[1] = w[off + 9];
|
||||
w2[2] = w[off + 10];
|
||||
w2[3] = w[off + 11];
|
||||
w3[0] = w[off + 12];
|
||||
w3[1] = w[off + 13];
|
||||
w3[2] = w[off + 14];
|
||||
w3[3] = w[off + 15];
|
||||
|
||||
off += 16;
|
||||
len -= 64;
|
||||
|
||||
streebog512_update_vector_64 (ctx, w0, w1, w2, w3, 64);
|
||||
}
|
||||
|
||||
if (len > 0)
|
||||
{
|
||||
w0[0] = w[off + 0];
|
||||
w0[1] = w[off + 1];
|
||||
w0[2] = w[off + 2];
|
||||
w0[3] = w[off + 3];
|
||||
w1[0] = w[off + 4];
|
||||
w1[1] = w[off + 5];
|
||||
w1[2] = w[off + 6];
|
||||
w1[3] = w[off + 7];
|
||||
w2[0] = w[off + 8];
|
||||
w2[1] = w[off + 9];
|
||||
w2[2] = w[off + 10];
|
||||
w2[3] = w[off + 11];
|
||||
w3[0] = w[off + 12];
|
||||
w3[1] = w[off + 13];
|
||||
w3[2] = w[off + 14];
|
||||
w3[3] = w[off + 15];
|
||||
|
||||
streebog512_update_vector_64 (ctx, w0, w1, w2, w3, len);
|
||||
}
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_update_vector_swap (streebog512_ctx_vector_t *ctx, const u32x *w, int len)
|
||||
{
|
||||
u32x w0[4];
|
||||
@ -1359,3 +1684,157 @@ DECLSPEC void streebog512_final_vector (streebog512_ctx_vector_t *ctx)
|
||||
|
||||
streebog512_g_vector (ctx->h, nullbuf, ctx->s);
|
||||
}
|
||||
|
||||
typedef struct streebog512_hmac_ctx_vector
|
||||
{
|
||||
streebog512_ctx_vector_t ipad;
|
||||
streebog512_ctx_vector_t opad;
|
||||
|
||||
} streebog512_hmac_ctx_vector_t;
|
||||
|
||||
DECLSPEC void streebog512_hmac_init_vector_64 (streebog512_hmac_ctx_vector_t *ctx, const u32x *w0, const u32x *w1, const u32x *w2, const u32x *w3)
|
||||
{
|
||||
u32x t0[4];
|
||||
u32x t1[4];
|
||||
u32x t2[4];
|
||||
u32x t3[4];
|
||||
|
||||
// ipad
|
||||
|
||||
t0[0] = w0[0] ^ 0x36363636;
|
||||
t0[1] = w0[1] ^ 0x36363636;
|
||||
t0[2] = w0[2] ^ 0x36363636;
|
||||
t0[3] = w0[3] ^ 0x36363636;
|
||||
t1[0] = w1[0] ^ 0x36363636;
|
||||
t1[1] = w1[1] ^ 0x36363636;
|
||||
t1[2] = w1[2] ^ 0x36363636;
|
||||
t1[3] = w1[3] ^ 0x36363636;
|
||||
t2[0] = w2[0] ^ 0x36363636;
|
||||
t2[1] = w2[1] ^ 0x36363636;
|
||||
t2[2] = w2[2] ^ 0x36363636;
|
||||
t2[3] = w2[3] ^ 0x36363636;
|
||||
t3[0] = w3[0] ^ 0x36363636;
|
||||
t3[1] = w3[1] ^ 0x36363636;
|
||||
t3[2] = w3[2] ^ 0x36363636;
|
||||
t3[3] = w3[3] ^ 0x36363636;
|
||||
|
||||
streebog512_init_vector (&ctx->ipad);
|
||||
|
||||
streebog512_update_vector_64 (&ctx->ipad, t0, t1, t2, t3, 64);
|
||||
|
||||
// opad
|
||||
|
||||
t0[0] = w0[0] ^ 0x5c5c5c5c;
|
||||
t0[1] = w0[1] ^ 0x5c5c5c5c;
|
||||
t0[2] = w0[2] ^ 0x5c5c5c5c;
|
||||
t0[3] = w0[3] ^ 0x5c5c5c5c;
|
||||
t1[0] = w1[0] ^ 0x5c5c5c5c;
|
||||
t1[1] = w1[1] ^ 0x5c5c5c5c;
|
||||
t1[2] = w1[2] ^ 0x5c5c5c5c;
|
||||
t1[3] = w1[3] ^ 0x5c5c5c5c;
|
||||
t2[0] = w2[0] ^ 0x5c5c5c5c;
|
||||
t2[1] = w2[1] ^ 0x5c5c5c5c;
|
||||
t2[2] = w2[2] ^ 0x5c5c5c5c;
|
||||
t2[3] = w2[3] ^ 0x5c5c5c5c;
|
||||
t3[0] = w3[0] ^ 0x5c5c5c5c;
|
||||
t3[1] = w3[1] ^ 0x5c5c5c5c;
|
||||
t3[2] = w3[2] ^ 0x5c5c5c5c;
|
||||
t3[3] = w3[3] ^ 0x5c5c5c5c;
|
||||
|
||||
streebog512_init_vector (&ctx->opad);
|
||||
|
||||
streebog512_update_vector_64 (&ctx->opad, t0, t1, t2, t3, 64);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_init_vector_swap (streebog512_hmac_ctx_vector_t *ctx, const u32x *w, const int len)
|
||||
{
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
u32x w2[4];
|
||||
u32x w3[4];
|
||||
|
||||
if (len > 64)
|
||||
{
|
||||
streebog512_ctx_vector_t tmp;
|
||||
|
||||
streebog512_init_vector (&tmp);
|
||||
|
||||
streebog512_update_vector (&tmp, w, len);
|
||||
|
||||
streebog512_final_vector (&tmp);
|
||||
|
||||
w0[0] = h32_from_64 (tmp.h[0]);
|
||||
w0[1] = l32_from_64 (tmp.h[0]);
|
||||
w0[2] = h32_from_64 (tmp.h[1]);
|
||||
w0[3] = l32_from_64 (tmp.h[1]);
|
||||
w1[0] = h32_from_64 (tmp.h[2]);
|
||||
w1[1] = l32_from_64 (tmp.h[2]);
|
||||
w1[2] = h32_from_64 (tmp.h[3]);
|
||||
w1[3] = l32_from_64 (tmp.h[3]);
|
||||
w2[0] = h32_from_64 (tmp.h[4]);
|
||||
w2[1] = l32_from_64 (tmp.h[4]);
|
||||
w2[2] = h32_from_64 (tmp.h[5]);
|
||||
w2[3] = l32_from_64 (tmp.h[5]);
|
||||
w3[0] = h32_from_64 (tmp.h[6]);
|
||||
w3[1] = l32_from_64 (tmp.h[6]);
|
||||
w3[2] = h32_from_64 (tmp.h[7]);
|
||||
w3[3] = l32_from_64 (tmp.h[7]);
|
||||
}
|
||||
else
|
||||
{
|
||||
w0[0] = swap32 (w[ 0]);
|
||||
w0[1] = swap32 (w[ 1]);
|
||||
w0[2] = swap32 (w[ 2]);
|
||||
w0[3] = swap32 (w[ 3]);
|
||||
w1[0] = swap32 (w[ 4]);
|
||||
w1[1] = swap32 (w[ 5]);
|
||||
w1[2] = swap32 (w[ 6]);
|
||||
w1[3] = swap32 (w[ 7]);
|
||||
w2[0] = swap32 (w[ 8]);
|
||||
w2[1] = swap32 (w[ 9]);
|
||||
w2[2] = swap32 (w[10]);
|
||||
w2[3] = swap32 (w[11]);
|
||||
w3[0] = swap32 (w[12]);
|
||||
w3[1] = swap32 (w[13]);
|
||||
w3[2] = swap32 (w[14]);
|
||||
w3[3] = swap32 (w[15]);
|
||||
}
|
||||
|
||||
streebog512_hmac_init_vector_64 (ctx, w0, w1, w2, w3);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_update_vector (streebog512_hmac_ctx_vector_t *ctx, const u32x *w, const int len)
|
||||
{
|
||||
streebog512_update_vector (&ctx->ipad, w, len);
|
||||
}
|
||||
|
||||
DECLSPEC void streebog512_hmac_final_vector (streebog512_hmac_ctx_vector_t *ctx)
|
||||
{
|
||||
streebog512_final_vector (&ctx->ipad);
|
||||
|
||||
u32x t0[4];
|
||||
u32x t1[4];
|
||||
u32x t2[4];
|
||||
u32x t3[4];
|
||||
|
||||
t0[0] = h32_from_64 (ctx->ipad.h[7]);
|
||||
t0[1] = l32_from_64 (ctx->ipad.h[7]);
|
||||
t0[2] = h32_from_64 (ctx->ipad.h[6]);
|
||||
t0[3] = l32_from_64 (ctx->ipad.h[6]);
|
||||
t1[0] = h32_from_64 (ctx->ipad.h[5]);
|
||||
t1[1] = l32_from_64 (ctx->ipad.h[5]);
|
||||
t1[2] = h32_from_64 (ctx->ipad.h[4]);
|
||||
t1[3] = l32_from_64 (ctx->ipad.h[4]);
|
||||
t2[0] = h32_from_64 (ctx->ipad.h[3]);
|
||||
t2[1] = l32_from_64 (ctx->ipad.h[3]);
|
||||
t2[2] = h32_from_64 (ctx->ipad.h[2]);
|
||||
t2[3] = l32_from_64 (ctx->ipad.h[2]);
|
||||
t3[0] = h32_from_64 (ctx->ipad.h[1]);
|
||||
t3[1] = l32_from_64 (ctx->ipad.h[1]);
|
||||
t3[2] = h32_from_64 (ctx->ipad.h[0]);
|
||||
t3[3] = l32_from_64 (ctx->ipad.h[0]);
|
||||
|
||||
streebog512_update_vector_64 (&ctx->opad, t0, t1, t2, t3, 64);
|
||||
|
||||
streebog512_final_vector (&ctx->opad);
|
||||
}
|
||||
|
||||
@ -55,7 +55,7 @@ __kernel void m11700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
streebog256_init_vector (&ctx);
|
||||
|
||||
streebog256_update_vector (&ctx, w, pw_len);
|
||||
streebog256_update_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog256_final_vector (&ctx);
|
||||
|
||||
@ -122,7 +122,7 @@ __kernel void m11700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
streebog256_init_vector (&ctx);
|
||||
|
||||
streebog256_update_vector (&ctx, w, pw_len);
|
||||
streebog256_update_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog256_final_vector (&ctx);
|
||||
|
||||
|
||||
@ -55,7 +55,7 @@ __kernel void m11800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
streebog512_init_vector (&ctx);
|
||||
|
||||
streebog512_update_vector (&ctx, w, pw_len);
|
||||
streebog512_update_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog512_final_vector (&ctx);
|
||||
|
||||
@ -122,7 +122,7 @@ __kernel void m11800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
streebog512_init_vector (&ctx);
|
||||
|
||||
streebog512_update_vector (&ctx, w, pw_len);
|
||||
streebog512_update_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog512_final_vector (&ctx);
|
||||
|
||||
|
||||
134
OpenCL/m11850_a0-pure.cl
Normal file
134
OpenCL/m11850_a0-pure.cl
Normal file
@ -0,0 +1,134 @@
|
||||
/**
|
||||
* 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_rp.h"
|
||||
#include "inc_rp.cl"
|
||||
#include "inc_scalar.cl"
|
||||
#include "inc_hash_streebog512.cl"
|
||||
|
||||
__kernel void m11850_mxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
COPY_PW (pws[gid]);
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
pw_t tmp = PASTE_PW;
|
||||
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
streebog512_hmac_ctx_t ctx;
|
||||
|
||||
streebog512_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
streebog512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r2 = l32_from_64_S (ctx.opad.h[1]);
|
||||
const u32 r3 = h32_from_64_S (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m11850_sxx (__global pw_t *pws, __constant const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
||||
};
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
COPY_PW (pws[gid]);
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
pw_t tmp = PASTE_PW;
|
||||
|
||||
tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
|
||||
|
||||
streebog512_hmac_ctx_t ctx;
|
||||
|
||||
streebog512_hmac_init_swap (&ctx, tmp.i, tmp.pw_len);
|
||||
|
||||
streebog512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r2 = l32_from_64_S (ctx.opad.h[1]);
|
||||
const u32 r3 = h32_from_64_S (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
182
OpenCL/m11850_a1-pure.cl
Normal file
182
OpenCL/m11850_a1-pure.cl
Normal file
@ -0,0 +1,182 @@
|
||||
/**
|
||||
* 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_scalar.cl"
|
||||
#include "inc_hash_streebog512.cl"
|
||||
|
||||
__kernel void m11850_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
const u32 comb_len = combs_buf[il_pos].pw_len;
|
||||
|
||||
u32 c[64];
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int idx = 0; idx < 64; idx++)
|
||||
{
|
||||
c[idx] = swap32_S (combs_buf[il_pos].i[idx]);
|
||||
}
|
||||
|
||||
switch_buffer_by_offset_1x64_be_S (c, pw_len);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 64; i++)
|
||||
{
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
streebog512_hmac_ctx_t ctx;
|
||||
|
||||
streebog512_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
streebog512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r2 = l32_from_64_S (ctx.opad.h[1]);
|
||||
const u32 r3 = h32_from_64_S (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m11850_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
||||
};
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
||||
{
|
||||
const u32 comb_len = combs_buf[il_pos].pw_len;
|
||||
|
||||
u32 c[64];
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int idx = 0; idx < 64; idx++)
|
||||
{
|
||||
c[idx] = swap32_S (combs_buf[il_pos].i[idx]);
|
||||
}
|
||||
|
||||
switch_buffer_by_offset_1x64_be_S (c, pw_len);
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (int i = 0; i < 64; i++)
|
||||
{
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
streebog512_hmac_ctx_t ctx;
|
||||
|
||||
streebog512_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
streebog512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[0]);
|
||||
const u32 r2 = l32_from_64_S (ctx.opad.h[1]);
|
||||
const u32 r3 = h32_from_64_S (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
154
OpenCL/m11850_a3-pure.cl
Normal file
154
OpenCL/m11850_a3-pure.cl
Normal file
@ -0,0 +1,154 @@
|
||||
/**
|
||||
* Author......: See docs/credits.txt
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
#define NEW_SIMD_CODE
|
||||
|
||||
#include "inc_vendor.cl"
|
||||
#include "inc_hash_constants.h"
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_simd.cl"
|
||||
#include "inc_hash_streebog512.cl"
|
||||
|
||||
__kernel void m11850_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
u32x w[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32x w0l = w[0];
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
|
||||
const u32x w0 = w0l | w0r;
|
||||
|
||||
w[0] = w0;
|
||||
|
||||
streebog512_hmac_ctx_vector_t ctx;
|
||||
|
||||
streebog512_hmac_init_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog512_hmac_update_vector (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final_vector (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.opad.h[0]);
|
||||
const u32x r1 = h32_from_64 (ctx.opad.h[0]);
|
||||
const u32x r2 = l32_from_64 (ctx.opad.h[1]);
|
||||
const u32x r3 = h32_from_64 (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_M_SIMD (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m11850_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u64 lid = get_local_id (0);
|
||||
const u64 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
/**
|
||||
* digest
|
||||
*/
|
||||
|
||||
const u32 search[4] =
|
||||
{
|
||||
digests_buf[digests_offset].digest_buf[DGST_R0],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R1],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R2],
|
||||
digests_buf[digests_offset].digest_buf[DGST_R3]
|
||||
};
|
||||
|
||||
/**
|
||||
* base
|
||||
*/
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
u32x w[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
|
||||
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
}
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32x w0l = w[0];
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
|
||||
const u32x w0 = w0l | w0r;
|
||||
|
||||
w[0] = w0;
|
||||
|
||||
streebog512_hmac_ctx_vector_t ctx;
|
||||
|
||||
streebog512_hmac_init_vector_swap (&ctx, w, pw_len);
|
||||
|
||||
streebog512_hmac_update_vector (&ctx, s, salt_len);
|
||||
|
||||
streebog512_hmac_final_vector (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.opad.h[0]);
|
||||
const u32x r1 = h32_from_64 (ctx.opad.h[0]);
|
||||
const u32x r2 = l32_from_64 (ctx.opad.h[1]);
|
||||
const u32x r3 = h32_from_64 (ctx.opad.h[1]);
|
||||
|
||||
COMPARE_S_SIMD (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
@ -13,12 +13,14 @@
|
||||
- Added pure kernels for hash-mode 11700 (Streebog-256)
|
||||
- Added pure kernels for hash-mode 11800 (Streebog-512)
|
||||
- Added hash-modes 18200 (Kerberos 5 AS-REP etype 23)
|
||||
- Added hash-mode 11850 (HMAC-Streebog-512 (key = $pass), big-endian)
|
||||
|
||||
##
|
||||
## Improvements
|
||||
##
|
||||
|
||||
- Tests: Added hash-modes 11700 (Streebog-256) and 11800 (Streebog-512)
|
||||
- Tests: Added hash-mode 11850 (HMAC-Streebog-512 (key = $pass), big-endian)
|
||||
|
||||
##
|
||||
## Bugs
|
||||
|
||||
@ -176,7 +176,7 @@ _hashcat ()
|
||||
{
|
||||
local VERSION=5.0.0
|
||||
|
||||
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 600 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 2501 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 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 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16801 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100 18200"
|
||||
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 600 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 2501 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 3910 4010 4110 4300 4400 4500 4520 4521 4522 4700 4800 4900 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 11850 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 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16801 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100 18200"
|
||||
local ATTACK_MODES="0 1 3 6 7"
|
||||
local HCCAPX_MESSAGE_PAIRS="0 1 2 3 4 5"
|
||||
local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15"
|
||||
|
||||
@ -1093,8 +1093,8 @@ typedef enum hash_type
|
||||
HASH_TYPE_PBKDF2_SHA256 = 39,
|
||||
HASH_TYPE_BITCOIN_WALLET = 40,
|
||||
HASH_TYPE_CRC32 = 41,
|
||||
HASH_TYPE_GOST_2012SBOG_256 = 42,
|
||||
HASH_TYPE_GOST_2012SBOG_512 = 43,
|
||||
HASH_TYPE_STREEBOG_256 = 42,
|
||||
HASH_TYPE_STREEBOG_512 = 43,
|
||||
HASH_TYPE_PBKDF2_MD5 = 44,
|
||||
HASH_TYPE_PBKDF2_SHA1 = 45,
|
||||
HASH_TYPE_PBKDF2_SHA512 = 46,
|
||||
@ -1270,8 +1270,9 @@ typedef enum kern_type
|
||||
KERN_TYPE_SIP_AUTH = 11400,
|
||||
KERN_TYPE_CRC32 = 11500,
|
||||
KERN_TYPE_SEVEN_ZIP = 11600,
|
||||
KERN_TYPE_GOST_2012SBOG_256 = 11700,
|
||||
KERN_TYPE_GOST_2012SBOG_512 = 11800,
|
||||
KERN_TYPE_STREEBOG_256 = 11700,
|
||||
KERN_TYPE_STREEBOG_512 = 11800,
|
||||
KERN_TYPE_HMAC_STREEBOG_512_PW = 11850,
|
||||
KERN_TYPE_PBKDF2_MD5 = 11900,
|
||||
KERN_TYPE_PBKDF2_SHA1 = 12000,
|
||||
KERN_TYPE_ECRYPTFS = 12200,
|
||||
@ -1559,8 +1560,8 @@ int bitcoin_wallet_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu
|
||||
int sip_auth_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int crc32_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int seven_zip_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int gost2012sbog_256_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int gost2012sbog_512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int streebog_256_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int streebog_512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int pbkdf2_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int pbkdf2_sha1_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
int pbkdf2_sha512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig);
|
||||
|
||||
@ -223,6 +223,7 @@ static const char *ST_HASH_11500 = "c762de4a:00000000";
|
||||
static const char *ST_HASH_11600 = "$7z$0$14$0$$11$33363437353138333138300000000000$2365089182$16$12$d00321533b483f54a523f624a5f63269";
|
||||
static const char *ST_HASH_11700 = "57e9e50caec93d72e9498c211d6dc4f4d328248b48ecf46ba7abfa874f666e36";
|
||||
static const char *ST_HASH_11800 = "5d5bdba48c8f89ee6c0a0e11023540424283e84902de08013aeeb626e819950bb32842903593a1d2e8f71897ff7fe72e17ac9ba8ce1d1d2f7e9c4359ea63bdc3";
|
||||
static const char *ST_HASH_11850 = "ddf70f583b785137a5aa272549278d5f8830b8344e16ce8058e1605bc49c83063eb3b5d65af9223bfbef097de1f4ff7bd80609c2277f8b033ad72dab89e9d1f9:11112222";
|
||||
static const char *ST_HASH_11900 = "md5:1000:NjAxMDY4MQ==:a00DtIW9hP9voC85fmEA5uVhgdDx67nSPSm9yADHjkI=";
|
||||
static const char *ST_HASH_12000 = "sha1:1000:MTYwNTM4MDU4Mzc4MzA=:aGghFQBtQ8+WVlMk5GEaMw==";
|
||||
static const char *ST_HASH_12001 = "{PKCS5S2}NTczNTY0NDY2NjQyNzU1Mx8gGiRGobaZYwumctGHbn2ZOHB8LkwzH+Z1gkWfy1zD";
|
||||
@ -491,6 +492,7 @@ static const char *HT_11500 = "CRC32";
|
||||
static const char *HT_11600 = "7-Zip";
|
||||
static const char *HT_11700 = "GOST R 34.11-2012 (Streebog) 256-bit, big-endian";
|
||||
static const char *HT_11800 = "GOST R 34.11-2012 (Streebog) 512-bit, big-endian";
|
||||
static const char *HT_11850 = "HMAC-Streebog-512 (key = $pass), big-endian";
|
||||
static const char *HT_11900 = "PBKDF2-HMAC-MD5";
|
||||
static const char *HT_12000 = "PBKDF2-HMAC-SHA1";
|
||||
static const char *HT_12100 = "PBKDF2-HMAC-SHA512";
|
||||
@ -13368,7 +13370,7 @@ int seven_zip_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int gost2012sbog_256_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
int streebog_256_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
{
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
@ -13399,7 +13401,7 @@ int gost2012sbog_256_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf,
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int gost2012sbog_512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
int streebog_512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
{
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
@ -13438,6 +13440,67 @@ int gost2012sbog_512_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf,
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int streebog_512s_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
{
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
salt_t *salt = hash_buf->salt;
|
||||
|
||||
token_t token;
|
||||
|
||||
token.token_cnt = 2;
|
||||
|
||||
token.sep[0] = hashconfig->separator;
|
||||
token.len_min[0] = 128;
|
||||
token.len_max[0] = 128;
|
||||
token.attr[0] = TOKEN_ATTR_VERIFY_LENGTH
|
||||
| TOKEN_ATTR_VERIFY_HEX;
|
||||
|
||||
token.len_min[1] = SALT_MIN;
|
||||
token.len_max[1] = SALT_MAX;
|
||||
token.attr[1] = TOKEN_ATTR_VERIFY_LENGTH;
|
||||
|
||||
if (hashconfig->opts_type & OPTS_TYPE_ST_HEX)
|
||||
{
|
||||
token.len_min[1] *= 2;
|
||||
token.len_max[1] *= 2;
|
||||
|
||||
token.attr[1] |= TOKEN_ATTR_VERIFY_HEX;
|
||||
}
|
||||
|
||||
const int rc_tokenizer = input_tokenizer (input_buf, input_len, &token);
|
||||
|
||||
if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
|
||||
|
||||
u8 *hash_pos = token.buf[0];
|
||||
|
||||
digest[ 0] = hex_to_u32 (hash_pos + 0);
|
||||
digest[ 1] = hex_to_u32 (hash_pos + 8);
|
||||
digest[ 2] = hex_to_u32 (hash_pos + 16);
|
||||
digest[ 3] = hex_to_u32 (hash_pos + 24);
|
||||
digest[ 4] = hex_to_u32 (hash_pos + 32);
|
||||
digest[ 5] = hex_to_u32 (hash_pos + 40);
|
||||
digest[ 6] = hex_to_u32 (hash_pos + 48);
|
||||
digest[ 7] = hex_to_u32 (hash_pos + 56);
|
||||
digest[ 8] = hex_to_u32 (hash_pos + 64);
|
||||
digest[ 9] = hex_to_u32 (hash_pos + 72);
|
||||
digest[10] = hex_to_u32 (hash_pos + 80);
|
||||
digest[11] = hex_to_u32 (hash_pos + 88);
|
||||
digest[12] = hex_to_u32 (hash_pos + 96);
|
||||
digest[13] = hex_to_u32 (hash_pos + 104);
|
||||
digest[14] = hex_to_u32 (hash_pos + 112);
|
||||
digest[15] = hex_to_u32 (hash_pos + 120);
|
||||
|
||||
u8 *salt_pos = token.buf[1];
|
||||
int salt_len = token.len[1];
|
||||
|
||||
const bool parse_rc = parse_and_store_generic_salt ((u8 *) salt->salt_buf, (int *) &salt->salt_len, salt_pos, salt_len, hashconfig);
|
||||
|
||||
if (parse_rc == false) return (PARSER_SALT_LENGTH);
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int pbkdf2_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED hashconfig_t *hashconfig)
|
||||
{
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
@ -18608,6 +18671,7 @@ const char *strhashtype (const u32 hash_mode)
|
||||
case 11600: return HT_11600;
|
||||
case 11700: return HT_11700;
|
||||
case 11800: return HT_11800;
|
||||
case 11850: return HT_11850;
|
||||
case 11900: return HT_11900;
|
||||
case 12000: return HT_12000;
|
||||
case 12001: return HT_12001;
|
||||
@ -21149,7 +21213,7 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
digest_buf[6],
|
||||
digest_buf[7]);
|
||||
}
|
||||
else if (hash_mode == 11800)
|
||||
else if (hash_mode == 11800 || hash_mode == 11850)
|
||||
{
|
||||
snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x",
|
||||
byte_swap_32 (digest_buf[ 0]),
|
||||
@ -26410,14 +26474,14 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 11700: hashconfig->hash_type = HASH_TYPE_GOST_2012SBOG_256;
|
||||
case 11700: hashconfig->hash_type = HASH_TYPE_STREEBOG_256;
|
||||
hashconfig->salt_type = SALT_TYPE_NONE;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
|
||||
| OPTS_TYPE_PT_ADD01;
|
||||
hashconfig->kern_type = KERN_TYPE_GOST_2012SBOG_256;
|
||||
hashconfig->kern_type = KERN_TYPE_STREEBOG_256;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_8;
|
||||
hashconfig->parse_func = gost2012sbog_256_parse_hash;
|
||||
hashconfig->parse_func = streebog_256_parse_hash;
|
||||
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE;
|
||||
hashconfig->dgst_pos0 = 0;
|
||||
hashconfig->dgst_pos1 = 1;
|
||||
@ -26427,14 +26491,14 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 11800: hashconfig->hash_type = HASH_TYPE_GOST_2012SBOG_512;
|
||||
case 11800: hashconfig->hash_type = HASH_TYPE_STREEBOG_512;
|
||||
hashconfig->salt_type = SALT_TYPE_NONE;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
|
||||
| OPTS_TYPE_PT_ADD01;
|
||||
hashconfig->kern_type = KERN_TYPE_GOST_2012SBOG_512;
|
||||
hashconfig->kern_type = KERN_TYPE_STREEBOG_512;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_16;
|
||||
hashconfig->parse_func = gost2012sbog_512_parse_hash;
|
||||
hashconfig->parse_func = streebog_512_parse_hash;
|
||||
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE;
|
||||
hashconfig->dgst_pos0 = 0;
|
||||
hashconfig->dgst_pos1 = 1;
|
||||
@ -26444,6 +26508,23 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 11850: hashconfig->hash_type = HASH_TYPE_STREEBOG_512;
|
||||
hashconfig->salt_type = SALT_TYPE_GENERIC;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE
|
||||
| OPTS_TYPE_PT_ADD01;
|
||||
hashconfig->kern_type = KERN_TYPE_HMAC_STREEBOG_512_PW;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_16;
|
||||
hashconfig->parse_func = streebog_512s_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;
|
||||
hashconfig->st_hash = ST_HASH_11850;
|
||||
hashconfig->st_pass = ST_PASS_HASHCAT_PLAIN;
|
||||
break;
|
||||
|
||||
case 11900: hashconfig->hash_type = HASH_TYPE_PBKDF2_MD5;
|
||||
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
|
||||
|
||||
@ -192,6 +192,7 @@ static const char *const USAGE_BIG[] =
|
||||
" 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated",
|
||||
" 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated",
|
||||
" 1760 | HMAC-SHA512 (key = $salt) | Raw Hash, Authenticated",
|
||||
" 11850 | HMAC-Streebog-512 (key = $pass), big-endian | Raw Hash, Authenticated",
|
||||
" 14000 | DES (PT = $salt, key = $pass) | Raw Cipher, Known-Plaintext attack",
|
||||
" 14100 | 3DES (PT = $salt, key = $pass) | Raw Cipher, Known-Plaintext attack",
|
||||
" 14900 | Skip32 (PT = $salt, key = $pass) | Raw Cipher, Known-Plaintext attack",
|
||||
|
||||
@ -61,7 +61,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, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 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, 11700, 11800, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 17300, 17400, 17500, 17600, 17700, 17800, 17900, 18000, 18100, 18200, 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, 600, 900, 1000, 1100, 1300, 1400, 1410, 1411, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4520, 4521, 4522, 4600, 4700, 4800, 4900, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7000, 7100, 7200, 7300, 7400, 7500, 7700, 7701, 7800, 7801, 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, 11700, 11800, 11850, 11900, 12000, 12001, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000, 13100, 13200, 13300, 13400, 13500, 13600, 13800, 13900, 14000, 14100, 14400, 14700, 14800, 14900, 15000, 15100, 15200, 15300, 15400, 15500, 15600, 15700, 15900, 16000, 16100, 16200, 16300, 16400, 16500, 16600, 16700, 16800, 16900, 17300, 17400, 17500, 17600, 17700, 17800, 17900, 18000, 18100, 18200, 99999);
|
||||
|
||||
my %is_utf16le = map { $_ => 1 } qw (30 40 130 131 132 133 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800 11600 13500 13800);
|
||||
my %less_fifteen = map { $_ => 1 } qw (500 1600 1800 3200 6300 7400 10500 10700);
|
||||
@ -242,7 +242,7 @@ sub verify
|
||||
$word = substr ($line, $index + 1);
|
||||
}
|
||||
# hash:salt
|
||||
elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000 || $mode == 18100)
|
||||
elsif ($mode == 10 || $mode == 11 || $mode == 12 || $mode == 20 || $mode == 21 || $mode == 22 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 112 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1100 || $mode == 1410 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 2410 || $mode == 2611 || $mode == 2711 || $mode == 2811 || $mode == 3100 || $mode == 3610 || $mode == 3710 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4520 || $mode == 4521 || $mode == 4522 || $mode == 4900 || $mode == 5800 || $mode == 8400 || $mode == 11000 || $mode == 11850 || $mode == 12600 || $mode == 13500 || $mode == 13800 || $mode == 13900 || $mode == 14000 || $mode == 14100 || $mode == 14400 || $mode == 14900 || $mode == 15000 || $mode == 18100)
|
||||
{
|
||||
# get hash
|
||||
my $index1 = index ($line, ":");
|
||||
@ -3592,7 +3592,7 @@ sub passthrough
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, "");
|
||||
}
|
||||
elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 18100)
|
||||
elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3800 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 4900 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11850 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 18100)
|
||||
{
|
||||
my $salt_len = get_random_num (1, 15);
|
||||
|
||||
@ -4138,7 +4138,7 @@ sub single
|
||||
}
|
||||
}
|
||||
}
|
||||
elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500 || $mode == 18100)
|
||||
elsif ($mode == 10 || $mode == 20 || $mode == 23 || $mode == 30 || $mode == 40 || $mode == 50 || $mode == 60 || $mode == 110 || $mode == 120 || $mode == 121 || $mode == 130 || $mode == 140 || $mode == 150 || $mode == 160 || $mode == 1410 || $mode == 1411 || $mode == 1420 || $mode == 1430 || $mode == 1440 || $mode == 1450 || $mode == 1460 || $mode == 1710 || $mode == 1711 || $mode == 1720 || $mode == 1730 || $mode == 1740 || $mode == 1750 || $mode == 1760 || $mode == 3610 || $mode == 3710 || $mode == 3711 || $mode == 3720 || $mode == 3910 || $mode == 4010 || $mode == 4110 || $mode == 4210 || $mode == 8900 || $mode == 10000 || $mode == 10200 || $mode == 10900 || $mode == 11850 || $mode == 11900 || $mode == 12000 || $mode == 12100 || $mode == 16500 || $mode == 18100)
|
||||
{
|
||||
my $salt_len = get_random_num (1, 15);
|
||||
|
||||
@ -7907,6 +7907,9 @@ sub gen_hash
|
||||
}
|
||||
elsif ($mode == 11700)
|
||||
{
|
||||
# PyGOST outputs digests in little-endian order, while the kernels
|
||||
# expect them in big-endian; hence the digest[::-1] mirroring.
|
||||
# Using sys.stdout.write instead of print to disable \n character.
|
||||
my $python_code = <<"END_CODE";
|
||||
|
||||
import binascii
|
||||
@ -7933,6 +7936,27 @@ END_CODE
|
||||
|
||||
$tmp_hash = `python2 -c '$python_code'`;
|
||||
}
|
||||
elsif ($mode == 11850)
|
||||
{
|
||||
#todo: python hmac; look at salt value?
|
||||
#$hash_buf = hmac_hex ($salt_buf, $word_buf, \&sha512, 128);
|
||||
#$tmp_hash = sprintf ("%s:%s", $hash_buf, $salt_buf);
|
||||
my $python_code = <<"END_CODE";
|
||||
|
||||
import binascii
|
||||
import hmac
|
||||
import sys
|
||||
from pygost import gost34112012512
|
||||
key = b"$word_buf"
|
||||
msg = b"$salt_buf"
|
||||
digest = hmac.new(key, msg, gost34112012512).digest()
|
||||
sys.stdout.write(binascii.hexlify(digest[::-1]))
|
||||
|
||||
END_CODE
|
||||
|
||||
$hash_buf = `python2 -c '$python_code'`;
|
||||
$tmp_hash = sprintf ("%s:%s", $hash_buf, $salt_buf);
|
||||
}
|
||||
elsif ($mode == 11900)
|
||||
{
|
||||
my $iterations = 1000;
|
||||
|
||||
@ -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 600 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 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 7701 7800 7801 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 11700 11800 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100 18200 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 600 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 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 7701 7800 7801 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 11700 11800 11850 11900 12000 12001 12100 12200 12300 12400 12600 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 13900 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 15400 15500 15600 15700 15900 16000 16100 16200 16300 16400 16500 16600 16700 16800 16900 17300 17400 17500 17600 17700 17800 17900 18000 18100 18200 99999"
|
||||
|
||||
#ATTACK_MODES="0 1 3 6 7"
|
||||
ATTACK_MODES="0 1 3 7"
|
||||
|
||||
Loading…
Reference in New Issue
Block a user