From f9a4669cc8cca7a0d1e485d9a2289856d1163dec Mon Sep 17 00:00:00 2001 From: David Llewellyn-Jones Date: Sun, 14 Aug 2016 19:18:30 +0100 Subject: [PATCH] Added mangle code to hash-mode 0 Hash mode 0 is MD5 without salt. This adds the mangle code to all three mode types (0, 1, 3) for testing purposes. --- OpenCL/m00000_a0.cl | 303 +++++++++++++++++ OpenCL/m00000_a1.cl | 439 ++++++++++++++++++++++++ OpenCL/m00000_a3.cl | 787 ++++++++++++++++++++++++++++++++++++++++++++ hashcat-help.txt | 385 ++++++++++++++++++++++ 4 files changed, 1914 insertions(+) create mode 100644 OpenCL/m00000_a0.cl create mode 100644 OpenCL/m00000_a1.cl create mode 100644 OpenCL/m00000_a3.cl create mode 100644 hashcat-help.txt diff --git a/OpenCL/m00000_a0.cl b/OpenCL/m00000_a0.cl new file mode 100644 index 0000000..17ee25f --- /dev/null +++ b/OpenCL/m00000_a0.cl @@ -0,0 +1,303 @@ +/** + * Author......: Jens Steube + * License.....: MIT + */ + +#define _MD5_ + +#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" +#include "mangle.cl" + +__kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + + /** + * base + */ + + 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; + + /** + * loop + */ + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + out_len = mangle(w0, w1, out_len); + + append_0x80_2x4_VV (w0, w1, out_len); + + w3[2] = out_len * 8; + w3[3] = 0; + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03); + + MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13); + + MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33); + + COMPARE_M_SIMD (a, d, c, b); + } +} + +__kernel void m00000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + + /** + * base + */ + + 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; + + /** + * 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 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; + + u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); + + out_len = mangle(w0, w1, out_len); + + append_0x80_2x4_VV (w0, w1, out_len); + + w3[2] = out_len * 8; + w3[3] = 0; + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03); + + MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13); + + MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30); + + if (MATCHES_NONE_VS (a, search[0])) continue; + + MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33); + + COMPARE_S_SIMD (a, d, c, b); + } +} + +__kernel void m00000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} diff --git a/OpenCL/m00000_a1.cl b/OpenCL/m00000_a1.cl new file mode 100644 index 0000000..0647a03 --- /dev/null +++ b/OpenCL/m00000_a1.cl @@ -0,0 +1,439 @@ +/** + * Author......: Jens Steube + * License.....: MIT + */ + +#define _MD5_ + +#define NEW_SIMD_CODE + +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_simd.cl" +#include "mangle.cl" + +__kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + + /** + * base + */ + + 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; + + /** + * 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); + + u32x pw_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]; + u32x w2[4]; + u32x w3[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]; + + pw_len = mangle(w0, w1, pw_len); + append_0x80_2x4_VV (w0, w1, pw_len); + + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + + //w2[0] = wordl2[0] | wordr2[0]; + //w2[1] = wordl2[1] | wordr2[1]; + //w2[2] = wordl2[2] | wordr2[2]; + //w2[3] = wordl2[3] | wordr2[3]; + //w3[0] = wordl3[0] | wordr3[0]; + //w3[1] = wordl3[1] | wordr3[1]; + + w3[2] = pw_len * 8; + w3[3] = 0; + + /** + * md5 + */ + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03); + + MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13); + + MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33); + + COMPARE_M_SIMD (a, d, c, b); + } +} + +__kernel void m00000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + + /** + * base + */ + + 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; + + /** + * 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); + + u32x pw_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]; + u32x w2[4]; + u32x w3[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]; + + pw_len = mangle(w0, w1, pw_len); + append_0x80_2x4_VV (w0, w1, pw_len); + + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + + //w2[0] = wordl2[0] | wordr2[0]; + //w2[1] = wordl2[1] | wordr2[1]; + //w2[2] = wordl2[2] | wordr2[2]; + //w2[3] = wordl2[3] | wordr2[3]; + //w3[0] = wordl3[0] | wordr3[0]; + //w3[1] = wordl3[1] | wordr3[1]; + + w3[2] = pw_len * 8; + w3[3] = 0; + + /** + * md5 + */ + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03); + + MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13); + + MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30); + + if (MATCHES_NONE_VS (a, search[0])) continue; + + MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33); + + COMPARE_S_SIMD (a, d, c, b); + } +} + +__kernel void m00000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} + +__kernel void m00000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ +} diff --git a/OpenCL/m00000_a3.cl b/OpenCL/m00000_a3.cl new file mode 100644 index 0000000..06e5d3c --- /dev/null +++ b/OpenCL/m00000_a3.cl @@ -0,0 +1,787 @@ +/** + * Author......: Jens Steube + * License.....: MIT + */ + +#define _MD5_ + +#define NEW_SIMD_CODE + +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_simd.cl" +#include "mangle.cl" + +#define MD5_STEP_REV(f,a,b,c,d,x,t,s) \ +{ \ + a -= b; \ + a = rotr32_S (a, s); \ + a -= f (b, c, d); \ + a -= x; \ + a -= t; \ +} + +#define MD5_STEP_REV1(f,a,b,c,d,x,t,s) \ +{ \ + a -= b; \ + a = rotr32_S (a, s); \ + a -= x; \ + a -= t; \ +} + +void m00000m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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) +{ + /** + * modifier + */ + + const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + + /** + * loop + */ + + u32 w0l = w[0]; + + u32x w0[4]; + u32x w1[4]; + u32x w2[4]; + u32x w3[4]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + //const u32x w0 = w0l | w0r; + + w0[0] = w0l | w0r; + w0[1] = w[1]; + w0[2] = w[2]; + w0[3] = w[3]; + w1[0] = w[4]; + w1[1] = w[5]; + w1[2] = w[6]; + w1[3] = w[7]; + w2[0] = w[8]; + w2[1] = w[9]; + w2[2] = w[10]; + w2[3] = w[11]; + w3[0] = w[12]; + w3[1] = w[13]; + w3[2] = w[14]; + w3[3] = w[15]; + + const u32 out_len = mangle(w0, w1, pw_len); + append_0x80_2x4_VV (w0, w1, out_len); + + w3[2] = out_len * 8; + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03); + MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00); + MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01); + MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02); + MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03); + + MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13); + MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10); + MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11); + MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13); + + MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23); + MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21); + MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22); + MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33); + MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30); + MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31); + MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32); + MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33); + + COMPARE_M_SIMD (a, d, c, b); + } +} + + + + + + +void m00000m_orig (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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) +{ + /** + * modifier + */ + + const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + + /** + * base + */ + + const u32 F_w0c00 = 0 + MD5C00; + const u32 F_w1c01 = w[ 1] + MD5C01; + const u32 F_w2c02 = w[ 2] + MD5C02; + const u32 F_w3c03 = w[ 3] + MD5C03; + const u32 F_w4c04 = w[ 4] + MD5C04; + const u32 F_w5c05 = w[ 5] + MD5C05; + const u32 F_w6c06 = w[ 6] + MD5C06; + const u32 F_w7c07 = w[ 7] + MD5C07; + const u32 F_w8c08 = w[ 8] + MD5C08; + const u32 F_w9c09 = w[ 9] + MD5C09; + const u32 F_wac0a = w[10] + MD5C0a; + const u32 F_wbc0b = w[11] + MD5C0b; + const u32 F_wcc0c = w[12] + MD5C0c; + const u32 F_wdc0d = w[13] + MD5C0d; + const u32 F_wec0e = w[14] + MD5C0e; + const u32 F_wfc0f = w[15] + MD5C0f; + + const u32 G_w1c10 = w[ 1] + MD5C10; + const u32 G_w6c11 = w[ 6] + MD5C11; + const u32 G_wbc12 = w[11] + MD5C12; + const u32 G_w0c13 = 0 + MD5C13; + const u32 G_w5c14 = w[ 5] + MD5C14; + const u32 G_wac15 = w[10] + MD5C15; + const u32 G_wfc16 = w[15] + MD5C16; + const u32 G_w4c17 = w[ 4] + MD5C17; + const u32 G_w9c18 = w[ 9] + MD5C18; + const u32 G_wec19 = w[14] + MD5C19; + const u32 G_w3c1a = w[ 3] + MD5C1a; + const u32 G_w8c1b = w[ 8] + MD5C1b; + const u32 G_wdc1c = w[13] + MD5C1c; + const u32 G_w2c1d = w[ 2] + MD5C1d; + const u32 G_w7c1e = w[ 7] + MD5C1e; + const u32 G_wcc1f = w[12] + MD5C1f; + + const u32 H_w5c20 = w[ 5] + MD5C20; + const u32 H_w8c21 = w[ 8] + MD5C21; + const u32 H_wbc22 = w[11] + MD5C22; + const u32 H_wec23 = w[14] + MD5C23; + const u32 H_w1c24 = w[ 1] + MD5C24; + const u32 H_w4c25 = w[ 4] + MD5C25; + const u32 H_w7c26 = w[ 7] + MD5C26; + const u32 H_wac27 = w[10] + MD5C27; + const u32 H_wdc28 = w[13] + MD5C28; + const u32 H_w0c29 = 0 + MD5C29; + const u32 H_w3c2a = w[ 3] + MD5C2a; + const u32 H_w6c2b = w[ 6] + MD5C2b; + const u32 H_w9c2c = w[ 9] + MD5C2c; + const u32 H_wcc2d = w[12] + MD5C2d; + const u32 H_wfc2e = w[15] + MD5C2e; + const u32 H_w2c2f = w[ 2] + MD5C2f; + + const u32 I_w0c30 = 0 + MD5C30; + const u32 I_w7c31 = w[ 7] + MD5C31; + const u32 I_wec32 = w[14] + MD5C32; + const u32 I_w5c33 = w[ 5] + MD5C33; + const u32 I_wcc34 = w[12] + MD5C34; + const u32 I_w3c35 = w[ 3] + MD5C35; + const u32 I_wac36 = w[10] + MD5C36; + const u32 I_w1c37 = w[ 1] + MD5C37; + const u32 I_w8c38 = w[ 8] + MD5C38; + const u32 I_wfc39 = w[15] + MD5C39; + const u32 I_w6c3a = w[ 6] + MD5C3a; + const u32 I_wdc3b = w[13] + MD5C3b; + const u32 I_w4c3c = w[ 4] + MD5C3c; + const u32 I_wbc3d = w[11] + MD5C3d; + const u32 I_w2c3e = w[ 2] + MD5C3e; + const u32 I_w9c3f = w[ 9] + MD5C3f; + + /** + * loop + */ + + u32 w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + u32x w0 = w0l | w0r; + + w0 = ('o' << 24) + ('e' << 16) + ('h' << 8) + ('t'); + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w1c01, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_w2c02, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_w3c03, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_w4c04, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w5c05, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_w6c06, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_w7c07, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_w8c08, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w9c09, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_wac0a, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_wbc0b, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_wcc0c, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_wdc0d, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_wec0e, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_wfc0f, MD5S03); + + MD5_STEP0(MD5_Go, a, b, c, d, G_w1c10, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_w6c11, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_wbc12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_w5c14, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_wac15, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_wfc16, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_w4c17, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_w9c18, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_wec19, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_w3c1a, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_w8c1b, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_wdc1c, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_w2c1d, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_w7c1e, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_wcc1f, MD5S13); + + MD5_STEP0(MD5_H , a, b, c, d, H_w5c20, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_w8c21, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_wbc22, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_wec23, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_w1c24, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_w4c25, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_w7c26, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_wac27, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_wdc28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0, H_w0c29, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_w3c2a, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_w6c2b, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_w9c2c, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_wcc2d, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_wfc2e, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_w2c2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_w7c31, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_wec32, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w5c33, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_wcc34, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_w3c35, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_wac36, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w1c37, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_w8c38, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_wfc39, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_w6c3a, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_wdc3b, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_w4c3c, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_wbc3d, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_w2c3e, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w9c3f, MD5S33); + + COMPARE_M_SIMD (a, d, c, b); + } +} + +void m00000s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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) +{ + /** + * modifier + */ + + const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + + /** + * base + */ + + const u32 F_w0c00 = 0 + MD5C00; + const u32 F_w1c01 = w[ 1] + MD5C01; + const u32 F_w2c02 = w[ 2] + MD5C02; + const u32 F_w3c03 = w[ 3] + MD5C03; + const u32 F_w4c04 = w[ 4] + MD5C04; + const u32 F_w5c05 = w[ 5] + MD5C05; + const u32 F_w6c06 = w[ 6] + MD5C06; + const u32 F_w7c07 = w[ 7] + MD5C07; + const u32 F_w8c08 = w[ 8] + MD5C08; + const u32 F_w9c09 = w[ 9] + MD5C09; + const u32 F_wac0a = w[10] + MD5C0a; + const u32 F_wbc0b = w[11] + MD5C0b; + const u32 F_wcc0c = w[12] + MD5C0c; + const u32 F_wdc0d = w[13] + MD5C0d; + const u32 F_wec0e = w[14] + MD5C0e; + const u32 F_wfc0f = w[15] + MD5C0f; + + const u32 G_w1c10 = w[ 1] + MD5C10; + const u32 G_w6c11 = w[ 6] + MD5C11; + const u32 G_wbc12 = w[11] + MD5C12; + const u32 G_w0c13 = 0 + MD5C13; + const u32 G_w5c14 = w[ 5] + MD5C14; + const u32 G_wac15 = w[10] + MD5C15; + const u32 G_wfc16 = w[15] + MD5C16; + const u32 G_w4c17 = w[ 4] + MD5C17; + const u32 G_w9c18 = w[ 9] + MD5C18; + const u32 G_wec19 = w[14] + MD5C19; + const u32 G_w3c1a = w[ 3] + MD5C1a; + const u32 G_w8c1b = w[ 8] + MD5C1b; + const u32 G_wdc1c = w[13] + MD5C1c; + const u32 G_w2c1d = w[ 2] + MD5C1d; + const u32 G_w7c1e = w[ 7] + MD5C1e; + const u32 G_wcc1f = w[12] + MD5C1f; + + const u32 H_w5c20 = w[ 5] + MD5C20; + const u32 H_w8c21 = w[ 8] + MD5C21; + const u32 H_wbc22 = w[11] + MD5C22; + const u32 H_wec23 = w[14] + MD5C23; + const u32 H_w1c24 = w[ 1] + MD5C24; + const u32 H_w4c25 = w[ 4] + MD5C25; + const u32 H_w7c26 = w[ 7] + MD5C26; + const u32 H_wac27 = w[10] + MD5C27; + const u32 H_wdc28 = w[13] + MD5C28; + const u32 H_w0c29 = 0 + MD5C29; + const u32 H_w3c2a = w[ 3] + MD5C2a; + const u32 H_w6c2b = w[ 6] + MD5C2b; + const u32 H_w9c2c = w[ 9] + MD5C2c; + const u32 H_wcc2d = w[12] + MD5C2d; + const u32 H_wfc2e = w[15] + MD5C2e; + const u32 H_w2c2f = w[ 2] + MD5C2f; + + const u32 I_w0c30 = 0 + MD5C30; + const u32 I_w7c31 = w[ 7] + MD5C31; + const u32 I_wec32 = w[14] + MD5C32; + const u32 I_w5c33 = w[ 5] + MD5C33; + const u32 I_wcc34 = w[12] + MD5C34; + const u32 I_w3c35 = w[ 3] + MD5C35; + const u32 I_wac36 = w[10] + MD5C36; + const u32 I_w1c37 = w[ 1] + MD5C37; + const u32 I_w8c38 = w[ 8] + MD5C38; + const u32 I_wfc39 = w[15] + MD5C39; + const u32 I_w6c3a = w[ 6] + MD5C3a; + const u32 I_wdc3b = w[13] + MD5C3b; + const u32 I_w4c3c = w[ 4] + MD5C3c; + const u32 I_wbc3d = w[11] + MD5C3d; + const u32 I_w2c3e = w[ 2] + MD5C3e; + const u32 I_w9c3f = w[ 9] + MD5C3f; + + /** + * 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] + }; + + /** + * reverse + */ + + u32 a_rev = digests_buf[digests_offset].digest_buf[0]; + u32 b_rev = digests_buf[digests_offset].digest_buf[1]; + u32 c_rev = digests_buf[digests_offset].digest_buf[2]; + u32 d_rev = digests_buf[digests_offset].digest_buf[3]; + + MD5_STEP_REV (MD5_I_S, b_rev, c_rev, d_rev, a_rev, w[ 9], MD5C3f, MD5S33); + MD5_STEP_REV (MD5_I_S, c_rev, d_rev, a_rev, b_rev, w[ 2], MD5C3e, MD5S32); + MD5_STEP_REV (MD5_I_S, d_rev, a_rev, b_rev, c_rev, w[11], MD5C3d, MD5S31); + MD5_STEP_REV (MD5_I_S, a_rev, b_rev, c_rev, d_rev, w[ 4], MD5C3c, MD5S30); + MD5_STEP_REV (MD5_I_S, b_rev, c_rev, d_rev, a_rev, w[13], MD5C3b, MD5S33); + MD5_STEP_REV (MD5_I_S, c_rev, d_rev, a_rev, b_rev, w[ 6], MD5C3a, MD5S32); + MD5_STEP_REV (MD5_I_S, d_rev, a_rev, b_rev, c_rev, w[15], MD5C39, MD5S31); + MD5_STEP_REV (MD5_I_S, a_rev, b_rev, c_rev, d_rev, w[ 8], MD5C38, MD5S30); + MD5_STEP_REV (MD5_I_S, b_rev, c_rev, d_rev, a_rev, w[ 1], MD5C37, MD5S33); + MD5_STEP_REV (MD5_I_S, c_rev, d_rev, a_rev, b_rev, w[10], MD5C36, MD5S32); + MD5_STEP_REV (MD5_I_S, d_rev, a_rev, b_rev, c_rev, w[ 3], MD5C35, MD5S31); + MD5_STEP_REV (MD5_I_S, a_rev, b_rev, c_rev, d_rev, w[12], MD5C34, MD5S30); + MD5_STEP_REV (MD5_I_S, b_rev, c_rev, d_rev, a_rev, w[ 5], MD5C33, MD5S33); + MD5_STEP_REV (MD5_I_S, c_rev, d_rev, a_rev, b_rev, w[14], MD5C32, MD5S32); + MD5_STEP_REV (MD5_I_S, d_rev, a_rev, b_rev, c_rev, w[ 7], MD5C31, MD5S31); + MD5_STEP_REV (MD5_I_S, a_rev, b_rev, c_rev, d_rev, 0, MD5C30, MD5S30); + + const u32 pre_cd = c_rev ^ d_rev; + + MD5_STEP_REV1(MD5_H_S, b_rev, c_rev, d_rev, a_rev, w[ 2], MD5C2f, MD5S23); + MD5_STEP_REV1(MD5_H_S, c_rev, d_rev, a_rev, b_rev, w[15], MD5C2e, MD5S22); + + /** + * loop + */ + + u32 w0l = w[0]; + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) + { + const u32x w0r = words_buf_r[il_pos / VECT_SIZE]; + + const u32x w0 = w0l | w0r; + + const u32x pre_d = d_rev; + const u32x pre_a = a_rev - w0; + const u32x pre_b = b_rev - (pre_a ^ pre_cd); + const u32x pre_c = c_rev - (pre_a ^ pre_b ^ pre_d); + + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; + + MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w1c01, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_w2c02, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_w3c03, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_w4c04, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w5c05, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_w6c06, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_w7c07, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_w8c08, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_w9c09, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_wac0a, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_wbc0b, MD5S03); + MD5_STEP0(MD5_Fo, a, b, c, d, F_wcc0c, MD5S00); + MD5_STEP0(MD5_Fo, d, a, b, c, F_wdc0d, MD5S01); + MD5_STEP0(MD5_Fo, c, d, a, b, F_wec0e, MD5S02); + MD5_STEP0(MD5_Fo, b, c, d, a, F_wfc0f, MD5S03); + + MD5_STEP0(MD5_Go, a, b, c, d, G_w1c10, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_w6c11, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_wbc12, MD5S12); + MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_w5c14, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_wac15, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_wfc16, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_w4c17, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_w9c18, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_wec19, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_w3c1a, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_w8c1b, MD5S13); + MD5_STEP0(MD5_Go, a, b, c, d, G_wdc1c, MD5S10); + MD5_STEP0(MD5_Go, d, a, b, c, G_w2c1d, MD5S11); + MD5_STEP0(MD5_Go, c, d, a, b, G_w7c1e, MD5S12); + MD5_STEP0(MD5_Go, b, c, d, a, G_wcc1f, MD5S13); + + MD5_STEP0(MD5_H , a, b, c, d, H_w5c20, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_w8c21, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_wbc22, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_wec23, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_w1c24, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_w4c25, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_w7c26, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_wac27, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_wdc28, MD5S20); + MD5_STEP (MD5_H , d, a, b, c, w0, H_w0c29, MD5S21); + MD5_STEP0(MD5_H , c, d, a, b, H_w3c2a, MD5S22); + + if (MATCHES_NONE_VV (pre_c, c)) continue; + + MD5_STEP0(MD5_H , b, c, d, a, H_w6c2b, MD5S23); + MD5_STEP0(MD5_H , a, b, c, d, H_w9c2c, MD5S20); + MD5_STEP0(MD5_H , d, a, b, c, H_wcc2d, MD5S21); + + if (MATCHES_NONE_VV (pre_d, d)) continue; + + MD5_STEP0(MD5_H , c, d, a, b, H_wfc2e, MD5S22); + MD5_STEP0(MD5_H , b, c, d, a, H_w2c2f, MD5S23); + + MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_w7c31, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_wec32, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w5c33, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_wcc34, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_w3c35, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_wac36, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w1c37, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_w8c38, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_wfc39, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_w6c3a, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_wdc3b, MD5S33); + MD5_STEP0(MD5_I , a, b, c, d, I_w4c3c, MD5S30); + MD5_STEP0(MD5_I , d, a, b, c, I_wbc3d, MD5S31); + MD5_STEP0(MD5_I , c, d, a, b, I_w2c3e, MD5S32); + MD5_STEP0(MD5_I , b, c, d, a, I_w9c3f, MD5S33); + + COMPARE_S_SIMD (a, d, c, b); + } +} + +__kernel void m00000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000m (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} + +__kernel void m00000_m08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000m (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} + +__kernel void m00000_m16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000m (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} + +__kernel void m00000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000s (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} + +__kernel void m00000_s08 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000s (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} + +__kernel void m00000_s16 (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __constant u32x * words_buf_r, __global void *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global 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 u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + 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 + */ + + m00000s (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_scryptV0_buf, d_scryptV1_buf, d_scryptV2_buf, d_scryptV3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset); +} diff --git a/hashcat-help.txt b/hashcat-help.txt new file mode 100644 index 0000000..abb9d9d --- /dev/null +++ b/hashcat-help.txt @@ -0,0 +1,385 @@ +hashcat, advanced password recovery + +Usage: hashcat [options]... hash|hashfile|hccapfile [dictionary|mask|directory]... + +- [ Options ] - + + Options Short / Long | Type | Description | Example +===============================+======+======================================================+======================= + -m, --hash-type | Num | Hash-type, see references below | -m 1000 + -a, --attack-mode | Num | Attack-mode, see references below | -a 3 + -V, --version | | Print version | + -h, --help | | Print help | + --quiet | | Suppress output | + --hex-charset | | Assume charset is given in hex | + --hex-salt | | Assume salt is given in hex | + --hex-wordlist | | Assume words in wordlist is given in hex | + --force | | Ignore warnings | + --status | | Enable automatic update of the status-screen | + --status-timer | Num | Sets seconds between status-screen update to X | --status-timer=1 + --machine-readable | | Display the status view in a machine readable format | + --loopback | | Add new plains to induct directory | + --weak-hash-threshold | Num | Threshold X when to stop checking for weak hashes | --weak=0 + --markov-hcstat | File | Specify hcstat file to use | --markov-hc=my.hcstat + --markov-disable | | Disables markov-chains, emulates classic brute-force | + --markov-classic | | Enables classic markov-chains, no per-position | + -t, --markov-threshold | Num | Threshold X when to stop accepting new markov-chains | -t 50 + --runtime | Num | Abort session after X seconds of runtime | --runtime=10 + --session | Str | Define specific session name | --session=mysession + --restore | | Restore session from --session | + --restore-disable | | Do not write restore file | + -o, --outfile | File | Define outfile for recovered hash | -o outfile.txt + --outfile-format | Num | Define outfile-format X for recovered hash | --outfile-format=7 + --outfile-autohex-disable | | Disable the use of $HEX[] in output plains | + --outfile-check-timer | Num | Sets seconds between outfile checks to X | --outfile-check=30 + -p, --separator | Char | Separator char for hashlists and outfile | -p : + --stdout | | Do not crack a hash, instead print candidates only | + --show | | Compare hashlist with potfile; Show cracked hashes | + --left | | Compare hashlist with potfile; Show uncracked hashes | + --username | | Enable ignoring of usernames in hashfile | + --remove | | Enable remove of hash once it is cracked | + --remove-timer | Num | Update input hash file each X seconds | --remove-timer=30 + --potfile-disable | | Do not write potfile | + --potfile-path | Dir | Specific path to potfile | --potfile-path=my.pot + --debug-mode | Num | Defines the debug mode (hybrid only by using rules) | --debug-mode=4 + --debug-file | File | Output file for debugging rules | --debug-file=good.log + --induction-dir | Dir | Specify the induction directory to use for loopback | --induction=inducts + --outfile-check-dir | Dir | Specify the outfile directory to monitor for plains | --outfile-check-dir=x + --logfile-disable | | Disable the logfile | + --truecrypt-keyfiles | File | Keyfiles used, separate with comma | --truecrypt-key=x.png + --veracrypt-keyfiles | File | Keyfiles used, separate with comma | --veracrypt-key=x.txt + --veracrypt-pim | Num | VeraCrypt personal iterations multiplier | --veracrypt-pim=1000 + -b, --benchmark | | Run benchmark | + -c, --segment-size | Num | Sets size in MB to cache from the wordfile to X | -c 32 + --bitmap-min | Num | Sets minimum bits allowed for bitmaps to X | --bitmap-min=24 + --bitmap-max | Num | Sets maximum bits allowed for bitmaps to X | --bitmap-min=24 + --cpu-affinity | Str | Locks to CPU devices, separate with comma | --cpu-affinity=1,2,3 + --opencl-platforms | Str | OpenCL platforms to use, separate with comma | --opencl-platforms=2 + -d, --opencl-devices | Str | OpenCL devices to use, separate with comma | -d 1 + -D, --opencl-device-types | Str | OpenCL device-types to use, separate with comma | -D 1 + --opencl-vector-width | Num | Manual override OpenCL vector-width to X | --opencl-vector=4 + -w, --workload-profile | Num | Enable a specific workload profile, see pool below | -w 3 + -n, --kernel-accel | Num | Manual workload tuning, set outerloop step size to X | -n 64 + -u, --kernel-loops | Num | Manual workload tuning, set innerloop step size to X | -u 256 + --nvidia-spin-damp | Num | Workaround NVidias CPU burning loop bug, in percent | --nvidia-spin-damp=50 + --gpu-temp-disable | | Disable temperature and fanspeed reads and triggers | + --gpu-temp-abort | Num | Abort if GPU temperature reaches X degrees celsius | --gpu-temp-abort=100 + --gpu-temp-retain | Num | Try to retain GPU temperature at X degrees celsius | --gpu-temp-retain=95 + --powertune-enable | | Enable power tuning, restores settings when finished | + --scrypt-tmto | Num | Manually override TMTO value for scrypt to X | --scrypt-tmto=3 + -s, --skip | Num | Skip X words from the start | -s 1000000 + -l, --limit | Num | Limit X words from the start + skipped words | -l 1000000 + --keyspace | | Show keyspace base:mod values and quit | + -j, --rule-left | Rule | Single rule applied to each word from left wordlist | -j 'c' + -k, --rule-right | Rule | Single rule applied to each word from right wordlist | -k '^-' + -r, --rules-file | File | Multiple rules applied to each word from wordlists | -r rules/best64.rule + -g, --generate-rules | Num | Generate X random rules | -g 10000 + --generate-rules-func-min | Num | Force min X funcs per rule | + --generate-rules-func-max | Num | Force max X funcs per rule | + --generate-rules-seed | Num | Force RNG seed set to X | + -1, --custom-charset1 | CS | User-defined charset ?1 | -1 ?l?d?u + -2, --custom-charset2 | CS | User-defined charset ?2 | -2 ?l?d?s + -3, --custom-charset3 | CS | User-defined charset ?3 | + -4, --custom-charset4 | CS | User-defined charset ?4 | + -i, --increment | | Enable mask increment mode | + --increment-min | Num | Start mask incrementing at X | --increment-min=4 + --increment-max | Num | Stop mask incrementing at X | --increment-max=8 + +- [ Hash modes ] - + + # | Name | Category + ======+==================================================+====================================== + 900 | MD4 | Raw Hash + 0 | MD5 | Raw Hash + 5100 | Half MD5 | Raw Hash + 100 | SHA1 | Raw Hash + 10800 | SHA-384 | Raw Hash + 1400 | SHA-256 | Raw Hash + 1700 | SHA-512 | Raw Hash + 5000 | SHA-3(Keccak) | Raw Hash + 10100 | SipHash | Raw Hash + 6000 | RipeMD160 | Raw Hash + 6100 | Whirlpool | Raw Hash + 6900 | GOST R 34.11-94 | Raw Hash + 11700 | GOST R 34.11-2012 (Streebog) 256-bit | Raw Hash + 11800 | GOST R 34.11-2012 (Streebog) 512-bit | Raw Hash + 10 | md5($pass.$salt) | Raw Hash, Salted and / or Iterated + 20 | md5($salt.$pass) | Raw Hash, Salted and / or Iterated + 30 | md5(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated + 40 | md5($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated + 3800 | md5($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated + 3710 | md5($salt.md5($pass)) | Raw Hash, Salted and / or Iterated + 2600 | md5(md5($pass)) | Raw Hash, Salted and / or Iterated + 4300 | md5(strtoupper(md5($pass))) | Raw Hash, Salted and / or Iterated + 4400 | md5(sha1($pass)) | Raw Hash, Salted and / or Iterated + 110 | sha1($pass.$salt) | Raw Hash, Salted and / or Iterated + 120 | sha1($salt.$pass) | Raw Hash, Salted and / or Iterated + 130 | sha1(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated + 140 | sha1($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated + 4500 | sha1(sha1($pass)) | Raw Hash, Salted and / or Iterated + 4700 | sha1(md5($pass)) | Raw Hash, Salted and / or Iterated + 4900 | sha1($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated + 1410 | sha256($pass.$salt) | Raw Hash, Salted and / or Iterated + 1420 | sha256($salt.$pass) | Raw Hash, Salted and / or Iterated + 1430 | sha256(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated + 1440 | sha256($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated + 1710 | sha512($pass.$salt) | Raw Hash, Salted and / or Iterated + 1720 | sha512($salt.$pass) | Raw Hash, Salted and / or Iterated + 1730 | sha512(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated + 1740 | sha512($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated + 50 | HMAC-MD5 (key = $pass) | Raw Hash, Authenticated + 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated + 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated + 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated + 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated + 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated + 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated + 1760 | HMAC-SHA512 (key = $salt) | Raw Hash, Authenticated + 400 | phpass | Generic KDF + 8900 | scrypt | Generic KDF + 11900 | PBKDF2-HMAC-MD5 | Generic KDF + 12000 | PBKDF2-HMAC-SHA1 | Generic KDF + 10900 | PBKDF2-HMAC-SHA256 | Generic KDF + 12100 | PBKDF2-HMAC-SHA512 | Generic KDF + 23 | Skype | Network protocols + 2500 | WPA/WPA2 | Network protocols + 4800 | iSCSI CHAP authentication, MD5(Chap) | Network protocols + 5300 | IKE-PSK MD5 | Network protocols + 5400 | IKE-PSK SHA1 | Network protocols + 5500 | NetNTLMv1 | Network protocols + 5500 | NetNTLMv1 + ESS | Network protocols + 5600 | NetNTLMv2 | Network protocols + 7300 | IPMI2 RAKP HMAC-SHA1 | Network protocols + 7500 | Kerberos 5 AS-REQ Pre-Auth etype 23 | Network protocols + 8300 | DNSSEC (NSEC3) | Network protocols + 10200 | Cram MD5 | Network protocols + 11100 | PostgreSQL CRAM (MD5) | Network protocols + 11200 | MySQL CRAM (SHA1) | Network protocols + 11400 | SIP digest authentication (MD5) | Network protocols + 13100 | Kerberos 5 TGS-REP etype 23 | Network protocols + 121 | SMF (Simple Machines Forum) | Forums, CMS, E-Commerce, Frameworks + 400 | phpBB3 | Forums, CMS, E-Commerce, Frameworks + 2611 | vBulletin < v3.8.5 | Forums, CMS, E-Commerce, Frameworks + 2711 | vBulletin > v3.8.5 | Forums, CMS, E-Commerce, Frameworks + 2811 | MyBB | Forums, CMS, E-Commerce, Frameworks + 2811 | IPB (Invison Power Board) | Forums, CMS, E-Commerce, Frameworks + 8400 | WBB3 (Woltlab Burning Board) | Forums, CMS, E-Commerce, Frameworks + 11 | Joomla < 2.5.18 | Forums, CMS, E-Commerce, Frameworks + 400 | Joomla > 2.5.18 | Forums, CMS, E-Commerce, Frameworks + 400 | Wordpress | Forums, CMS, E-Commerce, Frameworks + 2612 | PHPS | Forums, CMS, E-Commerce, Frameworks + 7900 | Drupal7 | Forums, CMS, E-Commerce, Frameworks + 21 | osCommerce | Forums, CMS, E-Commerce, Frameworks + 21 | xt:Commerce | Forums, CMS, E-Commerce, Frameworks + 11000 | PrestaShop | Forums, CMS, E-Commerce, Frameworks + 124 | Django (SHA-1) | Forums, CMS, E-Commerce, Frameworks + 10000 | Django (PBKDF2-SHA256) | Forums, CMS, E-Commerce, Frameworks + 3711 | Mediawiki B type | Forums, CMS, E-Commerce, Frameworks + 7600 | Redmine | Forums, CMS, E-Commerce, Frameworks + 12 | PostgreSQL | Database Server + 131 | MSSQL(2000) | Database Server + 132 | MSSQL(2005) | Database Server + 1731 | MSSQL(2012) | Database Server + 1731 | MSSQL(2014) | Database Server + 200 | MySQL323 | Database Server + 300 | MySQL4.1/MySQL5 | Database Server + 3100 | Oracle H: Type (Oracle 7+) | Database Server + 112 | Oracle S: Type (Oracle 11+) | Database Server + 12300 | Oracle T: Type (Oracle 12+) | Database Server + 8000 | Sybase ASE | Database Server + 141 | EPiServer 6.x < v4 | HTTP, SMTP, LDAP Server + 1441 | EPiServer 6.x > v4 | HTTP, SMTP, LDAP Server + 1600 | Apache $apr1$ | HTTP, SMTP, LDAP Server + 12600 | ColdFusion 10+ | HTTP, SMTP, LDAP Server + 1421 | hMailServer | HTTP, SMTP, LDAP Server + 101 | nsldap, SHA-1(Base64), Netscape LDAP SHA | HTTP, SMTP, LDAP Server + 111 | nsldaps, SSHA-1(Base64), Netscape LDAP SSHA | HTTP, SMTP, LDAP Server + 1711 | SSHA-512(Base64), LDAP {SSHA512} | HTTP, SMTP, LDAP Server + 11500 | CRC32 | Checksums + 3000 | LM | Operating-Systems + 1000 | NTLM | Operating-Systems + 1100 | Domain Cached Credentials (DCC), MS Cache | Operating-Systems + 2100 | Domain Cached Credentials 2 (DCC2), MS Cache 2 | Operating-Systems + 12800 | MS-AzureSync PBKDF2-HMAC-SHA256 | Operating-Systems + 1500 | descrypt, DES(Unix), Traditional DES | Operating-Systems + 12400 | BSDiCrypt, Extended DES | Operating-Systems + 500 | md5crypt $1$, MD5(Unix) | Operating-Systems + 3200 | bcrypt $2*$, Blowfish(Unix) | Operating-Systems + 7400 | sha256crypt $5$, SHA256(Unix) | Operating-Systems + 1800 | sha512crypt $6$, SHA512(Unix) | Operating-Systems + 122 | OSX v10.4, OSX v10.5, OSX v10.6 | Operating-Systems + 1722 | OSX v10.7 | Operating-Systems + 7100 | OSX v10.8, OSX v10.9, OSX v10.10 | Operating-Systems + 6300 | AIX {smd5} | Operating-Systems + 6700 | AIX {ssha1} | Operating-Systems + 6400 | AIX {ssha256} | Operating-Systems + 6500 | AIX {ssha512} | Operating-Systems + 2400 | Cisco-PIX | Operating-Systems + 2410 | Cisco-ASA | Operating-Systems + 500 | Cisco-IOS $1$ | Operating-Systems + 5700 | Cisco-IOS $4$ | Operating-Systems + 9200 | Cisco-IOS $8$ | Operating-Systems + 9300 | Cisco-IOS $9$ | Operating-Systems + 22 | Juniper Netscreen/SSG (ScreenOS) | Operating-Systems + 501 | Juniper IVE | Operating-Systems + 5800 | Android PIN | Operating-Systems + 13800 | Windows 8+ phone PIN/Password | Operating-Systems + 8100 | Citrix Netscaler | Operating-Systems + 8500 | RACF | Operating-Systems + 7200 | GRUB 2 | Operating-Systems + 9900 | Radmin2 | Operating-Systems + 125 | ArubaOS | Operating-Systems + 7700 | SAP CODVN B (BCODE) | Enterprise Application Software (EAS) + 7800 | SAP CODVN F/G (PASSCODE) | Enterprise Application Software (EAS) + 10300 | SAP CODVN H (PWDSALTEDHASH) iSSHA-1 | Enterprise Application Software (EAS) + 8600 | Lotus Notes/Domino 5 | Enterprise Application Software (EAS) + 8700 | Lotus Notes/Domino 6 | Enterprise Application Software (EAS) + 9100 | Lotus Notes/Domino 8 | Enterprise Application Software (EAS) + 133 | PeopleSoft | Enterprise Application Software (EAS) + 13500 | PeopleSoft Token | Enterprise Application Software (EAS) + 11600 | 7-Zip | Archives + 12500 | RAR3-hp | Archives + 13000 | RAR5 | Archives + 13200 | AxCrypt | Archives + 13300 | AxCrypt in memory SHA1 | Archives + 13600 | WinZip | Archives + 62XY | TrueCrypt | Full-Disk encryptions (FDE) + X | 1 = PBKDF2-HMAC-RipeMD160 | Full-Disk encryptions (FDE) + X | 2 = PBKDF2-HMAC-SHA512 | Full-Disk encryptions (FDE) + X | 3 = PBKDF2-HMAC-Whirlpool | Full-Disk encryptions (FDE) + X | 4 = PBKDF2-HMAC-RipeMD160 + boot-mode | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure AES | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure Serpent | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure AES | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure Serpent | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded AES-Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded Serpent-AES | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded Twofish-Serpent | Full-Disk encryptions (FDE) + Y | 3 = XTS 1536 bit all | Full-Disk encryptions (FDE) + 8800 | Android FDE < v4.3 | Full-Disk encryptions (FDE) + 12900 | Android FDE (Samsung DEK) | Full-Disk encryptions (FDE) + 12200 | eCryptfs | Full-Disk encryptions (FDE) + 137XY | VeraCrypt | Full-Disk encryptions (FDE) + X | 1 = PBKDF2-HMAC-RipeMD160 | Full-Disk encryptions (FDE) + X | 2 = PBKDF2-HMAC-SHA512 | Full-Disk encryptions (FDE) + X | 3 = PBKDF2-HMAC-Whirlpool | Full-Disk encryptions (FDE) + X | 4 = PBKDF2-HMAC-RipeMD160 + boot-mode | Full-Disk encryptions (FDE) + X | 5 = PBKDF2-HMAC-SHA256 | Full-Disk encryptions (FDE) + X | 6 = PBKDF2-HMAC-SHA256 + boot-mode | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure AES | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure Serpent | Full-Disk encryptions (FDE) + Y | 1 = XTS 512 bit pure Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure AES | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure Serpent | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit pure Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded AES-Twofish | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded Serpent-AES | Full-Disk encryptions (FDE) + Y | 2 = XTS 1024 bit cascaded Twofish-Serpent | Full-Disk encryptions (FDE) + Y | 3 = XTS 1536 bit all | Full-Disk encryptions (FDE) + 9700 | MS Office <= 2003 $0|$1, MD5 + RC4 | Documents + 9710 | MS Office <= 2003 $0|$1, MD5 + RC4, collider #1 | Documents + 9720 | MS Office <= 2003 $0|$1, MD5 + RC4, collider #2 | Documents + 9800 | MS Office <= 2003 $3|$4, SHA1 + RC4 | Documents + 9810 | MS Office <= 2003 $3|$4, SHA1 + RC4, collider #1 | Documents + 9820 | MS Office <= 2003 $3|$4, SHA1 + RC4, collider #2 | Documents + 9400 | MS Office 2007 | Documents + 9500 | MS Office 2010 | Documents + 9600 | MS Office 2013 | Documents + 10400 | PDF 1.1 - 1.3 (Acrobat 2 - 4) | Documents + 10410 | PDF 1.1 - 1.3 (Acrobat 2 - 4), collider #1 | Documents + 10420 | PDF 1.1 - 1.3 (Acrobat 2 - 4), collider #2 | Documents + 10500 | PDF 1.4 - 1.6 (Acrobat 5 - 8) | Documents + 10600 | PDF 1.7 Level 3 (Acrobat 9) | Documents + 10700 | PDF 1.7 Level 8 (Acrobat 10 - 11) | Documents + 9000 | Password Safe v2 | Password Managers + 5200 | Password Safe v3 | Password Managers + 6800 | Lastpass + Lastpass sniffed | Password Managers + 6600 | 1Password, agilekeychain | Password Managers + 8200 | 1Password, cloudkeychain | Password Managers + 11300 | Bitcoin/Litecoin wallet.dat | Password Managers + 12700 | Blockchain, My Wallet | Password Managers + 13400 | Keepass 1 (AES/Twofish) and Keepass 2 (AES) | Password Managers + +- [ Outfile Formats ] - + + # | Format + ===+======== + 1 | hash[:salt] + 2 | plain + 3 | hash[:salt]:plain + 4 | hex_plain + 5 | hash[:salt]:hex_plain + 6 | plain:hex_plain + 7 | hash[:salt]:plain:hex_plain + 8 | crackpos + 9 | hash[:salt]:crack_pos + 10 | plain:crack_pos + 11 | hash[:salt]:plain:crack_pos + 12 | hex_plain:crack_pos + 13 | hash[:salt]:hex_plain:crack_pos + 14 | plain:hex_plain:crack_pos + 15 | hash[:salt]:plain:hex_plain:crack_pos + +- [ Rule Debugging Modes ] - + + # | Format + ===+======== + 1 | Finding-Rule + 2 | Original-Word + 3 | Original-Word:Finding-Rule + 4 | Original-Word:Finding-Rule:Processed-Word + +- [ Attack Modes ] - + + # | Mode + ===+====== + 0 | Straight + 1 | Combination + 3 | Brute-force + 6 | Hybrid Wordlist + Mask + 7 | Hybrid Mask + Wordlist + +- [ Built-in Charsets ] - + + ? | Charset + ===+========= + l | abcdefghijklmnopqrstuvwxyz + u | ABCDEFGHIJKLMNOPQRSTUVWXYZ + d | 0123456789 + s | !"#$%&'()*+,-./:;<=>?@[\]^_`{|}~ + a | ?l?u?d?s + b | 0x00 - 0xff + +- [ OpenCL Device Types ] - + + # | Device Type + ===+============= + 1 | CPU + 2 | GPU + 3 | FPGA, DSP, Co-Processor + +- [ Workload Profiles ] - + + # | Performance | Runtime | Power Consumption | Desktop Impact + ===+=============+=========+===================+================= + 1 | Low | 2 ms | Low | Minimal + 2 | Default | 12 ms | Economic | Noticeable + 3 | High | 96 ms | High | Unresponsive + 4 | Nightmare | 480 ms | Insane | Headless + +- [ Basic Examples ] - + + Attack- | Hash- | + Mode | Type | Example command + ==================+=======+================================================================== + Wordlist | $P$ | hashcat -a 0 -m 400 example400.hash example.dict + Wordlist + Rules | MD5 | hashcat -a 0 -m 0 example0.hash example.dict -r rules/best64.rule + Brute-Force | MD5 | hashcat -a 3 -m 0 example0.hash ?a?a?a?a?a?a + Combinator | MD5 | hashcat -a 1 -m 0 example0.hash example.dict example.dict + +If you still have no idea what just happened try following pages: + +* https://hashcat.net/wiki/#howtos_videos_papers_articles_etc_in_the_wild +* https://hashcat.net/wiki/#frequently_asked_questions -- 2.25.1