mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-23 07:08:19 +00:00
Jens Steube
GitHub
commit
f2ad095191
@ -779,6 +779,15 @@ typedef struct luks_tmp
|
||||
|
||||
} luks_tmp_t;
|
||||
|
||||
typedef struct chacha20
|
||||
{
|
||||
u32 iv[2];
|
||||
u32 plain[2];
|
||||
u32 position[2];
|
||||
u32 offset;
|
||||
|
||||
} chacha20_t;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
int V;
|
||||
|
||||
364
OpenCL/m15400_a0.cl
Normal file
364
OpenCL/m15400_a0.cl
Normal file
@ -0,0 +1,364 @@
|
||||
/**
|
||||
* 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_simd.cl"
|
||||
|
||||
#define CHACHA_CONST_00 0x61707865
|
||||
#define CHACHA_CONST_01 0x3320646e
|
||||
#define CHACHA_CONST_02 0x79622d32
|
||||
#define CHACHA_CONST_03 0x6b206574
|
||||
|
||||
#define QR(a, b, c, d) \
|
||||
do { \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 16); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 12); \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 8); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 7); \
|
||||
} while (0);
|
||||
|
||||
void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[2], const u32 offset, const u32 iv[2], const u32 plain[4], u32x digest[4])
|
||||
{
|
||||
/**
|
||||
* Key expansion
|
||||
*/
|
||||
|
||||
u32x ctx[16];
|
||||
|
||||
ctx[ 0] = CHACHA_CONST_00;
|
||||
ctx[ 1] = CHACHA_CONST_01;
|
||||
ctx[ 2] = CHACHA_CONST_02;
|
||||
ctx[ 3] = CHACHA_CONST_03;
|
||||
ctx[ 4] = w0[0];
|
||||
ctx[ 5] = w0[1];
|
||||
ctx[ 6] = w0[2];
|
||||
ctx[ 7] = w0[3];
|
||||
ctx[ 8] = w1[0];
|
||||
ctx[ 9] = w1[1];
|
||||
ctx[10] = w1[2];
|
||||
ctx[11] = w1[3];
|
||||
ctx[12] = position[0];
|
||||
ctx[13] = position[1];
|
||||
ctx[14] = iv[1];
|
||||
ctx[15] = iv[0];
|
||||
|
||||
/**
|
||||
* Generate 64 byte keystream
|
||||
*/
|
||||
|
||||
u32x x[32];
|
||||
|
||||
x[ 0] = ctx[ 0];
|
||||
x[ 1] = ctx[ 1];
|
||||
x[ 2] = ctx[ 2];
|
||||
x[ 3] = ctx[ 3];
|
||||
x[ 4] = ctx[ 4];
|
||||
x[ 5] = ctx[ 5];
|
||||
x[ 6] = ctx[ 6];
|
||||
x[ 7] = ctx[ 7];
|
||||
x[ 8] = ctx[ 8];
|
||||
x[ 9] = ctx[ 9];
|
||||
x[10] = ctx[10];
|
||||
x[11] = ctx[11];
|
||||
x[12] = ctx[12];
|
||||
x[13] = ctx[13];
|
||||
x[14] = ctx[14];
|
||||
x[15] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(0, 4, 8, 12);
|
||||
QR(1, 5, 9, 13);
|
||||
QR(2, 6, 10, 14);
|
||||
QR(3, 7, 11, 15);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(0, 5, 10, 15);
|
||||
QR(1, 6, 11, 12);
|
||||
QR(2, 7, 8, 13);
|
||||
QR(3, 4, 9, 14);
|
||||
}
|
||||
|
||||
x[ 0] += ctx[ 0];
|
||||
x[ 1] += ctx[ 1];
|
||||
x[ 2] += ctx[ 2];
|
||||
x[ 3] += ctx[ 3];
|
||||
x[ 4] += ctx[ 4];
|
||||
x[ 5] += ctx[ 5];
|
||||
x[ 6] += ctx[ 6];
|
||||
x[ 7] += ctx[ 7];
|
||||
x[ 8] += ctx[ 8];
|
||||
x[ 9] += ctx[ 9];
|
||||
x[10] += ctx[10];
|
||||
x[11] += ctx[11];
|
||||
x[12] += ctx[12];
|
||||
x[13] += ctx[13];
|
||||
x[14] += ctx[14];
|
||||
x[15] += ctx[15];
|
||||
|
||||
if (offset > 56)
|
||||
{
|
||||
/**
|
||||
* Generate a second 64 byte keystream
|
||||
*/
|
||||
|
||||
ctx[12]++;
|
||||
|
||||
if (all(ctx[12] == 0)) ctx[13]++;
|
||||
|
||||
x[16] = ctx[ 0];
|
||||
x[17] = ctx[ 1];
|
||||
x[18] = ctx[ 2];
|
||||
x[19] = ctx[ 3];
|
||||
x[20] = ctx[ 4];
|
||||
x[21] = ctx[ 5];
|
||||
x[22] = ctx[ 6];
|
||||
x[23] = ctx[ 7];
|
||||
x[24] = ctx[ 8];
|
||||
x[25] = ctx[ 9];
|
||||
x[26] = ctx[10];
|
||||
x[27] = ctx[11];
|
||||
x[28] = ctx[12];
|
||||
x[29] = ctx[13];
|
||||
x[30] = ctx[14];
|
||||
x[31] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(16, 20, 24, 28);
|
||||
QR(17, 21, 25, 29);
|
||||
QR(18, 22, 26, 30);
|
||||
QR(19, 23, 27, 31);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(16, 21, 26, 31);
|
||||
QR(17, 22, 27, 28);
|
||||
QR(18, 23, 24, 29);
|
||||
QR(19, 20, 25, 30);
|
||||
}
|
||||
|
||||
x[16] += ctx[ 0];
|
||||
x[17] += ctx[ 1];
|
||||
x[18] += ctx[ 2];
|
||||
x[19] += ctx[ 3];
|
||||
x[20] += ctx[ 4];
|
||||
x[21] += ctx[ 5];
|
||||
x[22] += ctx[ 6];
|
||||
x[23] += ctx[ 7];
|
||||
x[24] += ctx[ 8];
|
||||
x[25] += ctx[ 9];
|
||||
x[26] += ctx[10];
|
||||
x[27] += ctx[11];
|
||||
x[28] += ctx[12];
|
||||
x[29] += ctx[13];
|
||||
x[30] += ctx[14];
|
||||
x[31] += ctx[15];
|
||||
}
|
||||
|
||||
/**
|
||||
* Encrypt plaintext with keystream
|
||||
*/
|
||||
|
||||
const u32 index = offset / 4;
|
||||
const u32 remain = offset % 4;
|
||||
|
||||
digest[0] = plain[1];
|
||||
digest[1] = plain[0];
|
||||
|
||||
if (remain > 0)
|
||||
{
|
||||
digest[1] ^= x[index + 0] >> ( 0 + remain * 8);
|
||||
digest[1] ^= x[index + 1] << (32 - remain * 8);
|
||||
|
||||
digest[0] ^= x[index + 1] >> ( 0 + remain * 8);
|
||||
digest[0] ^= x[index + 2] << (32 - remain * 8);
|
||||
}
|
||||
else
|
||||
{
|
||||
digest[1] ^= x[index + 0];
|
||||
digest[0] ^= x[index + 1];
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[0];
|
||||
pw_buf0[1] = pws[gid].i[1];
|
||||
pw_buf0[2] = pws[gid].i[2];
|
||||
pw_buf0[3] = pws[gid].i[3];
|
||||
pw_buf1[0] = pws[gid].i[4];
|
||||
pw_buf1[1] = pws[gid].i[5];
|
||||
pw_buf1[2] = pws[gid].i[6];
|
||||
pw_buf1[3] = pws[gid].i[7];
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
|
||||
const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0, w1, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_M_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[0];
|
||||
pw_buf0[1] = pws[gid].i[1];
|
||||
pw_buf0[2] = pws[gid].i[2];
|
||||
pw_buf0[3] = pws[gid].i[3];
|
||||
pw_buf1[0] = pws[gid].i[4];
|
||||
pw_buf1[1] = pws[gid].i[5];
|
||||
pw_buf1[2] = pws[gid].i[6];
|
||||
pw_buf1[3] = pws[gid].i[7];
|
||||
|
||||
const u32 pw_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* 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]
|
||||
};
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
u32x w0[4] = { 0 };
|
||||
u32x w1[4] = { 0 };
|
||||
|
||||
const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0, w1, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_S_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
468
OpenCL/m15400_a1.cl
Normal file
468
OpenCL/m15400_a1.cl
Normal file
@ -0,0 +1,468 @@
|
||||
/**
|
||||
* 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_simd.cl"
|
||||
|
||||
#define CHACHA_CONST_00 0x61707865
|
||||
#define CHACHA_CONST_01 0x3320646e
|
||||
#define CHACHA_CONST_02 0x79622d32
|
||||
#define CHACHA_CONST_03 0x6b206574
|
||||
|
||||
#define QR(a, b, c, d) \
|
||||
do { \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 16); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 12); \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 8); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 7); \
|
||||
} while (0);
|
||||
|
||||
void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[2], const u32 offset, const u32 iv[2], const u32 plain[4], u32x digest[4])
|
||||
{
|
||||
/**
|
||||
* Key expansion
|
||||
*/
|
||||
|
||||
u32x ctx[16];
|
||||
|
||||
ctx[ 0] = CHACHA_CONST_00;
|
||||
ctx[ 1] = CHACHA_CONST_01;
|
||||
ctx[ 2] = CHACHA_CONST_02;
|
||||
ctx[ 3] = CHACHA_CONST_03;
|
||||
ctx[ 4] = w0[0];
|
||||
ctx[ 5] = w0[1];
|
||||
ctx[ 6] = w0[2];
|
||||
ctx[ 7] = w0[3];
|
||||
ctx[ 8] = w1[0];
|
||||
ctx[ 9] = w1[1];
|
||||
ctx[10] = w1[2];
|
||||
ctx[11] = w1[3];
|
||||
ctx[12] = position[0];
|
||||
ctx[13] = position[1];
|
||||
ctx[14] = iv[1];
|
||||
ctx[15] = iv[0];
|
||||
|
||||
/**
|
||||
* Generate 64 byte keystream
|
||||
*/
|
||||
|
||||
u32x x[32];
|
||||
|
||||
x[ 0] = ctx[ 0];
|
||||
x[ 1] = ctx[ 1];
|
||||
x[ 2] = ctx[ 2];
|
||||
x[ 3] = ctx[ 3];
|
||||
x[ 4] = ctx[ 4];
|
||||
x[ 5] = ctx[ 5];
|
||||
x[ 6] = ctx[ 6];
|
||||
x[ 7] = ctx[ 7];
|
||||
x[ 8] = ctx[ 8];
|
||||
x[ 9] = ctx[ 9];
|
||||
x[10] = ctx[10];
|
||||
x[11] = ctx[11];
|
||||
x[12] = ctx[12];
|
||||
x[13] = ctx[13];
|
||||
x[14] = ctx[14];
|
||||
x[15] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(0, 4, 8, 12);
|
||||
QR(1, 5, 9, 13);
|
||||
QR(2, 6, 10, 14);
|
||||
QR(3, 7, 11, 15);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(0, 5, 10, 15);
|
||||
QR(1, 6, 11, 12);
|
||||
QR(2, 7, 8, 13);
|
||||
QR(3, 4, 9, 14);
|
||||
}
|
||||
|
||||
x[ 0] += ctx[ 0];
|
||||
x[ 1] += ctx[ 1];
|
||||
x[ 2] += ctx[ 2];
|
||||
x[ 3] += ctx[ 3];
|
||||
x[ 4] += ctx[ 4];
|
||||
x[ 5] += ctx[ 5];
|
||||
x[ 6] += ctx[ 6];
|
||||
x[ 7] += ctx[ 7];
|
||||
x[ 8] += ctx[ 8];
|
||||
x[ 9] += ctx[ 9];
|
||||
x[10] += ctx[10];
|
||||
x[11] += ctx[11];
|
||||
x[12] += ctx[12];
|
||||
x[13] += ctx[13];
|
||||
x[14] += ctx[14];
|
||||
x[15] += ctx[15];
|
||||
|
||||
if (offset > 56)
|
||||
{
|
||||
/**
|
||||
* Generate a second 64 byte keystream
|
||||
*/
|
||||
|
||||
ctx[12]++;
|
||||
|
||||
if (all(ctx[12] == 0)) ctx[13]++;
|
||||
|
||||
x[16] = ctx[ 0];
|
||||
x[17] = ctx[ 1];
|
||||
x[18] = ctx[ 2];
|
||||
x[19] = ctx[ 3];
|
||||
x[20] = ctx[ 4];
|
||||
x[21] = ctx[ 5];
|
||||
x[22] = ctx[ 6];
|
||||
x[23] = ctx[ 7];
|
||||
x[24] = ctx[ 8];
|
||||
x[25] = ctx[ 9];
|
||||
x[26] = ctx[10];
|
||||
x[27] = ctx[11];
|
||||
x[28] = ctx[12];
|
||||
x[29] = ctx[13];
|
||||
x[30] = ctx[14];
|
||||
x[31] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(16, 20, 24, 28);
|
||||
QR(17, 21, 25, 29);
|
||||
QR(18, 22, 26, 30);
|
||||
QR(19, 23, 27, 31);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(16, 21, 26, 31);
|
||||
QR(17, 22, 27, 28);
|
||||
QR(18, 23, 24, 29);
|
||||
QR(19, 20, 25, 30);
|
||||
}
|
||||
|
||||
x[16] += ctx[ 0];
|
||||
x[17] += ctx[ 1];
|
||||
x[18] += ctx[ 2];
|
||||
x[19] += ctx[ 3];
|
||||
x[20] += ctx[ 4];
|
||||
x[21] += ctx[ 5];
|
||||
x[22] += ctx[ 6];
|
||||
x[23] += ctx[ 7];
|
||||
x[24] += ctx[ 8];
|
||||
x[25] += ctx[ 9];
|
||||
x[26] += ctx[10];
|
||||
x[27] += ctx[11];
|
||||
x[28] += ctx[12];
|
||||
x[29] += ctx[13];
|
||||
x[30] += ctx[14];
|
||||
x[31] += ctx[15];
|
||||
}
|
||||
|
||||
/**
|
||||
* Encrypt plaintext with keystream
|
||||
*/
|
||||
|
||||
const u32 index = offset / 4;
|
||||
const u32 remain = offset % 4;
|
||||
|
||||
digest[0] = plain[1];
|
||||
digest[1] = plain[0];
|
||||
|
||||
if (remain > 0)
|
||||
{
|
||||
digest[1] ^= x[index + 0] >> ( 0 + remain * 8);
|
||||
digest[1] ^= x[index + 1] << (32 - remain * 8);
|
||||
|
||||
digest[0] ^= x[index + 1] >> ( 0 + remain * 8);
|
||||
digest[0] ^= x[index + 2] << (32 - remain * 8);
|
||||
}
|
||||
else
|
||||
{
|
||||
digest[1] ^= x[index + 0];
|
||||
digest[0] ^= x[index + 1];
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[0];
|
||||
pw_buf0[1] = pws[gid].i[1];
|
||||
pw_buf0[2] = pws[gid].i[2];
|
||||
pw_buf0[3] = pws[gid].i[3];
|
||||
pw_buf1[0] = pws[gid].i[4];
|
||||
pw_buf1[1] = pws[gid].i[5];
|
||||
pw_buf1[2] = pws[gid].i[6];
|
||||
pw_buf1[3] = pws[gid].i[7];
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x out_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
w0[2] = wordl0[2] | wordr0[2];
|
||||
w0[3] = wordl0[3] | wordr0[3];
|
||||
w1[0] = wordl1[0] | wordr1[0];
|
||||
w1[1] = wordl1[1] | wordr1[1];
|
||||
w1[2] = wordl1[2] | wordr1[2];
|
||||
w1[3] = wordl1[3] | wordr1[3];
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0, w1, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_M_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
|
||||
if (gid >= gid_max) return;
|
||||
|
||||
u32 pw_buf0[4];
|
||||
u32 pw_buf1[4];
|
||||
|
||||
pw_buf0[0] = pws[gid].i[0];
|
||||
pw_buf0[1] = pws[gid].i[1];
|
||||
pw_buf0[2] = pws[gid].i[2];
|
||||
pw_buf0[3] = pws[gid].i[3];
|
||||
pw_buf1[0] = pws[gid].i[4];
|
||||
pw_buf1[1] = pws[gid].i[5];
|
||||
pw_buf1[2] = pws[gid].i[6];
|
||||
pw_buf1[3] = pws[gid].i[7];
|
||||
|
||||
const u32 pw_l_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* 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]
|
||||
};
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos);
|
||||
|
||||
const u32x out_len = pw_l_len + pw_r_len;
|
||||
|
||||
/**
|
||||
* concat password candidate
|
||||
*/
|
||||
|
||||
u32x wordl0[4] = { 0 };
|
||||
u32x wordl1[4] = { 0 };
|
||||
u32x wordl2[4] = { 0 };
|
||||
u32x wordl3[4] = { 0 };
|
||||
|
||||
wordl0[0] = pw_buf0[0];
|
||||
wordl0[1] = pw_buf0[1];
|
||||
wordl0[2] = pw_buf0[2];
|
||||
wordl0[3] = pw_buf0[3];
|
||||
wordl1[0] = pw_buf1[0];
|
||||
wordl1[1] = pw_buf1[1];
|
||||
wordl1[2] = pw_buf1[2];
|
||||
wordl1[3] = pw_buf1[3];
|
||||
|
||||
u32x wordr0[4] = { 0 };
|
||||
u32x wordr1[4] = { 0 };
|
||||
u32x wordr2[4] = { 0 };
|
||||
u32x wordr3[4] = { 0 };
|
||||
|
||||
wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
|
||||
wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
|
||||
wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
|
||||
wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
|
||||
wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
|
||||
wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
|
||||
wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
|
||||
wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
|
||||
|
||||
if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
|
||||
}
|
||||
else
|
||||
{
|
||||
switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
|
||||
}
|
||||
|
||||
u32x w0[4];
|
||||
u32x w1[4];
|
||||
|
||||
w0[0] = wordl0[0] | wordr0[0];
|
||||
w0[1] = wordl0[1] | wordr0[1];
|
||||
w0[2] = wordl0[2] | wordr0[2];
|
||||
w0[3] = wordl0[3] | wordr0[3];
|
||||
w1[0] = wordl1[0] | wordr1[0];
|
||||
w1[1] = wordl1[1] | wordr1[1];
|
||||
w1[2] = wordl1[2] | wordr1[2];
|
||||
w1[3] = wordl1[3] | wordr1[3];
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0, w1, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_S_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
383
OpenCL/m15400_a3.cl
Normal file
383
OpenCL/m15400_a3.cl
Normal file
@ -0,0 +1,383 @@
|
||||
/**
|
||||
* 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 CHACHA_CONST_00 0x61707865
|
||||
#define CHACHA_CONST_01 0x3320646e
|
||||
#define CHACHA_CONST_02 0x79622d32
|
||||
#define CHACHA_CONST_03 0x6b206574
|
||||
|
||||
#define QR(a, b, c, d) \
|
||||
do { \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 16); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 12); \
|
||||
x[a] = x[a] + x[b]; \
|
||||
x[d] = rotl32(x[d] ^ x[a], 8); \
|
||||
x[c] = x[c] + x[d]; \
|
||||
x[b] = rotl32(x[b] ^ x[c], 7); \
|
||||
} while (0);
|
||||
|
||||
void chacha20_transform (const u32x w0[4], const u32x w1[4], const u32 position[2], const u32 offset, const u32 iv[2], const u32 plain[4], u32x digest[4])
|
||||
{
|
||||
/**
|
||||
* Key expansion
|
||||
*/
|
||||
|
||||
u32x ctx[16];
|
||||
|
||||
ctx[ 0] = CHACHA_CONST_00;
|
||||
ctx[ 1] = CHACHA_CONST_01;
|
||||
ctx[ 2] = CHACHA_CONST_02;
|
||||
ctx[ 3] = CHACHA_CONST_03;
|
||||
ctx[ 4] = w0[0];
|
||||
ctx[ 5] = w0[1];
|
||||
ctx[ 6] = w0[2];
|
||||
ctx[ 7] = w0[3];
|
||||
ctx[ 8] = w1[0];
|
||||
ctx[ 9] = w1[1];
|
||||
ctx[10] = w1[2];
|
||||
ctx[11] = w1[3];
|
||||
ctx[12] = position[0];
|
||||
ctx[13] = position[1];
|
||||
ctx[14] = iv[1];
|
||||
ctx[15] = iv[0];
|
||||
|
||||
/**
|
||||
* Generate 64 byte keystream
|
||||
*/
|
||||
|
||||
u32x x[32];
|
||||
|
||||
x[ 0] = ctx[ 0];
|
||||
x[ 1] = ctx[ 1];
|
||||
x[ 2] = ctx[ 2];
|
||||
x[ 3] = ctx[ 3];
|
||||
x[ 4] = ctx[ 4];
|
||||
x[ 5] = ctx[ 5];
|
||||
x[ 6] = ctx[ 6];
|
||||
x[ 7] = ctx[ 7];
|
||||
x[ 8] = ctx[ 8];
|
||||
x[ 9] = ctx[ 9];
|
||||
x[10] = ctx[10];
|
||||
x[11] = ctx[11];
|
||||
x[12] = ctx[12];
|
||||
x[13] = ctx[13];
|
||||
x[14] = ctx[14];
|
||||
x[15] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(0, 4, 8, 12);
|
||||
QR(1, 5, 9, 13);
|
||||
QR(2, 6, 10, 14);
|
||||
QR(3, 7, 11, 15);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(0, 5, 10, 15);
|
||||
QR(1, 6, 11, 12);
|
||||
QR(2, 7, 8, 13);
|
||||
QR(3, 4, 9, 14);
|
||||
}
|
||||
|
||||
x[ 0] += ctx[ 0];
|
||||
x[ 1] += ctx[ 1];
|
||||
x[ 2] += ctx[ 2];
|
||||
x[ 3] += ctx[ 3];
|
||||
x[ 4] += ctx[ 4];
|
||||
x[ 5] += ctx[ 5];
|
||||
x[ 6] += ctx[ 6];
|
||||
x[ 7] += ctx[ 7];
|
||||
x[ 8] += ctx[ 8];
|
||||
x[ 9] += ctx[ 9];
|
||||
x[10] += ctx[10];
|
||||
x[11] += ctx[11];
|
||||
x[12] += ctx[12];
|
||||
x[13] += ctx[13];
|
||||
x[14] += ctx[14];
|
||||
x[15] += ctx[15];
|
||||
|
||||
if (offset > 56)
|
||||
{
|
||||
/**
|
||||
* Generate a second 64 byte keystream
|
||||
*/
|
||||
|
||||
ctx[12]++;
|
||||
|
||||
if (all(ctx[12] == 0)) ctx[13]++;
|
||||
|
||||
x[16] = ctx[ 0];
|
||||
x[17] = ctx[ 1];
|
||||
x[18] = ctx[ 2];
|
||||
x[19] = ctx[ 3];
|
||||
x[20] = ctx[ 4];
|
||||
x[21] = ctx[ 5];
|
||||
x[22] = ctx[ 6];
|
||||
x[23] = ctx[ 7];
|
||||
x[24] = ctx[ 8];
|
||||
x[25] = ctx[ 9];
|
||||
x[26] = ctx[10];
|
||||
x[27] = ctx[11];
|
||||
x[28] = ctx[12];
|
||||
x[29] = ctx[13];
|
||||
x[30] = ctx[14];
|
||||
x[31] = ctx[15];
|
||||
|
||||
#pragma unroll
|
||||
for (u8 i = 0; i < 10; i++)
|
||||
{
|
||||
/* Column round */
|
||||
QR(16, 20, 24, 28);
|
||||
QR(17, 21, 25, 29);
|
||||
QR(18, 22, 26, 30);
|
||||
QR(19, 23, 27, 31);
|
||||
|
||||
/* Diagonal round */
|
||||
QR(16, 21, 26, 31);
|
||||
QR(17, 22, 27, 28);
|
||||
QR(18, 23, 24, 29);
|
||||
QR(19, 20, 25, 30);
|
||||
}
|
||||
|
||||
x[16] += ctx[ 0];
|
||||
x[17] += ctx[ 1];
|
||||
x[18] += ctx[ 2];
|
||||
x[19] += ctx[ 3];
|
||||
x[20] += ctx[ 4];
|
||||
x[21] += ctx[ 5];
|
||||
x[22] += ctx[ 6];
|
||||
x[23] += ctx[ 7];
|
||||
x[24] += ctx[ 8];
|
||||
x[25] += ctx[ 9];
|
||||
x[26] += ctx[10];
|
||||
x[27] += ctx[11];
|
||||
x[28] += ctx[12];
|
||||
x[29] += ctx[13];
|
||||
x[30] += ctx[14];
|
||||
x[31] += ctx[15];
|
||||
}
|
||||
|
||||
/**
|
||||
* Encrypt plaintext with keystream
|
||||
*/
|
||||
|
||||
const u32 index = offset / 4;
|
||||
const u32 remain = offset % 4;
|
||||
|
||||
digest[0] = plain[1];
|
||||
digest[1] = plain[0];
|
||||
|
||||
if (remain > 0)
|
||||
{
|
||||
digest[1] ^= x[index + 0] >> ( 0 + remain * 8);
|
||||
digest[1] ^= x[index + 1] << (32 - remain * 8);
|
||||
|
||||
digest[0] ^= x[index + 1] >> ( 0 + remain * 8);
|
||||
digest[0] ^= x[index + 2] << (32 - remain * 8);
|
||||
}
|
||||
else
|
||||
{
|
||||
digest[1] ^= x[index + 0];
|
||||
digest[0] ^= x[index + 1];
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global 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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
|
||||
w0[0] = pws[gid].i[0];
|
||||
w0[1] = pws[gid].i[1];
|
||||
w0[2] = pws[gid].i[2];
|
||||
w0[3] = pws[gid].i[3];
|
||||
w1[0] = pws[gid].i[4];
|
||||
w1[1] = pws[gid].i[5];
|
||||
w1[2] = pws[gid].i[6];
|
||||
w1[3] = pws[gid].i[7];
|
||||
|
||||
u32x out_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32 w0l = pws[gid].i[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 w0x = w0l | w0r;
|
||||
|
||||
u32x w0_t[4];
|
||||
u32x w1_t[4];
|
||||
|
||||
w0_t[0] = w0x;
|
||||
w0_t[1] = w0[1];
|
||||
w0_t[2] = w0[2];
|
||||
w0_t[3] = w0[3];
|
||||
w1_t[0] = w1[0];
|
||||
w1_t[1] = w1[1];
|
||||
w1_t[2] = w1[2];
|
||||
w1_t[3] = w1[3];
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0_t, w1_t, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_M_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global 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 chacha20_t *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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
*/
|
||||
|
||||
const u32 gid = get_global_id (0);
|
||||
const u32 lid = get_local_id (0);
|
||||
|
||||
u32 w0[4];
|
||||
u32 w1[4];
|
||||
|
||||
w0[0] = pws[gid].i[0];
|
||||
w0[1] = pws[gid].i[1];
|
||||
w0[2] = pws[gid].i[2];
|
||||
w0[3] = pws[gid].i[3];
|
||||
w1[0] = pws[gid].i[4];
|
||||
w1[1] = pws[gid].i[5];
|
||||
w1[2] = pws[gid].i[6];
|
||||
w1[3] = pws[gid].i[7];
|
||||
|
||||
u32 out_len = pws[gid].pw_len;
|
||||
|
||||
/**
|
||||
* Salt prep
|
||||
*/
|
||||
|
||||
u32 iv[2] = { 0 };
|
||||
u32 plain[2] = { 0 };
|
||||
u32 position[2] = { 0 };
|
||||
u32 offset = 0;
|
||||
|
||||
position[0] = esalt_bufs[digests_offset].position[0];
|
||||
position[1] = esalt_bufs[digests_offset].position[1];
|
||||
|
||||
offset = esalt_bufs[digests_offset].offset;
|
||||
|
||||
iv[0] = esalt_bufs[digests_offset].iv[0];
|
||||
iv[1] = esalt_bufs[digests_offset].iv[1];
|
||||
|
||||
plain[0] = esalt_bufs[digests_offset].plain[0];
|
||||
plain[1] = esalt_bufs[digests_offset].plain[1];
|
||||
|
||||
/**
|
||||
* 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]
|
||||
};
|
||||
|
||||
/**
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32 w0l = pws[gid].i[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 w0x = w0l | w0r;
|
||||
|
||||
u32x w0_t[4];
|
||||
u32x w1_t[4];
|
||||
|
||||
w0_t[0] = w0x;
|
||||
w0_t[1] = w0[1];
|
||||
w0_t[2] = w0[2];
|
||||
w0_t[3] = w0[3];
|
||||
w1_t[0] = w1[0];
|
||||
w1_t[1] = w1[1];
|
||||
w1_t[2] = w1[2];
|
||||
w1_t[3] = w1[3];
|
||||
|
||||
u32x digest[4] = { 0 };
|
||||
|
||||
chacha20_transform (w0_t, w1_t, position, offset, iv, plain, digest);
|
||||
|
||||
const u32x r0 = digest[0];
|
||||
const u32x r1 = digest[1];
|
||||
const u32x r2 = digest[2];
|
||||
const u32x r3 = digest[3];
|
||||
|
||||
COMPARE_S_SIMD(r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m15400_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
|
||||
__kernel void m15400_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_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 chacha20_t *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)
|
||||
{
|
||||
}
|
||||
@ -7,6 +7,7 @@
|
||||
- Added hash-mode 600 = Blake2-512
|
||||
- Added hash-mode 15200 = Blockchain, My Wallet, V2
|
||||
- Added hash-mode 15300 = DPAPI masterkey file v1 and v2
|
||||
- Added hash-mode 15400 = Chacha20
|
||||
|
||||
##
|
||||
## Features
|
||||
|
||||
@ -53,6 +53,7 @@ NVIDIA GPUs require "NVIDIA Driver" (367.x or later)
|
||||
- Whirlpool
|
||||
- DES (PT = $salt, key = $pass)
|
||||
- 3DES (PT = $salt, key = $pass)
|
||||
- Chacha20
|
||||
- GOST R 34.11-94
|
||||
- GOST R 34.11-2012 (Streebog) 256-bit
|
||||
- GOST R 34.11-2012 (Streebog) 512-bit
|
||||
|
||||
@ -176,7 +176,7 @@ _hashcat ()
|
||||
{
|
||||
local VERSION=3.5.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 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 15200 15300"
|
||||
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 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 15200 15300 15400"
|
||||
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"
|
||||
|
||||
@ -1226,6 +1226,8 @@ typedef enum display_len
|
||||
DISPLAY_LEN_MAX_15200 = 1 + 10 + 1 + 2 + 1 + 8 + 1 + 5 + 1 + 20000,
|
||||
DISPLAY_LEN_MIN_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 10 + 1 + 4 + 1 + 4 + 1 + 1 + 1 + 32 + 1 + 3 + 1 + 128,
|
||||
DISPLAY_LEN_MAX_15300 = 1 + 7 + 1 + 1 + 1 + 1 + 1 + 100 + 1 + 6 + 1 + 6 + 1 + 10 + 1 + 32 + 1 + 4 + 1 + 512,
|
||||
DISPLAY_LEN_MIN_15400 = 10 + 1 + 16 + 1 + 1 + 1 + 16 + 1 + 16 + 1 + 16,
|
||||
DISPLAY_LEN_MAX_15400 = 10 + 1 + 16 + 1 + 2 + 1 + 16 + 1 + 16 + 1 + 16,
|
||||
DISPLAY_LEN_MIN_99999 = 1,
|
||||
DISPLAY_LEN_MAX_99999 = 55,
|
||||
|
||||
@ -1557,6 +1559,7 @@ typedef enum kern_type
|
||||
KERN_TYPE_FILEZILLA_SERVER = 15000,
|
||||
KERN_TYPE_NETBSD_SHA1CRYPT = 15100,
|
||||
KERN_TYPE_DPAPIMK = 15300,
|
||||
KERN_TYPE_CHACHA20 = 15400,
|
||||
KERN_TYPE_PLAINTEXT = 99999,
|
||||
|
||||
} kern_type_t;
|
||||
@ -1649,6 +1652,7 @@ int postgresql_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_bu
|
||||
int netscreen_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int keccak_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int blake2b_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int lm_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int md4_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
int md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
|
||||
|
||||
@ -693,6 +693,15 @@ typedef struct
|
||||
|
||||
} blake2_t;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
u32 iv[2];
|
||||
u32 plain[2];
|
||||
u32 position[2];
|
||||
u32 offset;
|
||||
|
||||
} chacha20_t;
|
||||
|
||||
typedef struct
|
||||
{
|
||||
u32 salt_buf[16];
|
||||
|
||||
@ -6,7 +6,7 @@
|
||||
#include "common.h"
|
||||
#include "benchmark.h"
|
||||
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 151;
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_CNT = 152;
|
||||
|
||||
const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] =
|
||||
{
|
||||
@ -160,5 +160,6 @@ const unsigned int DEFAULT_BENCHMARK_ALGORITHMS_BUF[] =
|
||||
12700,
|
||||
15200,
|
||||
13400,
|
||||
125
|
||||
125,
|
||||
15400
|
||||
};
|
||||
|
||||
107
src/interface.c
107
src/interface.c
@ -241,6 +241,8 @@ static const char HT_15000[] = "FileZilla Server >= 0.9.55";
|
||||
static const char HT_15100[] = "Juniper/NetBSD sha1crypt";
|
||||
static const char HT_15200[] = "Blockchain, My Wallet, V2";
|
||||
static const char HT_15300[] = "DPAPI masterkey file v1 and v2";
|
||||
static const char HT_15400[] = "Chacha20";
|
||||
|
||||
static const char HT_99999[] = "Plaintext";
|
||||
|
||||
static const char HT_00011[] = "Joomla < 2.5.18";
|
||||
@ -381,6 +383,7 @@ static const char SIGNATURE_FORTIGATE[] = "AK1";
|
||||
static const char SIGNATURE_ATLASSIAN[] = "{PKCS5S2}";
|
||||
static const char SIGNATURE_NETBSD_SHA1CRYPT[] = "$sha1$";
|
||||
static const char SIGNATURE_BLAKE2B[] = "$BLAKE2$";
|
||||
static const char SIGNATURE_CHACHA20[] = "$Chacha20$";
|
||||
|
||||
/**
|
||||
* decoder / encoder
|
||||
@ -5468,6 +5471,71 @@ int blake2b_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_15400) || (input_len > DISPLAY_LEN_MAX_15400)) return (PARSER_GLOBAL_LENGTH);
|
||||
|
||||
if (memcmp (SIGNATURE_CHACHA20, input_buf, 10)) return (PARSER_SIGNATURE_UNMATCHED);
|
||||
|
||||
u32 *digest = (u32 *) hash_buf->digest;
|
||||
|
||||
chacha20_t *chacha20 = (chacha20_t *) hash_buf->esalt;
|
||||
|
||||
salt_t *salt = (salt_t *) hash_buf->salt;
|
||||
|
||||
u8 *position_marker = (u8 *) strchr ((const char *) input_buf, '*') + 1;
|
||||
if (position_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
if (is_valid_hex_string (position_marker, 16) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
u8 *offset_marker = (u8 *) strchr ((const char *) position_marker, '*') + 1;
|
||||
if (offset_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
|
||||
int offset = atoi ((char*) offset_marker);
|
||||
if (offset > 63) return (PARSER_SALT_VALUE);
|
||||
|
||||
u8 *iv_marker = (u8 *) strchr ((const char *) offset_marker, '*') + 1;
|
||||
if (iv_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
if (is_valid_hex_string (iv_marker, 16) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
u8 *plain_marker = (u8 *) strchr ((const char *) iv_marker, '*') + 1;
|
||||
if (plain_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
if (is_valid_hex_string (plain_marker, 16) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
u8 *cipher_marker = (u8 *) strchr ((const char *) plain_marker, '*') + 1;
|
||||
if (cipher_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
|
||||
if (is_valid_hex_string (cipher_marker, 16) == false) return (PARSER_SALT_ENCODING);
|
||||
|
||||
chacha20->iv[0] = hex_to_u32 ((const u8 *) iv_marker + 8);
|
||||
chacha20->iv[1] = hex_to_u32 ((const u8 *) iv_marker + 0);
|
||||
|
||||
chacha20->plain[0] = hex_to_u32 ((const u8 *) plain_marker + 0);
|
||||
chacha20->plain[1] = hex_to_u32 ((const u8 *) plain_marker + 8);
|
||||
|
||||
chacha20->position[0] = hex_to_u32 ((const u8 *) position_marker + 0);
|
||||
chacha20->position[1] = hex_to_u32 ((const u8 *) position_marker + 8);
|
||||
|
||||
chacha20->offset = offset;
|
||||
|
||||
/* some fake salt for the sorting mechanisms */
|
||||
|
||||
salt->salt_buf[0] = chacha20->iv[0];
|
||||
salt->salt_buf[1] = chacha20->iv[1];
|
||||
salt->salt_buf[2] = chacha20->plain[0];
|
||||
salt->salt_buf[3] = chacha20->plain[1];
|
||||
salt->salt_buf[4] = chacha20->position[0];
|
||||
salt->salt_buf[5] = chacha20->position[1];
|
||||
salt->salt_buf[6] = chacha20->offset;
|
||||
salt->salt_buf[7] = 0;
|
||||
salt->salt_len = 32;
|
||||
|
||||
/* Store cipher for search mechanism */
|
||||
digest[0] = hex_to_u32 ((const u8 *) cipher_marker + 8);
|
||||
digest[1] = hex_to_u32 ((const u8 *) cipher_marker + 0);
|
||||
|
||||
return (PARSER_OK);
|
||||
}
|
||||
|
||||
|
||||
int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
|
||||
{
|
||||
if ((input_len < DISPLAY_LEN_MIN_5300) || (input_len > DISPLAY_LEN_MAX_5300)) return (PARSER_GLOBAL_LENGTH);
|
||||
@ -15176,6 +15244,7 @@ char *strhashtype (const u32 hash_mode)
|
||||
case 15100: return ((char *) HT_15100);
|
||||
case 15200: return ((char *) HT_15200);
|
||||
case 15300: return ((char *) HT_15300);
|
||||
case 15400: return ((char *) HT_15400);
|
||||
case 99999: return ((char *) HT_99999);
|
||||
}
|
||||
|
||||
@ -18534,6 +18603,25 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
|
||||
byte_swap_32(ptr[14]),
|
||||
byte_swap_32(ptr[15]));
|
||||
}
|
||||
else if (hash_type == HASH_TYPE_CHACHA20)
|
||||
{
|
||||
u32 *ptr = digest_buf;
|
||||
|
||||
const chacha20_t *chacha20_tmp = (const chacha20_t *) esalts_buf;
|
||||
const chacha20_t *chacha20 = &chacha20_tmp[digest_cur];
|
||||
|
||||
snprintf (out_buf, out_len - 1, "%s*%08x%08x*%d*%08x%08x*%08x%08x*%08x%08x",
|
||||
SIGNATURE_CHACHA20,
|
||||
byte_swap_32(chacha20->position[0]),
|
||||
byte_swap_32(chacha20->position[1]),
|
||||
chacha20->offset,
|
||||
byte_swap_32(chacha20->iv[1]),
|
||||
byte_swap_32(chacha20->iv[0]),
|
||||
byte_swap_32(chacha20->plain[0]),
|
||||
byte_swap_32(chacha20->plain[1]),
|
||||
ptr[1],
|
||||
ptr[0]);
|
||||
}
|
||||
else if (hash_type == HASH_TYPE_RIPEMD160)
|
||||
{
|
||||
snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x%08x",
|
||||
@ -22685,6 +22773,22 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
hashconfig->dgst_pos3 = 3;
|
||||
break;
|
||||
|
||||
case 15400: hashconfig->hash_type = HASH_TYPE_CHACHA20;
|
||||
hashconfig->salt_type = SALT_TYPE_EMBEDDED;
|
||||
hashconfig->attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
|
||||
hashconfig->opts_type = OPTS_TYPE_PT_GENERATE_LE;
|
||||
hashconfig->kern_type = KERN_TYPE_CHACHA20;
|
||||
hashconfig->dgst_size = DGST_SIZE_4_4;
|
||||
hashconfig->parse_func = chacha20_parse_hash;
|
||||
hashconfig->opti_type = OPTI_TYPE_ZERO_BYTE
|
||||
| OPTI_TYPE_USES_BITS_32
|
||||
| OPTI_TYPE_RAW_HASH;
|
||||
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;
|
||||
@ -22822,6 +22926,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
case 14700: hashconfig->esalt_size = sizeof (itunes_backup_t); break;
|
||||
case 14800: hashconfig->esalt_size = sizeof (itunes_backup_t); break;
|
||||
case 15300: hashconfig->esalt_size = sizeof (dpapimk_t); break;
|
||||
case 15400: hashconfig->esalt_size = sizeof (chacha20_t); break;
|
||||
}
|
||||
|
||||
// hook_salt_size
|
||||
@ -23032,6 +23137,8 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
|
||||
break;
|
||||
case 14900: hashconfig->pw_max = 10;
|
||||
break;
|
||||
case 15400: hashconfig->pw_max = 32;
|
||||
break;
|
||||
}
|
||||
|
||||
return 0;
|
||||
|
||||
@ -172,6 +172,7 @@ static const char *USAGE_BIG[] =
|
||||
" 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",
|
||||
" 15400 | Chacha20 | Raw Cipher, Known-Plaintext attack",
|
||||
" 400 | phpass | Generic KDF",
|
||||
" 8900 | scrypt | Generic KDF",
|
||||
" 11900 | PBKDF2-HMAC-MD5 | Generic KDF",
|
||||
|
||||
@ -31,6 +31,7 @@ use Crypt::Twofish;
|
||||
use Crypt::Mode::ECB;
|
||||
use Crypt::UnixCrypt_XS qw (crypt_rounds fold_password base64_to_int24 block_to_base64 int24_to_base64);
|
||||
use Crypt::Skip32;
|
||||
use Crypt::OpenSSH::ChachaPoly;
|
||||
use MIME::Base64;
|
||||
use Authen::Passphrase::NTHash;
|
||||
use Authen::Passphrase::MySQL323;
|
||||
@ -47,7 +48,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, 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, 15200, 15300, 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, 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, 15200, 15300, 15400, 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);
|
||||
@ -3137,7 +3138,7 @@ sub passthrough
|
||||
|
||||
my $tmp_hash;
|
||||
|
||||
if ($mode == 0 || $mode == 100 || $mode == 101 || $mode == 133 || $mode == 200 || $mode == 300 || $mode == 600 || $mode == 900 || $mode == 1000 || $mode == 1300 || $mode == 1400 || $mode == 1700 || $mode == 2400 || $mode == 2600 || $mode == 3500 || $mode == 4300 || $mode == 4400 || $mode == 4500 || $mode == 4600 || $mode == 4700 || $mode == 5000 || $mode == 5100 || $mode == 6000 || $mode == 6100 || $mode == 6900 || $mode == 5700 || $mode == 9900 || $mode == 10800 || $mode == 11500 || $mode == 13300 || $mode == 99999)
|
||||
if ($mode == 0 || $mode == 100 || $mode == 101 || $mode == 133 || $mode == 200 || $mode == 300 || $mode == 600 || $mode == 900 || $mode == 1000 || $mode == 1300 || $mode == 1400 || $mode == 1700 || $mode == 2400 || $mode == 2600 || $mode == 3500 || $mode == 4300 || $mode == 4400 || $mode == 4500 || $mode == 4600 || $mode == 4700 || $mode == 5000 || $mode == 5100 || $mode == 6000 || $mode == 6100 || $mode == 6900 || $mode == 5700 || $mode == 9900 || $mode == 10800 || $mode == 11500 || $mode == 13300 || $mode == 15400 || $mode == 99999)
|
||||
{
|
||||
$tmp_hash = gen_hash ($mode, $word_buf, "");
|
||||
}
|
||||
@ -3602,7 +3603,7 @@ sub single
|
||||
{
|
||||
my $mode = $modes[$j];
|
||||
|
||||
if ($mode == 0 || $mode == 100 || $mode == 101 || $mode == 133 || $mode == 200 || $mode == 300 || $mode == 600 || $mode == 900 || $mode == 1000 || $mode == 1300 || $mode == 1400 || $mode == 1700 || $mode == 2400 || $mode == 2600 || $mode == 3500 || $mode == 4300 || $mode == 4400 || $mode == 4500 || $mode == 4600 || $mode == 4700 || $mode == 5000 || $mode == 5100 || $mode == 5300 || $mode == 5400 || $mode == 6000 || $mode == 6100 || $mode == 6600 || $mode == 6900 || $mode == 5700 || $mode == 8200 || $mode == 8300 || $mode == 9900 || $mode == 10800 || $mode == 11500 || $mode == 13300 || $mode == 99999)
|
||||
if ($mode == 0 || $mode == 100 || $mode == 101 || $mode == 133 || $mode == 200 || $mode == 300 || $mode == 600 || $mode == 900 || $mode == 1000 || $mode == 1300 || $mode == 1400 || $mode == 1700 || $mode == 2400 || $mode == 2600 || $mode == 3500 || $mode == 4300 || $mode == 4400 || $mode == 4500 || $mode == 4600 || $mode == 4700 || $mode == 5000 || $mode == 5100 || $mode == 5300 || $mode == 5400 || $mode == 6000 || $mode == 6100 || $mode == 6600 || $mode == 6900 || $mode == 5700 || $mode == 8200 || $mode == 8300 || $mode == 9900 || $mode == 10800 || $mode == 11500 || $mode == 13300 || $mode == 15400 || $mode == 99999)
|
||||
{
|
||||
for (my $i = 1; $i < 32; $i++)
|
||||
{
|
||||
@ -8495,6 +8496,26 @@ END_CODE
|
||||
$cipher_len,
|
||||
unpack ("H*", $cipher));
|
||||
}
|
||||
elsif ($mode == 15400)
|
||||
{
|
||||
my $iv = "0200000000000001";
|
||||
my $counter = "0400000000000003";
|
||||
my $plaintext = "ABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789abcdefghijklmnopqrstuvwxyz0a2b4c6d8e";
|
||||
my $eight_byte_iv = pack("H*", $iv);
|
||||
my $eight_byte_counter = pack("H*", $counter);
|
||||
my $offset = int(rand(63));
|
||||
my $pad_len = 32 - length $word_buf;
|
||||
my $key = $word_buf . "\0" x $pad_len;
|
||||
my $cipher = Crypt::OpenSSH::ChachaPoly->new($key);
|
||||
|
||||
$cipher->ivsetup($eight_byte_iv, $eight_byte_counter);
|
||||
|
||||
my $enc = $cipher->encrypt($plaintext);
|
||||
my $enc_offset = substr($enc, $offset, 8);
|
||||
$hash_buf = $enc_offset;
|
||||
|
||||
$tmp_hash = sprintf ("\$Chacha20\$\*%s\*%d\*%s\*%s\*%s", $counter, $offset, $iv, unpack("H*", substr($plaintext, $offset, 8)), unpack("H*", $enc_offset));
|
||||
}
|
||||
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 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 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 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 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 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 12700 12800 12900 13000 13100 13200 13300 13400 13500 13600 13800 14000 14100 14400 14600 14700 14800 14900 15000 15100 15200 15300 15400 99999"
|
||||
|
||||
#ATTACK_MODES="0 1 3 6 7"
|
||||
ATTACK_MODES="0 1 3 7"
|
||||
|
||||
Loading…
Reference in New Issue
Block a user