nv/m00120_a3.cu

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