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