mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-12 18:59:03 +00:00
224315dd62
Complete Streebog support with pure kernels that allow for passwords longer than 64 characters. Provide generic inc_hash_streebog files for future Streebog-based hash modes (HMAC, PBKDF2, VeraCrypt). Include streebog support in the test suite. For this, python module PyGOST is needed. Also add clarification to hash mode description stating that Streebog hashes are expected in big-endian byte order. There are several implementations, including PyGOST, which default to little-endian byte order, while the RFC examples are big-endian. - Add pure kernels for hash-mode 11700 (Streebog-256) - Add pure kernels for hash-mode 11800 (Streebog-512) - Tests: Add hash-modes 11700 (Streebog-256) and 11800 (Streebog-512)
111 lines
4.1 KiB
Common Lisp
111 lines
4.1 KiB
Common Lisp
/**
|
|
* 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 m11800_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
|
|
*/
|
|
|
|
streebog512_ctx_t ctx0;
|
|
|
|
streebog512_init (&ctx0);
|
|
|
|
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
|
|
|
/**
|
|
* loop
|
|
*/
|
|
|
|
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
|
{
|
|
streebog512_ctx_t ctx = ctx0;
|
|
|
|
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
|
|
|
|
streebog512_final (&ctx);
|
|
|
|
const u32 r0 = l32_from_64_S (ctx.h[0]);
|
|
const u32 r1 = h32_from_64_S (ctx.h[0]);
|
|
const u32 r2 = l32_from_64_S (ctx.h[1]);
|
|
const u32 r3 = h32_from_64_S (ctx.h[1]);
|
|
|
|
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
|
}
|
|
}
|
|
|
|
__kernel void m11800_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
|
|
*/
|
|
|
|
streebog512_ctx_t ctx0;
|
|
|
|
streebog512_init (&ctx0);
|
|
|
|
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
|
|
|
/**
|
|
* loop
|
|
*/
|
|
|
|
for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
|
|
{
|
|
streebog512_ctx_t ctx = ctx0;
|
|
|
|
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
|
|
|
|
streebog512_final (&ctx);
|
|
|
|
const u32 r0 = l32_from_64_S (ctx.h[0]);
|
|
const u32 r1 = h32_from_64_S (ctx.h[0]);
|
|
const u32 r2 = l32_from_64_S (ctx.h[1]);
|
|
const u32 r3 = h32_from_64_S (ctx.h[1]);
|
|
|
|
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
|
}
|
|
}
|