nv/m07700_a0.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SAPB_
   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 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27 #include "include/rp_gpu.h"
  28 #include "rp_nv.c"
  29
  30 #ifdef  VECT_SIZE1
  31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  33 #endif
  34
  35 #define GETCHAR(a,p)  (((a)[(p) / 4] >> (((p) & 3) * 8)) & 0xff)
  36 #define PUTCHAR(a,p,c) ((a)[(p) / 4] = (((a)[(p) / 4] & ~(0xff << (((p) & 3) * 8))) | ((c) << (((p) & 3) * 8))))
  37
  38 __device__ __constant__ u32 sapb_trans_tbl[256] =
  39 {
  40   // first value hack for 0 byte as part of an optimization
  41   0x00, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  42   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  43   0x3f, 0x40, 0x41, 0x50, 0x43, 0x44, 0x45, 0x4b, 0x47, 0x48, 0x4d, 0x4e, 0x54, 0x51, 0x53, 0x46,
  44   0x35, 0x36, 0x37, 0x38, 0x39, 0x3a, 0x3b, 0x3c, 0x3d, 0x3e, 0x56, 0x55, 0x5c, 0x49, 0x5d, 0x4a,
  45   0x42, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
  46   0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x58, 0x5b, 0x59, 0xff, 0x52,
  47   0x4c, 0x01, 0x02, 0x03, 0x04, 0x05, 0x06, 0x07, 0x08, 0x09, 0x0a, 0x0b, 0x0c, 0x0d, 0x0e, 0x0f,
  48   0x10, 0x11, 0x12, 0x13, 0x14, 0x15, 0x16, 0x17, 0x18, 0x19, 0x1a, 0x57, 0x5e, 0x5a, 0x4f, 0xff,
  49   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  50   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  51   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  52   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  53   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  54   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  55   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff,
  56   0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff
  57 };
  58
  59 __device__ __constant__ u32 bcodeArray[48] =
  60 {
  61   0x14, 0x77, 0xf3, 0xd4, 0xbb, 0x71, 0x23, 0xd0, 0x03, 0xff, 0x47, 0x93, 0x55, 0xaa, 0x66, 0x91,
  62   0xf2, 0x88, 0x6b, 0x99, 0xbf, 0xcb, 0x32, 0x1a, 0x19, 0xd9, 0xa7, 0x82, 0x22, 0x49, 0xa2, 0x51,
  63   0xe2, 0xb7, 0x33, 0x71, 0x8b, 0x9f, 0x5d, 0x01, 0x44, 0x70, 0xae, 0x11, 0xef, 0x28, 0xf0, 0x0d
  64 };
  65
  66 __device__ static u32x sapb_trans (const u32x in)
  67 {
  68   u32x out = 0;
  69
  70   #ifdef VECT_SIZE1
  71   out |= (sapb_trans_tbl[(in >>  0) & 0xff]) <<  0;
  72   out |= (sapb_trans_tbl[(in >>  8) & 0xff]) <<  8;
  73   out |= (sapb_trans_tbl[(in >> 16) & 0xff]) << 16;
  74   out |= (sapb_trans_tbl[(in >> 24) & 0xff]) << 24;
  75   #endif
  76
  77   return out;
  78 }
  79
  80 __device__ static u32x walld0rf_magic (const u32x w0[4], const u32 pw_len, const u32x salt_buf0[4], const u32 salt_len, const u32x a, const u32x b, const u32x c, const u32x d, u32x t[16])
  81 {
  82   t[ 0] = 0;
  83   t[ 1] = 0;
  84   t[ 2] = 0;
  85   t[ 3] = 0;
  86   t[ 4] = 0;
  87   t[ 5] = 0;
  88   t[ 6] = 0;
  89   t[ 7] = 0;
  90   t[ 8] = 0;
  91   t[ 9] = 0;
  92   t[10] = 0;
  93   t[11] = 0;
  94   t[12] = 0;
  95   t[13] = 0;
  96   t[14] = 0;
  97   t[15] = 0;
  98
  99   u32 sum20 = ((a >> 24) & 3)
 100              + ((a >> 16) & 3)
 101              + ((a >>  8) & 3)
 102              + ((a >>  0) & 3)
 103              + ((b >>  8) & 3);
 104
 105   sum20 |= 0x20;
 106
 107   const u32 w[2] = { w0[0], w0[1] };
 108
 109   const u32 s[3] = { salt_buf0[0], salt_buf0[1], salt_buf0[2] };
 110
 111   u32 saved_key[4] = { a, b, c, d };
 112
 113   u32 i1 = 0;
 114   u32 i2 = 0;
 115   u32 i3 = 0;
 116
 117   // we can assume this because the password must be at least 3
 118   // and the username must be at least 1 so we can save the if ()
 119
 120   u32 t0 = 0;
 121
 122   if ((d >> 24) & 1)
 123   {
 124     t0 |= bcodeArray[47] <<  0;
 125     t0 |= (w[0] & 0xff)  <<  8;
 126     t0 |= (s[0] & 0xff)  << 16;
 127     t0 |= bcodeArray[ 1] << 24;
 128
 129     i1 = 1;
 130     i2 = 5;
 131     i3 = 1;
 132   }
 133   else
 134   {
 135     t0 |= (w[0] & 0xff)  <<  0;
 136     t0 |= (s[0] & 0xff)  <<  8;
 137     t0 |= bcodeArray[ 0] << 16;
 138
 139     i1 = 1;
 140     i2 = 4;
 141     i3 = 1;
 142   }
 143
 144   t[0] = t0;
 145
 146   // because the following code can increase i2 by a maximum of 5,
 147   // there is an overflow potential of 4 before it comes to the next test for i2 >= sum20
 148   // we need to truncate in that case
 149
 150   while ((i1 < pw_len) && (i3 < salt_len))
 151   {
 152     if (GETCHAR (saved_key, 15 - i1) & 1)
 153     {
 154       PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
 155
 156       i2++;
 157     }
 158
 159     PUTCHAR (t, i2, GETCHAR (w, i1));
 160
 161     i1++;
 162     i2++;
 163
 164     PUTCHAR (t, i2, GETCHAR (s, i3));
 165
 166     i2++;
 167     i3++;
 168
 169     PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
 170
 171     i2++;
 172     i2++;
 173
 174     if (i2 >= sum20)
 175     {
 176       PUTCHAR (t, sum20 + 0, 0x80);
 177       PUTCHAR (t, sum20 + 1, 0);
 178       PUTCHAR (t, sum20 + 2, 0);
 179       PUTCHAR (t, sum20 + 3, 0);
 180
 181       return sum20;
 182     }
 183   }
 184
 185   while ((i1 < pw_len) || (i3 < salt_len))
 186   {
 187     if (i1 < pw_len) // max 8
 188     {
 189       if (GETCHAR (saved_key, 15 - i1) & 1)
 190       {
 191         PUTCHAR (t, i2, bcodeArray[48 - 1 - i1]);
 192
 193         i2++;
 194       }
 195
 196       PUTCHAR (t, i2, GETCHAR (w, i1));
 197
 198       i1++;
 199       i2++;
 200     }
 201     else if (i3 < salt_len) // max 12
 202     {
 203       PUTCHAR (t, i2, GETCHAR (s, i3));
 204
 205       i2++;
 206       i3++;
 207     }
 208
 209     PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
 210
 211     i2++;
 212     i2++;
 213
 214     if (i2 >= sum20)
 215     {
 216       PUTCHAR (t, sum20 + 0, 0x80);
 217       PUTCHAR (t, sum20 + 1, 0);
 218       PUTCHAR (t, sum20 + 2, 0);
 219       PUTCHAR (t, sum20 + 3, 0);
 220
 221       return sum20;
 222     }
 223   }
 224
 225   while (i2 < sum20)
 226   {
 227     PUTCHAR (t, i2, bcodeArray[i2 - i1 - i3]);
 228
 229     i2++;
 230     i2++;
 231   }
 232
 233   PUTCHAR (t, sum20 + 0, 0x80);
 234   PUTCHAR (t, sum20 + 1, 0);
 235   PUTCHAR (t, sum20 + 2, 0);
 236   PUTCHAR (t, sum20 + 3, 0);
 237
 238   return sum20;
 239 }
 240
 241 __device__ __constant__ gpu_rule_t c_rules[1024];
 242
 243 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 244 {
 245   /**
 246    * modifier
 247    */
 248
 249   const u32 lid = threadIdx.x;
 250
 251   /**
 252    * base
 253    */
 254
 255   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 256
 257   if (gid >= gid_max) return;
 258
 259   u32x pw_buf0[4];
 260
 261   pw_buf0[0] = pws[gid].i[ 0];
 262   pw_buf0[1] = pws[gid].i[ 1];
 263   pw_buf0[2] = 0;
 264   pw_buf0[3] = 0;
 265
 266   const u32 pw_len = pws[gid].pw_len;
 267
 268   /**
 269    * salt
 270    */
 271
 272   u32 salt_buf0[4];
 273
 274   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
 275   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
 276   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
 277   salt_buf0[3] = 0;
 278
 279   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 280
 281   salt_buf0[0] = sapb_trans (salt_buf0[0]);
 282   salt_buf0[1] = sapb_trans (salt_buf0[1]);
 283   salt_buf0[2] = sapb_trans (salt_buf0[2]);
 284
 285   /**
 286    * loop
 287    */
 288
 289   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 290   {
 291     u32x w0[4];
 292
 293     w0[0] = pw_buf0[0];
 294     w0[1] = pw_buf0[1];
 295     w0[2] = 0;
 296     w0[3] = 0;
 297
 298     u32x w1[4];
 299
 300     w1[0] = 0;
 301     w1[1] = 0;
 302     w1[2] = 0;
 303     w1[3] = 0;
 304
 305     u32x w2[4];
 306
 307     w2[0] = 0;
 308     w2[1] = 0;
 309     w2[2] = 0;
 310     w2[3] = 0;
 311
 312     u32x w3[4];
 313
 314     w3[0] = 0;
 315     w3[1] = 0;
 316     w3[2] = 0;
 317     w3[3] = 0;
 318
 319     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 320
 321     if (out_len > 8) continue; // otherwise it overflows in waldorf function
 322
 323     w0[0] = sapb_trans (w0[0]);
 324     w0[1] = sapb_trans (w0[1]);
 325
 326     /**
 327      * append salt
 328      */
 329
 330     u32 s0[4];
 331
 332     s0[0] = salt_buf0[0];
 333     s0[1] = salt_buf0[1];
 334     s0[2] = salt_buf0[2];
 335     s0[3] = 0;
 336
 337     u32 s1[4];
 338
 339     s1[0] = 0;
 340     s1[1] = 0;
 341     s1[2] = 0;
 342     s1[3] = 0;
 343
 344     u32 s2[4];
 345
 346     s2[0] = 0;
 347     s2[1] = 0;
 348     s2[2] = 0;
 349     s2[3] = 0;
 350
 351     u32 s3[4];
 352
 353     s3[0] = 0;
 354     s3[1] = 0;
 355     s3[2] = 0;
 356     s3[3] = 0;
 357
 358     switch_buffer_by_offset (s0, s1, s2, s3, out_len);
 359
 360     const u32 pw_salt_len = out_len + salt_len;
 361
 362     u32x t[16];
 363
 364     t[ 0] = s0[0] | w0[0];
 365     t[ 1] = s0[1] | w0[1];
 366     t[ 2] = s0[2];
 367     t[ 3] = s0[3];
 368     t[ 4] = s1[0];
 369     t[ 5] = 0;
 370     t[ 6] = 0;
 371     t[ 7] = 0;
 372     t[ 8] = 0;
 373     t[ 9] = 0;
 374     t[10] = 0;
 375     t[11] = 0;
 376     t[12] = 0;
 377     t[13] = 0;
 378     t[14] = pw_salt_len * 8;
 379     t[15] = 0;
 380
 381     append_0x80_4 (&t[0], &t[4], &t[8], &t[12], pw_salt_len);
 382
 383     /**
 384      * md5
 385      */
 386
 387     u32x a = MD5M_A;
 388     u32x b = MD5M_B;
 389     u32x c = MD5M_C;
 390     u32x d = MD5M_D;
 391
 392     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 393     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 394     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 395     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 396     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 397     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 398     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 399     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 400     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 401     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 402     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 403     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 404     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 405     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 406     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 407     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 408
 409     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 410     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 411     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 412     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 413     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 414     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 415     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 416     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 417     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 418     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 419     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 420     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 421     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 422     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 423     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 424     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 425
 426     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 427     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 428     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 429     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 430     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 431     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 432     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 433     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 434     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 435     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 436     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 437     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 438     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 439     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 440     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 441     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 442
 443     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 444     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 445     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 446     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 447     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 448     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 449     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 450     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 451     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 452     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 453     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 454     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 455     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 456     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 457     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 458     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 459
 460     a += MD5M_A;
 461     b += MD5M_B;
 462     c += MD5M_C;
 463     d += MD5M_D;
 464
 465     const u32x sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
 466
 467     t[14] = sum20 * 8;
 468
 469     a = MD5M_A;
 470     b = MD5M_B;
 471     c = MD5M_C;
 472     d = MD5M_D;
 473
 474     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 475     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 476     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 477     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 478     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 479     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 480     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 481     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 482     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 483     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 484     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 485     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 486     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 487     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 488     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 489     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 490
 491     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 492     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 493     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 494     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 495     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 496     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 497     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 498     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 499     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 500     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 501     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 502     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 503     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 504     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 505     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 506     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 507
 508     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 509     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 510     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 511     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 512     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 513     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 514     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 515     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 516     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 517     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 518     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 519     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 520     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 521     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 522     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 523     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 524
 525     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 526     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 527     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 528     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 529     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 530     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 531     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 532     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 533     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 534     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 535     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 536     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 537     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 538     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 539     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 540     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 541
 542     a += MD5M_A;
 543     b += MD5M_B;
 544     c += MD5M_C;
 545     d += MD5M_D;
 546
 547     a ^= c;
 548     b ^= d;
 549
 550     const u32x r0 = a;
 551     const u32x r1 = b;
 552     const u32x r2 = 0;
 553     const u32x r3 = 0;
 554
 555     #include VECT_COMPARE_M
 556   }
 557 }
 558
 559 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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)
 560 {
 561 }
 562
 563 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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)
 564 {
 565 }
 566
 567 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 568 {
 569   /**
 570    * modifier
 571    */
 572
 573   const u32 lid = threadIdx.x;
 574
 575   /**
 576    * base
 577    */
 578
 579   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 580
 581   if (gid >= gid_max) return;
 582
 583   u32x pw_buf0[4];
 584
 585   pw_buf0[0] = pws[gid].i[ 0];
 586   pw_buf0[1] = pws[gid].i[ 1];
 587   pw_buf0[2] = 0;
 588   pw_buf0[3] = 0;
 589
 590   const u32 pw_len = pws[gid].pw_len;
 591
 592   /**
 593    * salt
 594    */
 595
 596   u32 salt_buf0[4];
 597
 598   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
 599   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
 600   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
 601   salt_buf0[3] = 0;
 602
 603   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 604
 605   salt_buf0[0] = sapb_trans (salt_buf0[0]);
 606   salt_buf0[1] = sapb_trans (salt_buf0[1]);
 607   salt_buf0[2] = sapb_trans (salt_buf0[2]);
 608
 609   /**
 610    * digest
 611    */
 612
 613   const u32 search[4] =
 614   {
 615     digests_buf[digests_offset].digest_buf[DGST_R0],
 616     digests_buf[digests_offset].digest_buf[DGST_R1],
 617     digests_buf[digests_offset].digest_buf[DGST_R2],
 618     digests_buf[digests_offset].digest_buf[DGST_R3]
 619   };
 620
 621   /**
 622    * loop
 623    */
 624
 625   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 626   {
 627     u32x w0[4];
 628
 629     w0[0] = pw_buf0[0];
 630     w0[1] = pw_buf0[1];
 631     w0[2] = 0;
 632     w0[3] = 0;
 633
 634     u32x w1[4];
 635
 636     w1[0] = 0;
 637     w1[1] = 0;
 638     w1[2] = 0;
 639     w1[3] = 0;
 640
 641     u32x w2[4];
 642
 643     w2[0] = 0;
 644     w2[1] = 0;
 645     w2[2] = 0;
 646     w2[3] = 0;
 647
 648     u32x w3[4];
 649
 650     w3[0] = 0;
 651     w3[1] = 0;
 652     w3[2] = 0;
 653     w3[3] = 0;
 654
 655     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 656
 657     if (out_len > 8) continue; // otherwise it overflows in waldorf function
 658
 659     w0[0] = sapb_trans (w0[0]);
 660     w0[1] = sapb_trans (w0[1]);
 661
 662     /**
 663      * append salt
 664      */
 665
 666     u32 s0[4];
 667
 668     s0[0] = salt_buf0[0];
 669     s0[1] = salt_buf0[1];
 670     s0[2] = salt_buf0[2];
 671     s0[3] = 0;
 672
 673     u32 s1[4];
 674
 675     s1[0] = 0;
 676     s1[1] = 0;
 677     s1[2] = 0;
 678     s1[3] = 0;
 679
 680     u32 s2[4];
 681
 682     s2[0] = 0;
 683     s2[1] = 0;
 684     s2[2] = 0;
 685     s2[3] = 0;
 686
 687     u32 s3[4];
 688
 689     s3[0] = 0;
 690     s3[1] = 0;
 691     s3[2] = 0;
 692     s3[3] = 0;
 693
 694     switch_buffer_by_offset (s0, s1, s2, s3, out_len);
 695
 696     const u32 pw_salt_len = out_len + salt_len;
 697
 698     u32x t[16];
 699
 700     t[ 0] = s0[0] | w0[0];
 701     t[ 1] = s0[1] | w0[1];
 702     t[ 2] = s0[2];
 703     t[ 3] = s0[3];
 704     t[ 4] = s1[0];
 705     t[ 5] = 0;
 706     t[ 6] = 0;
 707     t[ 7] = 0;
 708     t[ 8] = 0;
 709     t[ 9] = 0;
 710     t[10] = 0;
 711     t[11] = 0;
 712     t[12] = 0;
 713     t[13] = 0;
 714     t[14] = pw_salt_len * 8;
 715     t[15] = 0;
 716
 717     append_0x80_4 (&t[0], &t[4], &t[8], &t[12], pw_salt_len);
 718
 719     /**
 720      * md5
 721      */
 722
 723     u32x a = MD5M_A;
 724     u32x b = MD5M_B;
 725     u32x c = MD5M_C;
 726     u32x d = MD5M_D;
 727
 728     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 729     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 730     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 731     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 732     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 733     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 734     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 735     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 736     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 737     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 738     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 739     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 740     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 741     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 742     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 743     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 744
 745     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 746     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 747     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 748     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 749     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 750     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 751     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 752     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 753     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 754     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 755     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 756     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 757     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 758     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 759     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 760     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 761
 762     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 763     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 764     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 765     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 766     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 767     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 768     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 769     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 770     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 771     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 772     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 773     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 774     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 775     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 776     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 777     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 778
 779     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 780     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 781     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 782     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 783     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 784     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 785     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 786     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 787     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 788     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 789     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 790     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 791     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 792     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 793     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 794     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 795
 796     a += MD5M_A;
 797     b += MD5M_B;
 798     c += MD5M_C;
 799     d += MD5M_D;
 800
 801     const u32x sum20 = walld0rf_magic (w0, out_len, salt_buf0, salt_len, a, b, c, d, t);
 802
 803     t[14] = sum20 * 8;
 804
 805     a = MD5M_A;
 806     b = MD5M_B;
 807     c = MD5M_C;
 808     d = MD5M_D;
 809
 810     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 811     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 812     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 813     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 814     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 815     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 816     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 817     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 818     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 819     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 820     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 821     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 822     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 823     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 824     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 825     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 826
 827     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 828     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 829     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 830     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 831     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 832     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 833     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 834     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 835     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 836     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 837     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 838     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 839     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 840     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 841     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 842     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 843
 844     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 845     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 846     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 847     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 848     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 849     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 850     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 851     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 852     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 853     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 854     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 855     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 856     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 857     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 858     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 859     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 860
 861     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 862     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 863     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 864     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 865     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 866     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 867     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 868     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 869     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 870     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 871     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 872     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 873     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 874     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 875     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 876     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 877
 878     a += MD5M_A;
 879     b += MD5M_B;
 880     c += MD5M_C;
 881     d += MD5M_D;
 882
 883     a ^= c;
 884     b ^= d;
 885
 886     const u32x r0 = a;
 887     const u32x r1 = b;
 888     const u32x r2 = 0;
 889     const u32x r3 = 0;
 890
 891     #include VECT_COMPARE_S
 892   }
 893 }
 894
 895 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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)
 896 {
 897 }
 898
 899 extern "C" __global__ void __launch_bounds__ (256, 1) m07700_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)
 900 {
 901 }