nv/m04700_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA1_MD5_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 3
  20 #define DGST_R1 4
  21 #define DGST_R2 2
  22 #define DGST_R3 1
  23
  24 #include "include/kernel_functions.c"
  25 #undef _MD5_
  26 #include "types_nv.c"
  27 #include "common_nv.c"
  28
  29 #ifdef  VECT_SIZE1
  30 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  31 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  32 #endif
  33
  34 #ifdef  VECT_SIZE2
  35 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  36 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  37 #endif
  38
  39 #ifdef  VECT_SIZE4
  40 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  41 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  42 #endif
  43
  44 #ifdef VECT_SIZE1
  45 #define uint_to_hex_lower8_le(i) l_bin2asc[(i)]
  46 #endif
  47
  48 #ifdef VECT_SIZE2
  49 #define uint_to_hex_lower8_le(i) u32x (l_bin2asc[(i).x], l_bin2asc[(i).y])
  50 #endif
  51
  52 #ifdef VECT_SIZE4
  53 #define uint_to_hex_lower8_le(i) u32x (l_bin2asc[(i).x], l_bin2asc[(i).y], l_bin2asc[(i).z], l_bin2asc[(i).w])
  54 #endif
  55
  56 __device__ __constant__ char c_bin2asc[16] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', 'a', 'b', 'c', 'd', 'e', 'f' };
  57
  58 __device__ __shared__ short l_bin2asc[256];
  59
  60 __device__ __constant__ bf_t c_bfs[1024];
  61
  62 __device__ static void m04700m (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
  63 {
  64   /**
  65    * modifier
  66    */
  67
  68   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  69   const u32 lid = threadIdx.x;
  70
  71   /**
  72    * loop
  73    */
  74
  75   u32x w0l = w0[0];
  76
  77   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
  78   {
  79     const u32 w0r = c_bfs[il_pos].i;
  80
  81     w0[0] = w0l | w0r;
  82
  83     /**
  84      * md5
  85      */
  86
  87     u32x a = MD5M_A;
  88     u32x b = MD5M_B;
  89     u32x c = MD5M_C;
  90     u32x d = MD5M_D;
  91
  92     MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
  93     MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
  94     MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
  95     MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
  96     MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
  97     MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
  98     MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
  99     MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
 100     MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
 101     MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
 102     MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
 103     MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
 104     MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
 105     MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
 106     MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
 107     MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
 108
 109     MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
 110     MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
 111     MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
 112     MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
 113     MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
 114     MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
 115     MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
 116     MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
 117     MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
 118     MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
 119     MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
 120     MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
 121     MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
 122     MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
 123     MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
 124     MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
 125
 126     MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
 127     MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
 128     MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
 129     MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
 130     MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
 131     MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
 132     MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
 133     MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
 134     MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
 135     MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
 136     MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
 137     MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
 138     MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
 139     MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
 140     MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
 141     MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
 142
 143     MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
 144     MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
 145     MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
 146     MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
 147     MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
 148     MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
 149     MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
 150     MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
 151     MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
 152     MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
 153     MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
 154     MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
 155     MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
 156     MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
 157     MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
 158     MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
 159
 160     a += MD5M_A;
 161     b += MD5M_B;
 162     c += MD5M_C;
 163     d += MD5M_D;
 164
 165     /*
 166      * sha1
 167      */
 168
 169     u32x w0_t = uint_to_hex_lower8_le ((a >>  8) & 255) <<  0
 170                | uint_to_hex_lower8_le ((a >>  0) & 255) << 16;
 171     u32x w1_t = uint_to_hex_lower8_le ((a >> 24) & 255) <<  0
 172                | uint_to_hex_lower8_le ((a >> 16) & 255) << 16;
 173     u32x w2_t = uint_to_hex_lower8_le ((b >>  8) & 255) <<  0
 174                | uint_to_hex_lower8_le ((b >>  0) & 255) << 16;
 175     u32x w3_t = uint_to_hex_lower8_le ((b >> 24) & 255) <<  0
 176                | uint_to_hex_lower8_le ((b >> 16) & 255) << 16;
 177     u32x w4_t = uint_to_hex_lower8_le ((c >>  8) & 255) <<  0
 178                | uint_to_hex_lower8_le ((c >>  0) & 255) << 16;
 179     u32x w5_t = uint_to_hex_lower8_le ((c >> 24) & 255) <<  0
 180                | uint_to_hex_lower8_le ((c >> 16) & 255) << 16;
 181     u32x w6_t = uint_to_hex_lower8_le ((d >>  8) & 255) <<  0
 182                | uint_to_hex_lower8_le ((d >>  0) & 255) << 16;
 183     u32x w7_t = uint_to_hex_lower8_le ((d >> 24) & 255) <<  0
 184                | uint_to_hex_lower8_le ((d >> 16) & 255) << 16;
 185
 186     u32x w8_t = 0x80000000;
 187     u32x w9_t = 0;
 188     u32x wa_t = 0;
 189     u32x wb_t = 0;
 190     u32x wc_t = 0;
 191     u32x wd_t = 0;
 192     u32x we_t = 0;
 193     u32x wf_t = 32 * 8;
 194
 195     u32x e;
 196
 197     a = SHA1M_A;
 198     b = SHA1M_B;
 199     c = SHA1M_C;
 200     d = SHA1M_D;
 201     e = SHA1M_E;
 202
 203     #undef K
 204     #define K SHA1C00
 205
 206     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 207     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 208     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 209     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 210     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 211     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 212     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 213     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 214     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 215     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 216     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 217     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 218     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 219     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 220     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 221     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 222     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 223     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 224     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 225     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 226
 227     #undef K
 228     #define K SHA1C01
 229
 230     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 231     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 232     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 233     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 234     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 235     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 236     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 237     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 238     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 239     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 240     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 241     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 242     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 243     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 244     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 245     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 246     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 247     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 248     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 249     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 250
 251     #undef K
 252     #define K SHA1C02
 253
 254     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 255     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 256     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 257     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 258     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 259     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 260     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 261     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 262     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 263     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 264     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 265     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 266     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 267     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 268     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 269     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 270     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 271     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 272     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 273     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 274
 275     #undef K
 276     #define K SHA1C03
 277
 278     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 279     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 280     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 281     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 282     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 283     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 284     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 285     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 286     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 287     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 288     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 289     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 290     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 291     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 292     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 293     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 294     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 295     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 296     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 297     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 298
 299     const u32x r0 = d;
 300     const u32x r1 = e;
 301     const u32x r2 = c;
 302     const u32x r3 = b;
 303
 304     #include VECT_COMPARE_M
 305   }
 306 }
 307
 308 __device__ static void m04700s (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 309 {
 310   /**
 311    * modifier
 312    */
 313
 314   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 315   const u32 lid = threadIdx.x;
 316
 317   /**
 318    * digest
 319    */
 320
 321   const u32 search[4] =
 322   {
 323     digests_buf[digests_offset].digest_buf[DGST_R0],
 324     digests_buf[digests_offset].digest_buf[DGST_R1],
 325     digests_buf[digests_offset].digest_buf[DGST_R2],
 326     digests_buf[digests_offset].digest_buf[DGST_R3]
 327   };
 328
 329   /**
 330    * reverse
 331    */
 332
 333   const u32 e_rev = rotl32 (search[1], 2u);
 334
 335   /**
 336    * loop
 337    */
 338
 339   u32x w0l = w0[0];
 340
 341   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 342   {
 343     const u32 w0r = c_bfs[il_pos].i;
 344
 345     w0[0] = w0l | w0r;
 346
 347     /**
 348      * md5
 349      */
 350
 351     u32x a = MD5M_A;
 352     u32x b = MD5M_B;
 353     u32x c = MD5M_C;
 354     u32x d = MD5M_D;
 355
 356     MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
 357     MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
 358     MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
 359     MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
 360     MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
 361     MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
 362     MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
 363     MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
 364     MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
 365     MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
 366     MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
 367     MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
 368     MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
 369     MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
 370     MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
 371     MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
 372
 373     MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
 374     MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
 375     MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
 376     MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
 377     MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
 378     MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
 379     MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
 380     MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
 381     MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
 382     MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
 383     MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
 384     MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
 385     MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
 386     MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
 387     MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
 388     MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
 389
 390     MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
 391     MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
 392     MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
 393     MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
 394     MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
 395     MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
 396     MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
 397     MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
 398     MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
 399     MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
 400     MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
 401     MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
 402     MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
 403     MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
 404     MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
 405     MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
 406
 407     MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
 408     MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
 409     MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
 410     MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
 411     MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
 412     MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
 413     MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
 414     MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
 415     MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
 416     MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
 417     MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
 418     MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
 419     MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
 420     MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
 421     MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
 422     MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
 423
 424     a += MD5M_A;
 425     b += MD5M_B;
 426     c += MD5M_C;
 427     d += MD5M_D;
 428
 429     /*
 430      * sha1
 431      */
 432
 433     u32x w0_t = uint_to_hex_lower8_le ((a >>  8) & 255) <<  0
 434                | uint_to_hex_lower8_le ((a >>  0) & 255) << 16;
 435     u32x w1_t = uint_to_hex_lower8_le ((a >> 24) & 255) <<  0
 436                | uint_to_hex_lower8_le ((a >> 16) & 255) << 16;
 437     u32x w2_t = uint_to_hex_lower8_le ((b >>  8) & 255) <<  0
 438                | uint_to_hex_lower8_le ((b >>  0) & 255) << 16;
 439     u32x w3_t = uint_to_hex_lower8_le ((b >> 24) & 255) <<  0
 440                | uint_to_hex_lower8_le ((b >> 16) & 255) << 16;
 441     u32x w4_t = uint_to_hex_lower8_le ((c >>  8) & 255) <<  0
 442                | uint_to_hex_lower8_le ((c >>  0) & 255) << 16;
 443     u32x w5_t = uint_to_hex_lower8_le ((c >> 24) & 255) <<  0
 444                | uint_to_hex_lower8_le ((c >> 16) & 255) << 16;
 445     u32x w6_t = uint_to_hex_lower8_le ((d >>  8) & 255) <<  0
 446                | uint_to_hex_lower8_le ((d >>  0) & 255) << 16;
 447     u32x w7_t = uint_to_hex_lower8_le ((d >> 24) & 255) <<  0
 448                | uint_to_hex_lower8_le ((d >> 16) & 255) << 16;
 449
 450     u32x w8_t = 0x80000000;
 451     u32x w9_t = 0;
 452     u32x wa_t = 0;
 453     u32x wb_t = 0;
 454     u32x wc_t = 0;
 455     u32x wd_t = 0;
 456     u32x we_t = 0;
 457     u32x wf_t = 32 * 8;
 458
 459     u32x e;
 460
 461     a = SHA1M_A;
 462     b = SHA1M_B;
 463     c = SHA1M_C;
 464     d = SHA1M_D;
 465     e = SHA1M_E;
 466
 467     #undef K
 468     #define K SHA1C00
 469
 470     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 471     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 472     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 473     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 474     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 475     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 476     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 477     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 478     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 479     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 480     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 481     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 482     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 483     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 484     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 485     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 486     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 487     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 488     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 489     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 490
 491     #undef K
 492     #define K SHA1C01
 493
 494     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 495     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 496     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 497     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 498     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 499     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 500     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 501     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 502     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 503     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 504     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 505     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 506     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 507     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 508     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 509     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 510     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 511     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 512     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 513     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 514
 515     #undef K
 516     #define K SHA1C02
 517
 518     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 519     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 520     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 521     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 522     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 523     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 524     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 525     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 526     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 527     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 528     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 529     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 530     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 531     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 532     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 533     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 534     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 535     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 536     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 537     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 538
 539     #undef K
 540     #define K SHA1C03
 541
 542     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 543     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 544     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 545     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 546     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 547     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 548     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 549     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 550     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 551     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 552     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 553     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 554     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 555     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 556     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 557     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 558
 559     if (e != e_rev) continue;
 560
 561     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 562     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 563     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 564     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 565
 566     const u32x r0 = d;
 567     const u32x r1 = e;
 568     const u32x r2 = c;
 569     const u32x r3 = b;
 570
 571     #include VECT_COMPARE_S
 572   }
 573 }
 574
 575 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 576 {
 577   /**
 578    * base
 579    */
 580
 581   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 582
 583   /**
 584    * modifier
 585    */
 586
 587   const u32 lid = threadIdx.x;
 588
 589
 590   u32x w0[4];
 591
 592   w0[0] = pws[gid].i[ 0];
 593   w0[1] = pws[gid].i[ 1];
 594   w0[2] = pws[gid].i[ 2];
 595   w0[3] = pws[gid].i[ 3];
 596
 597   u32x w1[4];
 598
 599   w1[0] = 0;
 600   w1[1] = 0;
 601   w1[2] = 0;
 602   w1[3] = 0;
 603
 604   u32x w2[4];
 605
 606   w2[0] = 0;
 607   w2[1] = 0;
 608   w2[2] = 0;
 609   w2[3] = 0;
 610
 611   u32x w3[4];
 612
 613   w3[0] = 0;
 614   w3[1] = 0;
 615   w3[2] = pws[gid].i[14];
 616   w3[3] = 0;
 617
 618   const u32 pw_len = pws[gid].pw_len;
 619
 620   /**
 621    * bin2asc table
 622    */
 623
 624   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 625                  | c_bin2asc[(lid >> 4) & 15] << 8;
 626
 627   __syncthreads ();
 628
 629   if (gid >= gid_max) return;
 630
 631   /**
 632    * main
 633    */
 634
 635   m04700m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 636 }
 637
 638 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 639 {
 640   /**
 641    * base
 642    */
 643
 644   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 645
 646   /**
 647    * modifier
 648    */
 649
 650   const u32 lid = threadIdx.x;
 651
 652   u32x w0[4];
 653
 654   w0[0] = pws[gid].i[ 0];
 655   w0[1] = pws[gid].i[ 1];
 656   w0[2] = pws[gid].i[ 2];
 657   w0[3] = pws[gid].i[ 3];
 658
 659   u32x w1[4];
 660
 661   w1[0] = pws[gid].i[ 4];
 662   w1[1] = pws[gid].i[ 5];
 663   w1[2] = pws[gid].i[ 6];
 664   w1[3] = pws[gid].i[ 7];
 665
 666   u32x w2[4];
 667
 668   w2[0] = 0;
 669   w2[1] = 0;
 670   w2[2] = 0;
 671   w2[3] = 0;
 672
 673   u32x w3[4];
 674
 675   w3[0] = 0;
 676   w3[1] = 0;
 677   w3[2] = pws[gid].i[14];
 678   w3[3] = 0;
 679
 680   const u32 pw_len = pws[gid].pw_len;
 681
 682   /**
 683    * bin2asc table
 684    */
 685
 686   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 687                  | c_bin2asc[(lid >> 4) & 15] << 8;
 688
 689   __syncthreads ();
 690
 691   if (gid >= gid_max) return;
 692
 693   /**
 694    * main
 695    */
 696
 697   m04700m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 698 }
 699
 700 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 701 {
 702   /**
 703    * base
 704    */
 705
 706   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 707
 708   /**
 709    * modifier
 710    */
 711
 712   const u32 lid = threadIdx.x;
 713
 714   u32x w0[4];
 715
 716   w0[0] = pws[gid].i[ 0];
 717   w0[1] = pws[gid].i[ 1];
 718   w0[2] = pws[gid].i[ 2];
 719   w0[3] = pws[gid].i[ 3];
 720
 721   u32x w1[4];
 722
 723   w1[0] = pws[gid].i[ 4];
 724   w1[1] = pws[gid].i[ 5];
 725   w1[2] = pws[gid].i[ 6];
 726   w1[3] = pws[gid].i[ 7];
 727
 728   u32x w2[4];
 729
 730   w2[0] = pws[gid].i[ 8];
 731   w2[1] = pws[gid].i[ 9];
 732   w2[2] = pws[gid].i[10];
 733   w2[3] = pws[gid].i[11];
 734
 735   u32x w3[4];
 736
 737   w3[0] = pws[gid].i[12];
 738   w3[1] = pws[gid].i[13];
 739   w3[2] = pws[gid].i[14];
 740   w3[3] = pws[gid].i[15];
 741
 742   const u32 pw_len = pws[gid].pw_len;
 743
 744   /**
 745    * bin2asc table
 746    */
 747
 748   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 749                  | c_bin2asc[(lid >> 4) & 15] << 8;
 750
 751   __syncthreads ();
 752
 753   if (gid >= gid_max) return;
 754
 755   /**
 756    * main
 757    */
 758
 759   m04700m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 760 }
 761
 762 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 763 {
 764   /**
 765    * base
 766    */
 767
 768   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 769
 770   /**
 771    * modifier
 772    */
 773
 774   const u32 lid = threadIdx.x;
 775
 776   u32x w0[4];
 777
 778   w0[0] = pws[gid].i[ 0];
 779   w0[1] = pws[gid].i[ 1];
 780   w0[2] = pws[gid].i[ 2];
 781   w0[3] = pws[gid].i[ 3];
 782
 783   u32x w1[4];
 784
 785   w1[0] = 0;
 786   w1[1] = 0;
 787   w1[2] = 0;
 788   w1[3] = 0;
 789
 790   u32x w2[4];
 791
 792   w2[0] = 0;
 793   w2[1] = 0;
 794   w2[2] = 0;
 795   w2[3] = 0;
 796
 797   u32x w3[4];
 798
 799   w3[0] = 0;
 800   w3[1] = 0;
 801   w3[2] = pws[gid].i[14];
 802   w3[3] = 0;
 803
 804   const u32 pw_len = pws[gid].pw_len;
 805
 806   /**
 807    * bin2asc table
 808    */
 809
 810   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 811                  | c_bin2asc[(lid >> 4) & 15] << 8;
 812
 813   __syncthreads ();
 814
 815   if (gid >= gid_max) return;
 816
 817   /**
 818    * main
 819    */
 820
 821   m04700s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 822 }
 823
 824 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 825 {
 826   /**
 827    * base
 828    */
 829
 830   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 831
 832   /**
 833    * modifier
 834    */
 835
 836   const u32 lid = threadIdx.x;
 837
 838   u32x w0[4];
 839
 840   w0[0] = pws[gid].i[ 0];
 841   w0[1] = pws[gid].i[ 1];
 842   w0[2] = pws[gid].i[ 2];
 843   w0[3] = pws[gid].i[ 3];
 844
 845   u32x w1[4];
 846
 847   w1[0] = pws[gid].i[ 4];
 848   w1[1] = pws[gid].i[ 5];
 849   w1[2] = pws[gid].i[ 6];
 850   w1[3] = pws[gid].i[ 7];
 851
 852   u32x w2[4];
 853
 854   w2[0] = 0;
 855   w2[1] = 0;
 856   w2[2] = 0;
 857   w2[3] = 0;
 858
 859   u32x w3[4];
 860
 861   w3[0] = 0;
 862   w3[1] = 0;
 863   w3[2] = pws[gid].i[14];
 864   w3[3] = 0;
 865
 866   const u32 pw_len = pws[gid].pw_len;
 867
 868   /**
 869    * bin2asc table
 870    */
 871
 872   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 873                  | c_bin2asc[(lid >> 4) & 15] << 8;
 874
 875   __syncthreads ();
 876
 877   if (gid >= gid_max) return;
 878
 879   /**
 880    * main
 881    */
 882
 883   m04700s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 884 }
 885
 886 extern "C" __global__ void __launch_bounds__ (256, 1) m04700_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 887 {
 888   /**
 889    * base
 890    */
 891
 892   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 893
 894   /**
 895    * modifier
 896    */
 897
 898   const u32 lid = threadIdx.x;
 899
 900   u32x w0[4];
 901
 902   w0[0] = pws[gid].i[ 0];
 903   w0[1] = pws[gid].i[ 1];
 904   w0[2] = pws[gid].i[ 2];
 905   w0[3] = pws[gid].i[ 3];
 906
 907   u32x w1[4];
 908
 909   w1[0] = pws[gid].i[ 4];
 910   w1[1] = pws[gid].i[ 5];
 911   w1[2] = pws[gid].i[ 6];
 912   w1[3] = pws[gid].i[ 7];
 913
 914   u32x w2[4];
 915
 916   w2[0] = pws[gid].i[ 8];
 917   w2[1] = pws[gid].i[ 9];
 918   w2[2] = pws[gid].i[10];
 919   w2[3] = pws[gid].i[11];
 920
 921   u32x w3[4];
 922
 923   w3[0] = pws[gid].i[12];
 924   w3[1] = pws[gid].i[13];
 925   w3[2] = pws[gid].i[14];
 926   w3[3] = pws[gid].i[15];
 927
 928   const u32 pw_len = pws[gid].pw_len;
 929
 930   /**
 931    * bin2asc table
 932    */
 933
 934   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 935                  | c_bin2asc[(lid >> 4) & 15] << 8;
 936
 937   __syncthreads ();
 938
 939   if (gid >= gid_max) return;
 940
 941   /**
 942    * main
 943    */
 944
 945   m04700s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 946 }