nv/m04700_a1.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__ comb_t c_combs[1024];
  61
  62 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
  63 {
  64   /**
  65    * modifier
  66    */
  67
  68   const u32 lid = threadIdx.x;
  69
  70   /**
  71    * base
  72    */
  73
  74   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  75
  76   u32x wordl0[4];
  77
  78   wordl0[0] = pws[gid].i[ 0];
  79   wordl0[1] = pws[gid].i[ 1];
  80   wordl0[2] = pws[gid].i[ 2];
  81   wordl0[3] = pws[gid].i[ 3];
  82
  83   u32x wordl1[4];
  84
  85   wordl1[0] = pws[gid].i[ 4];
  86   wordl1[1] = pws[gid].i[ 5];
  87   wordl1[2] = pws[gid].i[ 6];
  88   wordl1[3] = pws[gid].i[ 7];
  89
  90   u32x wordl2[4];
  91
  92   wordl2[0] = 0;
  93   wordl2[1] = 0;
  94   wordl2[2] = 0;
  95   wordl2[3] = 0;
  96
  97   u32x wordl3[4];
  98
  99   wordl3[0] = 0;
 100   wordl3[1] = 0;
 101   wordl3[2] = 0;
 102   wordl3[3] = 0;
 103
 104   const u32 pw_l_len = pws[gid].pw_len;
 105
 106   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 107   {
 108     append_0x80_2 (wordl0, wordl1, pw_l_len);
 109
 110     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 111   }
 112
 113   /**
 114    * bin2asc table
 115    */
 116
 117   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 118                  | c_bin2asc[(lid >> 4) & 15] << 8;
 119
 120   __syncthreads ();
 121
 122   if (gid >= gid_max) return;
 123
 124   /**
 125    * loop
 126    */
 127
 128   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 129   {
 130     const u32 pw_r_len = c_combs[il_pos].pw_len;
 131
 132     const u32 pw_len = pw_l_len + pw_r_len;
 133
 134     u32 wordr0[4];
 135
 136     wordr0[0] = c_combs[il_pos].i[0];
 137     wordr0[1] = c_combs[il_pos].i[1];
 138     wordr0[2] = c_combs[il_pos].i[2];
 139     wordr0[3] = c_combs[il_pos].i[3];
 140
 141     u32 wordr1[4];
 142
 143     wordr1[0] = c_combs[il_pos].i[4];
 144     wordr1[1] = c_combs[il_pos].i[5];
 145     wordr1[2] = c_combs[il_pos].i[6];
 146     wordr1[3] = c_combs[il_pos].i[7];
 147
 148     u32 wordr2[4];
 149
 150     wordr2[0] = 0;
 151     wordr2[1] = 0;
 152     wordr2[2] = 0;
 153     wordr2[3] = 0;
 154
 155     u32 wordr3[4];
 156
 157     wordr3[0] = 0;
 158     wordr3[1] = 0;
 159     wordr3[2] = 0;
 160     wordr3[3] = 0;
 161
 162     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 163     {
 164       append_0x80_2 (wordr0, wordr1, pw_r_len);
 165
 166       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 167     }
 168
 169     u32x w0[4];
 170
 171     w0[0] = wordl0[0] | wordr0[0];
 172     w0[1] = wordl0[1] | wordr0[1];
 173     w0[2] = wordl0[2] | wordr0[2];
 174     w0[3] = wordl0[3] | wordr0[3];
 175
 176     u32x w1[4];
 177
 178     w1[0] = wordl1[0] | wordr1[0];
 179     w1[1] = wordl1[1] | wordr1[1];
 180     w1[2] = wordl1[2] | wordr1[2];
 181     w1[3] = wordl1[3] | wordr1[3];
 182
 183     u32x w2[4];
 184
 185     w2[0] = wordl2[0] | wordr2[0];
 186     w2[1] = wordl2[1] | wordr2[1];
 187     w2[2] = wordl2[2] | wordr2[2];
 188     w2[3] = wordl2[3] | wordr2[3];
 189
 190     u32x w3[4];
 191
 192     w3[0] = wordl3[0] | wordr3[0];
 193     w3[1] = wordl3[1] | wordr3[1];
 194     w3[2] = pw_len * 8;
 195     w3[3] = 0;
 196
 197     /**
 198      * md5
 199      */
 200
 201     u32x a = MD5M_A;
 202     u32x b = MD5M_B;
 203     u32x c = MD5M_C;
 204     u32x d = MD5M_D;
 205
 206     MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
 207     MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
 208     MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
 209     MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
 210     MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
 211     MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
 212     MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
 213     MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
 214     MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
 215     MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
 216     MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
 217     MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
 218     MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
 219     MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
 220     MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
 221     MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
 222
 223     MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
 224     MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
 225     MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
 226     MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
 227     MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
 228     MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
 229     MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
 230     MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
 231     MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
 232     MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
 233     MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
 234     MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
 235     MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
 236     MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
 237     MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
 238     MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
 239
 240     MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
 241     MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
 242     MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
 243     MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
 244     MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
 245     MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
 246     MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
 247     MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
 248     MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
 249     MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
 250     MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
 251     MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
 252     MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
 253     MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
 254     MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
 255     MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
 256
 257     MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
 258     MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
 259     MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
 260     MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
 261     MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
 262     MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
 263     MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
 264     MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
 265     MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
 266     MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
 267     MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
 268     MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
 269     MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
 270     MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
 271     MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
 272     MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
 273
 274     a += MD5M_A;
 275     b += MD5M_B;
 276     c += MD5M_C;
 277     d += MD5M_D;
 278
 279     /*
 280      * sha1
 281      */
 282
 283     u32x w0_t = uint_to_hex_lower8_le ((a >>  8) & 255) <<  0
 284                | uint_to_hex_lower8_le ((a >>  0) & 255) << 16;
 285     u32x w1_t = uint_to_hex_lower8_le ((a >> 24) & 255) <<  0
 286                | uint_to_hex_lower8_le ((a >> 16) & 255) << 16;
 287     u32x w2_t = uint_to_hex_lower8_le ((b >>  8) & 255) <<  0
 288                | uint_to_hex_lower8_le ((b >>  0) & 255) << 16;
 289     u32x w3_t = uint_to_hex_lower8_le ((b >> 24) & 255) <<  0
 290                | uint_to_hex_lower8_le ((b >> 16) & 255) << 16;
 291     u32x w4_t = uint_to_hex_lower8_le ((c >>  8) & 255) <<  0
 292                | uint_to_hex_lower8_le ((c >>  0) & 255) << 16;
 293     u32x w5_t = uint_to_hex_lower8_le ((c >> 24) & 255) <<  0
 294                | uint_to_hex_lower8_le ((c >> 16) & 255) << 16;
 295     u32x w6_t = uint_to_hex_lower8_le ((d >>  8) & 255) <<  0
 296                | uint_to_hex_lower8_le ((d >>  0) & 255) << 16;
 297     u32x w7_t = uint_to_hex_lower8_le ((d >> 24) & 255) <<  0
 298                | uint_to_hex_lower8_le ((d >> 16) & 255) << 16;
 299
 300     u32x w8_t = 0x80000000;
 301     u32x w9_t = 0;
 302     u32x wa_t = 0;
 303     u32x wb_t = 0;
 304     u32x wc_t = 0;
 305     u32x wd_t = 0;
 306     u32x we_t = 0;
 307     u32x wf_t = 32 * 8;
 308
 309     u32x e;
 310
 311     a = SHA1M_A;
 312     b = SHA1M_B;
 313     c = SHA1M_C;
 314     d = SHA1M_D;
 315     e = SHA1M_E;
 316
 317     #undef K
 318     #define K SHA1C00
 319
 320     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 321     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 322     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 323     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 324     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 325     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 326     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 327     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 328     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 329     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 330     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 331     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 332     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 333     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 334     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 335     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 336     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 337     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 338     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 339     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 340
 341     #undef K
 342     #define K SHA1C01
 343
 344     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 345     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 346     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 347     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 348     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 349     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 350     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 351     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 352     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 353     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 354     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 355     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 356     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 357     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 358     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 359     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 360     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 361     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 362     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 363     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 364
 365     #undef K
 366     #define K SHA1C02
 367
 368     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 369     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 370     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 371     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 372     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 373     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 374     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 375     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 376     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 377     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 378     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 379     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 380     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 381     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 382     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 383     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 384     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 385     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 386     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 387     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 388
 389     #undef K
 390     #define K SHA1C03
 391
 392     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 393     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 394     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 395     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 396     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 397     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 398     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 399     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 400     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 401     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 402     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 403     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 404     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 405     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 406     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 407     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 408     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 409     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 410     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 411     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 412
 413     const u32x r0 = d;
 414     const u32x r1 = e;
 415     const u32x r2 = c;
 416     const u32x r3 = b;
 417
 418     #include VECT_COMPARE_M
 419   }
 420 }
 421
 422 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 423 {
 424 }
 425
 426 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 427 {
 428 }
 429
 430 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 431 {
 432   /**
 433    * modifier
 434    */
 435
 436   const u32 lid = threadIdx.x;
 437
 438   /**
 439    * base
 440    */
 441
 442   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 443
 444   u32x wordl0[4];
 445
 446   wordl0[0] = pws[gid].i[ 0];
 447   wordl0[1] = pws[gid].i[ 1];
 448   wordl0[2] = pws[gid].i[ 2];
 449   wordl0[3] = pws[gid].i[ 3];
 450
 451   u32x wordl1[4];
 452
 453   wordl1[0] = pws[gid].i[ 4];
 454   wordl1[1] = pws[gid].i[ 5];
 455   wordl1[2] = pws[gid].i[ 6];
 456   wordl1[3] = pws[gid].i[ 7];
 457
 458   u32x wordl2[4];
 459
 460   wordl2[0] = 0;
 461   wordl2[1] = 0;
 462   wordl2[2] = 0;
 463   wordl2[3] = 0;
 464
 465   u32x wordl3[4];
 466
 467   wordl3[0] = 0;
 468   wordl3[1] = 0;
 469   wordl3[2] = 0;
 470   wordl3[3] = 0;
 471
 472   const u32 pw_l_len = pws[gid].pw_len;
 473
 474   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 475   {
 476     append_0x80_2 (wordl0, wordl1, pw_l_len);
 477
 478     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 479   }
 480
 481   /**
 482    * bin2asc table
 483    */
 484
 485   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 0
 486                  | c_bin2asc[(lid >> 4) & 15] << 8;
 487
 488   __syncthreads ();
 489
 490   if (gid >= gid_max) return;
 491
 492   /**
 493    * digest
 494    */
 495
 496   const u32 search[4] =
 497   {
 498     digests_buf[digests_offset].digest_buf[DGST_R0],
 499     digests_buf[digests_offset].digest_buf[DGST_R1],
 500     digests_buf[digests_offset].digest_buf[DGST_R2],
 501     digests_buf[digests_offset].digest_buf[DGST_R3]
 502   };
 503
 504   /**
 505    * reverse
 506    */
 507
 508   const u32 e_rev = rotl32 (search[1], 2u);
 509
 510   /**
 511    * loop
 512    */
 513
 514   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 515   {
 516     const u32 pw_r_len = c_combs[il_pos].pw_len;
 517
 518     const u32 pw_len = pw_l_len + pw_r_len;
 519
 520     u32 wordr0[4];
 521
 522     wordr0[0] = c_combs[il_pos].i[0];
 523     wordr0[1] = c_combs[il_pos].i[1];
 524     wordr0[2] = c_combs[il_pos].i[2];
 525     wordr0[3] = c_combs[il_pos].i[3];
 526
 527     u32 wordr1[4];
 528
 529     wordr1[0] = c_combs[il_pos].i[4];
 530     wordr1[1] = c_combs[il_pos].i[5];
 531     wordr1[2] = c_combs[il_pos].i[6];
 532     wordr1[3] = c_combs[il_pos].i[7];
 533
 534     u32 wordr2[4];
 535
 536     wordr2[0] = 0;
 537     wordr2[1] = 0;
 538     wordr2[2] = 0;
 539     wordr2[3] = 0;
 540
 541     u32 wordr3[4];
 542
 543     wordr3[0] = 0;
 544     wordr3[1] = 0;
 545     wordr3[2] = 0;
 546     wordr3[3] = 0;
 547
 548     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 549     {
 550       append_0x80_2 (wordr0, wordr1, pw_r_len);
 551
 552       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 553     }
 554
 555     u32x w0[4];
 556
 557     w0[0] = wordl0[0] | wordr0[0];
 558     w0[1] = wordl0[1] | wordr0[1];
 559     w0[2] = wordl0[2] | wordr0[2];
 560     w0[3] = wordl0[3] | wordr0[3];
 561
 562     u32x w1[4];
 563
 564     w1[0] = wordl1[0] | wordr1[0];
 565     w1[1] = wordl1[1] | wordr1[1];
 566     w1[2] = wordl1[2] | wordr1[2];
 567     w1[3] = wordl1[3] | wordr1[3];
 568
 569     u32x w2[4];
 570
 571     w2[0] = wordl2[0] | wordr2[0];
 572     w2[1] = wordl2[1] | wordr2[1];
 573     w2[2] = wordl2[2] | wordr2[2];
 574     w2[3] = wordl2[3] | wordr2[3];
 575
 576     u32x w3[4];
 577
 578     w3[0] = wordl3[0] | wordr3[0];
 579     w3[1] = wordl3[1] | wordr3[1];
 580     w3[2] = pw_len * 8;
 581     w3[3] = 0;
 582
 583     /**
 584      * md5
 585      */
 586
 587     u32x a = MD5M_A;
 588     u32x b = MD5M_B;
 589     u32x c = MD5M_C;
 590     u32x d = MD5M_D;
 591
 592     MD5_STEP (MD5_Fo, a, b, c, d, w0[0], MD5C00, MD5S00);
 593     MD5_STEP (MD5_Fo, d, a, b, c, w0[1], MD5C01, MD5S01);
 594     MD5_STEP (MD5_Fo, c, d, a, b, w0[2], MD5C02, MD5S02);
 595     MD5_STEP (MD5_Fo, b, c, d, a, w0[3], MD5C03, MD5S03);
 596     MD5_STEP (MD5_Fo, a, b, c, d, w1[0], MD5C04, MD5S00);
 597     MD5_STEP (MD5_Fo, d, a, b, c, w1[1], MD5C05, MD5S01);
 598     MD5_STEP (MD5_Fo, c, d, a, b, w1[2], MD5C06, MD5S02);
 599     MD5_STEP (MD5_Fo, b, c, d, a, w1[3], MD5C07, MD5S03);
 600     MD5_STEP (MD5_Fo, a, b, c, d, w2[0], MD5C08, MD5S00);
 601     MD5_STEP (MD5_Fo, d, a, b, c, w2[1], MD5C09, MD5S01);
 602     MD5_STEP (MD5_Fo, c, d, a, b, w2[2], MD5C0a, MD5S02);
 603     MD5_STEP (MD5_Fo, b, c, d, a, w2[3], MD5C0b, MD5S03);
 604     MD5_STEP (MD5_Fo, a, b, c, d, w3[0], MD5C0c, MD5S00);
 605     MD5_STEP (MD5_Fo, d, a, b, c, w3[1], MD5C0d, MD5S01);
 606     MD5_STEP (MD5_Fo, c, d, a, b, w3[2], MD5C0e, MD5S02);
 607     MD5_STEP (MD5_Fo, b, c, d, a, w3[3], MD5C0f, MD5S03);
 608
 609     MD5_STEP (MD5_Go, a, b, c, d, w0[1], MD5C10, MD5S10);
 610     MD5_STEP (MD5_Go, d, a, b, c, w1[2], MD5C11, MD5S11);
 611     MD5_STEP (MD5_Go, c, d, a, b, w2[3], MD5C12, MD5S12);
 612     MD5_STEP (MD5_Go, b, c, d, a, w0[0], MD5C13, MD5S13);
 613     MD5_STEP (MD5_Go, a, b, c, d, w1[1], MD5C14, MD5S10);
 614     MD5_STEP (MD5_Go, d, a, b, c, w2[2], MD5C15, MD5S11);
 615     MD5_STEP (MD5_Go, c, d, a, b, w3[3], MD5C16, MD5S12);
 616     MD5_STEP (MD5_Go, b, c, d, a, w1[0], MD5C17, MD5S13);
 617     MD5_STEP (MD5_Go, a, b, c, d, w2[1], MD5C18, MD5S10);
 618     MD5_STEP (MD5_Go, d, a, b, c, w3[2], MD5C19, MD5S11);
 619     MD5_STEP (MD5_Go, c, d, a, b, w0[3], MD5C1a, MD5S12);
 620     MD5_STEP (MD5_Go, b, c, d, a, w2[0], MD5C1b, MD5S13);
 621     MD5_STEP (MD5_Go, a, b, c, d, w3[1], MD5C1c, MD5S10);
 622     MD5_STEP (MD5_Go, d, a, b, c, w0[2], MD5C1d, MD5S11);
 623     MD5_STEP (MD5_Go, c, d, a, b, w1[3], MD5C1e, MD5S12);
 624     MD5_STEP (MD5_Go, b, c, d, a, w3[0], MD5C1f, MD5S13);
 625
 626     MD5_STEP (MD5_H , a, b, c, d, w1[1], MD5C20, MD5S20);
 627     MD5_STEP (MD5_H , d, a, b, c, w2[0], MD5C21, MD5S21);
 628     MD5_STEP (MD5_H , c, d, a, b, w2[3], MD5C22, MD5S22);
 629     MD5_STEP (MD5_H , b, c, d, a, w3[2], MD5C23, MD5S23);
 630     MD5_STEP (MD5_H , a, b, c, d, w0[1], MD5C24, MD5S20);
 631     MD5_STEP (MD5_H , d, a, b, c, w1[0], MD5C25, MD5S21);
 632     MD5_STEP (MD5_H , c, d, a, b, w1[3], MD5C26, MD5S22);
 633     MD5_STEP (MD5_H , b, c, d, a, w2[2], MD5C27, MD5S23);
 634     MD5_STEP (MD5_H , a, b, c, d, w3[1], MD5C28, MD5S20);
 635     MD5_STEP (MD5_H , d, a, b, c, w0[0], MD5C29, MD5S21);
 636     MD5_STEP (MD5_H , c, d, a, b, w0[3], MD5C2a, MD5S22);
 637     MD5_STEP (MD5_H , b, c, d, a, w1[2], MD5C2b, MD5S23);
 638     MD5_STEP (MD5_H , a, b, c, d, w2[1], MD5C2c, MD5S20);
 639     MD5_STEP (MD5_H , d, a, b, c, w3[0], MD5C2d, MD5S21);
 640     MD5_STEP (MD5_H , c, d, a, b, w3[3], MD5C2e, MD5S22);
 641     MD5_STEP (MD5_H , b, c, d, a, w0[2], MD5C2f, MD5S23);
 642
 643     MD5_STEP (MD5_I , a, b, c, d, w0[0], MD5C30, MD5S30);
 644     MD5_STEP (MD5_I , d, a, b, c, w1[3], MD5C31, MD5S31);
 645     MD5_STEP (MD5_I , c, d, a, b, w3[2], MD5C32, MD5S32);
 646     MD5_STEP (MD5_I , b, c, d, a, w1[1], MD5C33, MD5S33);
 647     MD5_STEP (MD5_I , a, b, c, d, w3[0], MD5C34, MD5S30);
 648     MD5_STEP (MD5_I , d, a, b, c, w0[3], MD5C35, MD5S31);
 649     MD5_STEP (MD5_I , c, d, a, b, w2[2], MD5C36, MD5S32);
 650     MD5_STEP (MD5_I , b, c, d, a, w0[1], MD5C37, MD5S33);
 651     MD5_STEP (MD5_I , a, b, c, d, w2[0], MD5C38, MD5S30);
 652     MD5_STEP (MD5_I , d, a, b, c, w3[3], MD5C39, MD5S31);
 653     MD5_STEP (MD5_I , c, d, a, b, w1[2], MD5C3a, MD5S32);
 654     MD5_STEP (MD5_I , b, c, d, a, w3[1], MD5C3b, MD5S33);
 655     MD5_STEP (MD5_I , a, b, c, d, w1[0], MD5C3c, MD5S30);
 656     MD5_STEP (MD5_I , d, a, b, c, w2[3], MD5C3d, MD5S31);
 657     MD5_STEP (MD5_I , c, d, a, b, w0[2], MD5C3e, MD5S32);
 658     MD5_STEP (MD5_I , b, c, d, a, w2[1], MD5C3f, MD5S33);
 659
 660     a += MD5M_A;
 661     b += MD5M_B;
 662     c += MD5M_C;
 663     d += MD5M_D;
 664
 665     /*
 666      * sha1
 667      */
 668
 669     u32x w0_t = uint_to_hex_lower8_le ((a >>  8) & 255) <<  0
 670                | uint_to_hex_lower8_le ((a >>  0) & 255) << 16;
 671     u32x w1_t = uint_to_hex_lower8_le ((a >> 24) & 255) <<  0
 672                | uint_to_hex_lower8_le ((a >> 16) & 255) << 16;
 673     u32x w2_t = uint_to_hex_lower8_le ((b >>  8) & 255) <<  0
 674                | uint_to_hex_lower8_le ((b >>  0) & 255) << 16;
 675     u32x w3_t = uint_to_hex_lower8_le ((b >> 24) & 255) <<  0
 676                | uint_to_hex_lower8_le ((b >> 16) & 255) << 16;
 677     u32x w4_t = uint_to_hex_lower8_le ((c >>  8) & 255) <<  0
 678                | uint_to_hex_lower8_le ((c >>  0) & 255) << 16;
 679     u32x w5_t = uint_to_hex_lower8_le ((c >> 24) & 255) <<  0
 680                | uint_to_hex_lower8_le ((c >> 16) & 255) << 16;
 681     u32x w6_t = uint_to_hex_lower8_le ((d >>  8) & 255) <<  0
 682                | uint_to_hex_lower8_le ((d >>  0) & 255) << 16;
 683     u32x w7_t = uint_to_hex_lower8_le ((d >> 24) & 255) <<  0
 684                | uint_to_hex_lower8_le ((d >> 16) & 255) << 16;
 685
 686     u32x w8_t = 0x80000000;
 687     u32x w9_t = 0;
 688     u32x wa_t = 0;
 689     u32x wb_t = 0;
 690     u32x wc_t = 0;
 691     u32x wd_t = 0;
 692     u32x we_t = 0;
 693     u32x wf_t = 32 * 8;
 694
 695     u32x e;
 696
 697     a = SHA1M_A;
 698     b = SHA1M_B;
 699     c = SHA1M_C;
 700     d = SHA1M_D;
 701     e = SHA1M_E;
 702
 703     #undef K
 704     #define K SHA1C00
 705
 706     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 707     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 708     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 709     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 710     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 711     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 712     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 713     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 714     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 715     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 716     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 717     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 718     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 719     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 720     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 721     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 722     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 723     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 724     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 725     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 726
 727     #undef K
 728     #define K SHA1C01
 729
 730     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 731     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 732     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 733     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 734     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 735     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 736     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 737     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 738     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 739     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 740     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 741     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 742     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 743     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 744     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 745     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 746     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 747     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 748     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 749     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 750
 751     #undef K
 752     #define K SHA1C02
 753
 754     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 755     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 756     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 757     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 758     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 759     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 760     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 761     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 762     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 763     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 764     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 765     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 766     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 767     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 768     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 769     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 770     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 771     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 772     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 773     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 774
 775     #undef K
 776     #define K SHA1C03
 777
 778     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 779     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 780     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 781     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 782     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 783     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 784     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 785     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 786     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 787     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 788     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 789     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 790     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 791     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 792     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 793     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 794
 795     if (e != e_rev) continue;
 796
 797     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 798     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 799     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 800     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 801
 802     const u32x r0 = d;
 803     const u32x r1 = e;
 804     const u32x r2 = c;
 805     const u32x r3 = b;
 806
 807     #include VECT_COMPARE_S
 808   }
 809 }
 810
 811 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 812 {
 813 }
 814
 815 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 816 {
 817 }