+++ /dev/null
-/**
- * Author......: Jens Steube <jens.steube@gmail.com>
- * License.....: MIT
- */
-
-#define _MD5_
-
-#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"
-#include "include/rp_gpu.h"
-#include "rp_nv.c"
-
-#ifdef VECT_SIZE1
-#define VECT_COMPARE_S "check_single_vect1_comp4.c"
-#define VECT_COMPARE_M "check_multi_vect1_comp4.c"
-#endif
-
-#ifdef VECT_SIZE2
-#define VECT_COMPARE_S "check_single_vect2_comp4.c"
-#define VECT_COMPARE_M "check_multi_vect2_comp4.c"
-#endif
-
-__device__ __constant__ u32 padding[8] =
-{
- 0x5e4ebf28,
- 0x418a754e,
- 0x564e0064,
- 0x0801faff,
- 0xb6002e2e,
- 0x803e68d0,
- 0xfea90c2f,
- 0x7a695364
-};
-
-typedef struct
-{
- u8 S[256];
-
- u32 wtf_its_faster;
-
-} RC4_KEY;
-
-__device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
-{
- u8 tmp;
-
- tmp = rc4_key->S[i];
- rc4_key->S[i] = rc4_key->S[j];
- rc4_key->S[j] = tmp;
-}
-
-__device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
-{
- u32 v = 0x03020100;
- u32 a = 0x04040404;
-
- u32 *ptr = (u32 *) rc4_key->S;
-
- #pragma unroll 64
- for (u32 i = 0; i < 64; i++)
- {
- *ptr++ = v; v += a;
- }
-
- const u32 d0 = data[0] >> 0;
- const u32 d1 = data[0] >> 8;
- const u32 d2 = data[0] >> 16;
- const u32 d3 = data[0] >> 24;
- const u32 d4 = data[1] >> 0;
-
- u32 i = 0;
- u32 j = 0;
-
- #pragma unroll 52
- for (i = 0; i < 255; i += 5)
- {
- j += rc4_key->S[i + 0]; j += d0; swap (rc4_key, i + 0, j);
- j += rc4_key->S[i + 1]; j += d1; swap (rc4_key, i + 1, j);
- j += rc4_key->S[i + 2]; j += d2; swap (rc4_key, i + 2, j);
- j += rc4_key->S[i + 3]; j += d3; swap (rc4_key, i + 3, j);
- j += rc4_key->S[i + 4]; j += d4; swap (rc4_key, i + 4, j);
- }
-
- j += rc4_key->S[i + 0]; j += d0; swap (rc4_key, i + 0, j);
-}
-
-__device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
-{
- for (u32 k = 0; k < 4; k++)
- {
- u32 xor4 = 0;
-
- u8 idx;
-
- i += 1;
- j += rc4_key->S[i];
-
- swap (rc4_key, i, j);
-
- idx = rc4_key->S[i] + rc4_key->S[j];
-
- xor4 |= rc4_key->S[idx] << 0;
-
- i += 1;
- j += rc4_key->S[i];
-
- swap (rc4_key, i, j);
-
- idx = rc4_key->S[i] + rc4_key->S[j];
-
- xor4 |= rc4_key->S[idx] << 8;
-
- i += 1;
- j += rc4_key->S[i];
-
- swap (rc4_key, i, j);
-
- idx = rc4_key->S[i] + rc4_key->S[j];
-
- xor4 |= rc4_key->S[idx] << 16;
-
- i += 1;
- j += rc4_key->S[i];
-
- swap (rc4_key, i, j);
-
- idx = rc4_key->S[i] + rc4_key->S[j];
-
- xor4 |= rc4_key->S[idx] << 24;
-
- out[k] = in[k] ^ xor4;
- }
-
- return j;
-}
-
-__device__ __constant__ gpu_rule_t c_rules[1024];
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * modifier
- */
-
- const u32 lid = threadIdx.x;
-
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32x pw_buf0[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];
-
- u32x pw_buf1[4];
-
- 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;
-
- /**
- * key
- */
-
- __shared__ RC4_KEY rc4_keys[64];
- RC4_KEY *rc4_key = &rc4_keys[lid];
-
- /**
- * loop
- */
-
- for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
- {
- u32x w0[4];
-
- w0[0] = pw_buf0[0];
- w0[1] = pw_buf0[1];
- w0[2] = pw_buf0[2];
- w0[3] = pw_buf0[3];
-
- u32x w1[4];
-
- w1[0] = pw_buf1[0];
- w1[1] = pw_buf1[1];
- w1[2] = pw_buf1[2];
- w1[3] = pw_buf1[3];
-
- u32x w2[4];
-
- w2[0] = 0;
- w2[1] = 0;
- w2[2] = 0;
- w2[3] = 0;
-
- u32x w3[4];
-
- w3[0] = 0;
- w3[1] = 0;
- w3[2] = 0;
- w3[3] = 0;
-
- const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
-
- // now the RC4 part
-
- u32x key[4];
-
- key[0] = w0[0];
- key[1] = w0[1];
- key[2] = 0;
- key[3] = 0;
-
- rc4_init_16 (rc4_key, key);
-
- u32x out[4];
-
- rc4_next_16 (rc4_key, 0, 0, padding, out);
-
- const u32x r0 = out[0];
- const u32x r1 = out[1];
- const u32x r2 = out[2];
- const u32x r3 = out[3];
-
- #include VECT_COMPARE_M
- }
-}
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
-}
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
-}
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
- /**
- * modifier
- */
-
- const u32 lid = threadIdx.x;
-
- /**
- * base
- */
-
- const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
-
- if (gid >= gid_max) return;
-
- u32x pw_buf0[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];
-
- u32x pw_buf1[4];
-
- 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;
-
- /**
- * 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]
- };
-
- /**
- * key
- */
-
- __shared__ RC4_KEY rc4_keys[64];
- RC4_KEY *rc4_key = &rc4_keys[lid];
-
- /**
- * loop
- */
-
- for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
- {
- u32x w0[4];
-
- w0[0] = pw_buf0[0];
- w0[1] = pw_buf0[1];
- w0[2] = pw_buf0[2];
- w0[3] = pw_buf0[3];
-
- u32x w1[4];
-
- w1[0] = pw_buf1[0];
- w1[1] = pw_buf1[1];
- w1[2] = pw_buf1[2];
- w1[3] = pw_buf1[3];
-
- u32x w2[4];
-
- w2[0] = 0;
- w2[1] = 0;
- w2[2] = 0;
- w2[3] = 0;
-
- u32x w3[4];
-
- w3[0] = 0;
- w3[1] = 0;
- w3[2] = 0;
- w3[3] = 0;
-
- const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
-
- // now the RC4 part
-
- u32x key[4];
-
- key[0] = w0[0];
- key[1] = w0[1];
- key[2] = 0;
- key[3] = 0;
-
- rc4_init_16 (rc4_key, key);
-
- u32x out[4];
-
- rc4_next_16 (rc4_key, 0, 0, padding, out);
-
- const u32x r0 = out[0];
- const u32x r1 = out[1];
- const u32x r2 = out[2];
- const u32x r3 = out[3];
-
- #include VECT_COMPARE_S
- }
-}
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
-}
-
-extern "C" __global__ void __launch_bounds__ (64, 1) m10410_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const 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 pdf_t *pdf_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
-{
-}
projects / hashcat.git / blobdiff