--- /dev/null
+/**
+ * Author......: Jens Steube <jens.steube@gmail.com>
+ * 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)
+{
+}
--- /dev/null
+/**
+ * Author......: Jens Steube <jens.steube@gmail.com>
+ * 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)
+{
+}
--- /dev/null
+/**
+ * Author......: Jens Steube <jens.steube@gmail.com>
+ * 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);
+}
--- /dev/null
+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
projects / pwdhash.git / commitdiff