+++ /dev/null
-/**
- * Author......: Jens Steube <jens.steube@gmail.com>
- * License.....: MIT
- */
-
-#define _SIPHASH_
-#define _SCALAR_
-
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
-
-#ifdef VLIW1
-#define VECT_SIZE1
-#endif
-
-#ifdef VLIW2
-#define VECT_SIZE1
-#endif
-
-#define DGST_R0 0
-#define DGST_R1 1
-#define DGST_R2 2
-#define DGST_R3 3
-
-#include "include/kernel_functions.c"
-#include "types_nv.c"
-#include "common_nv.c"
-
-#ifdef VECT_SIZE1
-#define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
-#define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
-#define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
-#endif
-
-#ifdef VECT_SIZE4
-#define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
-#define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
-#endif
-
-#define SIPROUND(v0,v1,v2,v3) \
- (v0) += (v1); \
- (v1) = rotl64 ((v1), 13); \
- (v1) ^= (v0); \
- (v0) = rotl64 ((v0), 32); \
- (v2) += (v3); \
- (v3) = rotl64 ((v3), 16); \
- (v3) ^= (v2); \
- (v0) += (v3); \
- (v3) = rotl64 ((v3), 21); \
- (v3) ^= (v0); \
- (v2) += (v1); \
- (v1) = rotl64 ((v1), 17); \
- (v1) ^= (v2); \
- (v2) = rotl64 ((v2), 32);
-
-__device__ __constant__ u32x c_bfs[1024];
-
-__device__ static void m10100m (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
-{
- /**
- * modifier
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
- const u32 lid = threadIdx.x;
-
- /**
- * base
- */
-
- u64 v0p = SIPHASHM_0;
- u64 v1p = SIPHASHM_1;
- u64 v2p = SIPHASHM_2;
- u64 v3p = SIPHASHM_3;
-
- v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
- v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
- v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
- v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
-
- u64 *w_ptr = (u64 *) w;
-
- w_ptr[pw_len / 8] |= (u64) pw_len << 56;
-
- /**
- * loop
- */
-
- const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
-
- u32x w0l = w[0];
-
- for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
- {
- const u32x w0r = c_bfs[il_pos];
-
- const u32x w0 = w0l | w0r;
-
- u64x v0 = v0p;
- u64x v1 = v1p;
- u64x v2 = v2p;
- u64x v3 = v3p;
-
- u64x m = hl32_to_64 (w[1], w0);
-
- v3 ^= m;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- v0 ^= m;
-
- int i;
- int j;
-
- for (i = 8, j = 2; i <= pw_len; i += 8, j += 2)
- {
- m = hl32_to_64 (w[j + 1], w[j + 0]);
-
- v3 ^= m;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- v0 ^= m;
- }
-
- v2 ^= 0xff;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- const u64x v = v0 ^ v1 ^ v2 ^ v3;
-
- const u32x a = l32_from_64 (v);
- const u32x b = h32_from_64 (v);
-
- const u32x r0 = a;
- const u32x r1 = b;
- const u32x r2 = 0;
- const u32x r3 = 0;
-
- #include VECT_COMPARE_M
- }
-}
-
-__device__ static void m10100s (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x * words_buf_r, void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
-{
- /**
- * modifier
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
- const u32 lid = threadIdx.x;
-
- /**
- * base
- */
-
- u64 v0p = SIPHASHM_0;
- u64 v1p = SIPHASHM_1;
- u64 v2p = SIPHASHM_2;
- u64 v3p = SIPHASHM_3;
-
- v0p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
- v1p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
- v2p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[1], salt_bufs[salt_pos].salt_buf[0]);
- v3p ^= hl32_to_64 (salt_bufs[salt_pos].salt_buf[3], salt_bufs[salt_pos].salt_buf[2]);
-
- u64 *w_ptr = (u64 *) w;
-
- w_ptr[pw_len / 8] |= (u64) pw_len << 56;
-
- /**
- * 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
- */
-
- const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
-
- u32x w0l = w[0];
-
- for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
- {
- const u32x w0r = c_bfs[il_pos];
-
- const u32x w0 = w0l | w0r;
-
- u64x v0 = v0p;
- u64x v1 = v1p;
- u64x v2 = v2p;
- u64x v3 = v3p;
-
- u64x m = hl32_to_64 (w[1], w0);
-
- v3 ^= m;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- v0 ^= m;
-
- int i;
- int j;
-
- for (i = 8, j = 2; i <= pw_len; i += 8, j += 2)
- {
- m = hl32_to_64 (w[j + 1], w[j + 0]);
-
- v3 ^= m;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- v0 ^= m;
- }
-
- v2 ^= 0xff;
-
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
- SIPROUND (v0, v1, v2, v3);
-
- const u64x v = v0 ^ v1 ^ v2 ^ v3;
-
- const u32x a = l32_from_64 (v);
- const u32x b = h32_from_64 (v);
-
- const u32x r0 = a;
- const u32x r1 = b;
- const u32x r2 = 0;
- const u32x r3 = 0;
-
- #include VECT_COMPARE_S
- }
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = 0;
- w[ 5] = 0;
- w[ 6] = 0;
- w[ 7] = 0;
- w[ 8] = 0;
- w[ 9] = 0;
- w[10] = 0;
- w[11] = 0;
- w[12] = 0;
- w[13] = 0;
- w[14] = pws[gid].i[14];
- w[15] = 0;
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = pws[gid].i[ 4];
- w[ 5] = pws[gid].i[ 5];
- w[ 6] = pws[gid].i[ 6];
- w[ 7] = pws[gid].i[ 7];
- w[ 8] = 0;
- w[ 9] = 0;
- w[10] = 0;
- w[11] = 0;
- w[12] = 0;
- w[13] = 0;
- w[14] = pws[gid].i[14];
- w[15] = 0;
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = pws[gid].i[ 4];
- w[ 5] = pws[gid].i[ 5];
- w[ 6] = pws[gid].i[ 6];
- w[ 7] = pws[gid].i[ 7];
- w[ 8] = pws[gid].i[ 8];
- w[ 9] = pws[gid].i[ 9];
- w[10] = pws[gid].i[10];
- w[11] = pws[gid].i[11];
- w[12] = pws[gid].i[12];
- w[13] = pws[gid].i[13];
- w[14] = pws[gid].i[14];
- w[15] = pws[gid].i[15];
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = 0;
- w[ 5] = 0;
- w[ 6] = 0;
- w[ 7] = 0;
- w[ 8] = 0;
- w[ 9] = 0;
- w[10] = 0;
- w[11] = 0;
- w[12] = 0;
- w[13] = 0;
- w[14] = pws[gid].i[14];
- w[15] = 0;
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = pws[gid].i[ 4];
- w[ 5] = pws[gid].i[ 5];
- w[ 6] = pws[gid].i[ 6];
- w[ 7] = pws[gid].i[ 7];
- w[ 8] = 0;
- w[ 9] = 0;
- w[10] = 0;
- w[11] = 0;
- w[12] = 0;
- w[13] = 0;
- w[14] = pws[gid].i[14];
- w[15] = 0;
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
-
-extern "C" __global__ void __launch_bounds__ (256, 1) m10100_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32 w[16];
-
- w[ 0] = pws[gid].i[ 0];
- w[ 1] = pws[gid].i[ 1];
- w[ 2] = pws[gid].i[ 2];
- w[ 3] = pws[gid].i[ 3];
- w[ 4] = pws[gid].i[ 4];
- w[ 5] = pws[gid].i[ 5];
- w[ 6] = pws[gid].i[ 6];
- w[ 7] = pws[gid].i[ 7];
- w[ 8] = pws[gid].i[ 8];
- w[ 9] = pws[gid].i[ 9];
- w[10] = pws[gid].i[10];
- w[11] = pws[gid].i[11];
- w[12] = pws[gid].i[12];
- w[13] = pws[gid].i[13];
- w[14] = pws[gid].i[14];
- w[15] = pws[gid].i[15];
-
- const u32 pw_len = pws[gid].pw_len;
-
- /**
- * main
- */
-
- m10100s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
-}
projects / hashcat.git / blobdiff