nv/m07700_a3.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__ bf_t c_bfs[1024];
 240
 241 __device__ static void m07700m (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 242 {
 243   /**
 244    * modifier
 245    */
 246
 247   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 248   const u32 lid = threadIdx.x;
 249
 250   w0[0] = sapb_trans (w0[0]);
 251   w0[1] = sapb_trans (w0[1]);
 252
 253   /**
 254    * salt
 255    */
 256
 257   u32 salt_buf0[3];
 258
 259   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
 260   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
 261   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
 262
 263   salt_buf0[0] = sapb_trans (salt_buf0[0]);
 264   salt_buf0[1] = sapb_trans (salt_buf0[1]);
 265   salt_buf0[2] = sapb_trans (salt_buf0[2]);
 266
 267   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 268
 269   u32 s0[4];
 270
 271   s0[0] = salt_buf0[0];
 272   s0[1] = salt_buf0[1];
 273   s0[2] = salt_buf0[2];
 274   s0[3] = 0;
 275
 276   u32 s1[4];
 277
 278   s1[0] = 0;
 279   s1[1] = 0;
 280   s1[2] = 0;
 281   s1[3] = 0;
 282
 283   u32 s2[4];
 284
 285   s2[0] = 0;
 286   s2[1] = 0;
 287   s2[2] = 0;
 288   s2[3] = 0;
 289
 290   u32 s3[4];
 291
 292   s3[0] = 0;
 293   s3[1] = 0;
 294   s3[2] = 0;
 295   s3[3] = 0;
 296
 297   switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
 298
 299   const u32 pw_salt_len = pw_len + salt_len;
 300
 301   /**
 302    * loop
 303    */
 304
 305   u32x w0l = w0[0];
 306
 307   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 308   {
 309     const u32 w0r = sapb_trans (c_bfs[il_pos].i);
 310
 311     w0[0] = w0l | w0r;
 312
 313     u32x t[16];
 314
 315     t[ 0] = s0[0] | w0[0];
 316     t[ 1] = s0[1] | w0[1];
 317     t[ 2] = s0[2];
 318     t[ 3] = s0[3];
 319     t[ 4] = s1[0];
 320     t[ 5] = 0;
 321     t[ 6] = 0;
 322     t[ 7] = 0;
 323     t[ 8] = 0;
 324     t[ 9] = 0;
 325     t[10] = 0;
 326     t[11] = 0;
 327     t[12] = 0;
 328     t[13] = 0;
 329     t[14] = pw_salt_len * 8;
 330     t[15] = 0;
 331
 332     append_0x80_4 (&t[0], &t[4], &t[8], &t[12], pw_salt_len);
 333
 334     /**
 335      * md5
 336      */
 337
 338     u32x a = MD5M_A;
 339     u32x b = MD5M_B;
 340     u32x c = MD5M_C;
 341     u32x d = MD5M_D;
 342
 343     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 344     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 345     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 346     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 347     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 348     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 349     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 350     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 351     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 352     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 353     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 354     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 355     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 356     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 357     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 358     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 359
 360     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 361     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 362     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 363     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 364     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 365     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 366     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 367     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 368     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 369     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 370     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 371     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 372     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 373     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 374     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 375     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 376
 377     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 378     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 379     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 380     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 381     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 382     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 383     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 384     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 385     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 386     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 387     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 388     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 389     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 390     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 391     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 392     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 393
 394     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 395     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 396     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 397     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 398     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 399     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 400     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 401     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 402     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 403     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 404     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 405     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 406     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 407     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 408     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 409     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 410
 411     a += MD5M_A;
 412     b += MD5M_B;
 413     c += MD5M_C;
 414     d += MD5M_D;
 415
 416     const u32x sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
 417
 418     t[14] = sum20 * 8;
 419
 420     a = MD5M_A;
 421     b = MD5M_B;
 422     c = MD5M_C;
 423     d = MD5M_D;
 424
 425     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 426     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 427     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 428     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 429     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 430     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 431     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 432     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 433     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 434     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 435     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 436     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 437     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 438     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 439     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 440     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 441
 442     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 443     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 444     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 445     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 446     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 447     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 448     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 449     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 450     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 451     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 452     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 453     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 454     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 455     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 456     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 457     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 458
 459     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 460     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 461     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 462     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 463     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 464     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 465     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 466     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 467     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 468     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 469     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 470     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 471     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 472     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 473     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 474     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 475
 476     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 477     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 478     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 479     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 480     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 481     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 482     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 483     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 484     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 485     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 486     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 487     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 488     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 489     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 490     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 491     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 492
 493     a += MD5M_A;
 494     b += MD5M_B;
 495     c += MD5M_C;
 496     d += MD5M_D;
 497
 498     a ^= c;
 499     b ^= d;
 500
 501     const u32x r0 = a;
 502     const u32x r1 = b;
 503     const u32x r2 = 0;
 504     const u32x r3 = 0;
 505
 506     #include VECT_COMPARE_M
 507   }
 508 }
 509
 510 __device__ static void m07700s (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 511 {
 512   /**
 513    * modifier
 514    */
 515
 516   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 517   const u32 lid = threadIdx.x;
 518
 519   w0[0] = sapb_trans (w0[0]);
 520   w0[1] = sapb_trans (w0[1]);
 521
 522   /**
 523    * salt
 524    */
 525
 526   u32 salt_buf0[3];
 527
 528   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
 529   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
 530   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
 531
 532   salt_buf0[0] = sapb_trans (salt_buf0[0]);
 533   salt_buf0[1] = sapb_trans (salt_buf0[1]);
 534   salt_buf0[2] = sapb_trans (salt_buf0[2]);
 535
 536   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 537
 538   u32 s0[4];
 539
 540   s0[0] = salt_buf0[0];
 541   s0[1] = salt_buf0[1];
 542   s0[2] = salt_buf0[2];
 543   s0[3] = 0;
 544
 545   u32 s1[4];
 546
 547   s1[0] = 0;
 548   s1[1] = 0;
 549   s1[2] = 0;
 550   s1[3] = 0;
 551
 552   u32 s2[4];
 553
 554   s2[0] = 0;
 555   s2[1] = 0;
 556   s2[2] = 0;
 557   s2[3] = 0;
 558
 559   u32 s3[4];
 560
 561   s3[0] = 0;
 562   s3[1] = 0;
 563   s3[2] = 0;
 564   s3[3] = 0;
 565
 566   switch_buffer_by_offset (s0, s1, s2, s3, pw_len);
 567
 568   const u32 pw_salt_len = pw_len + salt_len;
 569
 570   /**
 571    * digest
 572    */
 573
 574   const u32 search[4] =
 575   {
 576     digests_buf[digests_offset].digest_buf[DGST_R0],
 577     digests_buf[digests_offset].digest_buf[DGST_R1],
 578     digests_buf[digests_offset].digest_buf[DGST_R2],
 579     digests_buf[digests_offset].digest_buf[DGST_R3]
 580   };
 581
 582   /**
 583    * loop
 584    */
 585
 586   u32x w0l = w0[0];
 587
 588   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 589   {
 590     const u32 w0r = sapb_trans (c_bfs[il_pos].i);
 591
 592     w0[0] = w0l | w0r;
 593
 594     u32x t[16];
 595
 596     t[ 0] = s0[0] | w0[0];
 597     t[ 1] = s0[1] | w0[1];
 598     t[ 2] = s0[2];
 599     t[ 3] = s0[3];
 600     t[ 4] = s1[0];
 601     t[ 5] = 0;
 602     t[ 6] = 0;
 603     t[ 7] = 0;
 604     t[ 8] = 0;
 605     t[ 9] = 0;
 606     t[10] = 0;
 607     t[11] = 0;
 608     t[12] = 0;
 609     t[13] = 0;
 610     t[14] = pw_salt_len * 8;
 611     t[15] = 0;
 612
 613     append_0x80_4 (&t[0], &t[4], &t[8], &t[12], pw_salt_len);
 614
 615     /**
 616      * md5
 617      */
 618
 619     u32x a = MD5M_A;
 620     u32x b = MD5M_B;
 621     u32x c = MD5M_C;
 622     u32x d = MD5M_D;
 623
 624     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 625     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 626     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 627     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 628     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 629     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 630     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 631     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 632     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 633     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 634     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 635     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 636     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 637     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 638     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 639     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 640
 641     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 642     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 643     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 644     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 645     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 646     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 647     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 648     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 649     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 650     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 651     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 652     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 653     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 654     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 655     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 656     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 657
 658     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 659     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 660     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 661     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 662     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 663     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 664     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 665     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 666     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 667     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 668     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 669     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 670     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 671     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 672     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 673     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 674
 675     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 676     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 677     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 678     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 679     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 680     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 681     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 682     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 683     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 684     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 685     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 686     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 687     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 688     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 689     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 690     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 691
 692     a += MD5M_A;
 693     b += MD5M_B;
 694     c += MD5M_C;
 695     d += MD5M_D;
 696
 697     const u32x sum20 = walld0rf_magic (w0, pw_len, salt_buf0, salt_len, a, b, c, d, t);
 698
 699     t[14] = sum20 * 8;
 700
 701     a = MD5M_A;
 702     b = MD5M_B;
 703     c = MD5M_C;
 704     d = MD5M_D;
 705
 706     MD5_STEP (MD5_Fo, a, b, c, d, t[ 0], MD5C00, MD5S00);
 707     MD5_STEP (MD5_Fo, d, a, b, c, t[ 1], MD5C01, MD5S01);
 708     MD5_STEP (MD5_Fo, c, d, a, b, t[ 2], MD5C02, MD5S02);
 709     MD5_STEP (MD5_Fo, b, c, d, a, t[ 3], MD5C03, MD5S03);
 710     MD5_STEP (MD5_Fo, a, b, c, d, t[ 4], MD5C04, MD5S00);
 711     MD5_STEP (MD5_Fo, d, a, b, c, t[ 5], MD5C05, MD5S01);
 712     MD5_STEP (MD5_Fo, c, d, a, b, t[ 6], MD5C06, MD5S02);
 713     MD5_STEP (MD5_Fo, b, c, d, a, t[ 7], MD5C07, MD5S03);
 714     MD5_STEP (MD5_Fo, a, b, c, d, t[ 8], MD5C08, MD5S00);
 715     MD5_STEP (MD5_Fo, d, a, b, c, t[ 9], MD5C09, MD5S01);
 716     MD5_STEP (MD5_Fo, c, d, a, b, t[10], MD5C0a, MD5S02);
 717     MD5_STEP (MD5_Fo, b, c, d, a, t[11], MD5C0b, MD5S03);
 718     MD5_STEP (MD5_Fo, a, b, c, d, t[12], MD5C0c, MD5S00);
 719     MD5_STEP (MD5_Fo, d, a, b, c, t[13], MD5C0d, MD5S01);
 720     MD5_STEP (MD5_Fo, c, d, a, b, t[14], MD5C0e, MD5S02);
 721     MD5_STEP (MD5_Fo, b, c, d, a, t[15], MD5C0f, MD5S03);
 722
 723     MD5_STEP (MD5_Go, a, b, c, d, t[ 1], MD5C10, MD5S10);
 724     MD5_STEP (MD5_Go, d, a, b, c, t[ 6], MD5C11, MD5S11);
 725     MD5_STEP (MD5_Go, c, d, a, b, t[11], MD5C12, MD5S12);
 726     MD5_STEP (MD5_Go, b, c, d, a, t[ 0], MD5C13, MD5S13);
 727     MD5_STEP (MD5_Go, a, b, c, d, t[ 5], MD5C14, MD5S10);
 728     MD5_STEP (MD5_Go, d, a, b, c, t[10], MD5C15, MD5S11);
 729     MD5_STEP (MD5_Go, c, d, a, b, t[15], MD5C16, MD5S12);
 730     MD5_STEP (MD5_Go, b, c, d, a, t[ 4], MD5C17, MD5S13);
 731     MD5_STEP (MD5_Go, a, b, c, d, t[ 9], MD5C18, MD5S10);
 732     MD5_STEP (MD5_Go, d, a, b, c, t[14], MD5C19, MD5S11);
 733     MD5_STEP (MD5_Go, c, d, a, b, t[ 3], MD5C1a, MD5S12);
 734     MD5_STEP (MD5_Go, b, c, d, a, t[ 8], MD5C1b, MD5S13);
 735     MD5_STEP (MD5_Go, a, b, c, d, t[13], MD5C1c, MD5S10);
 736     MD5_STEP (MD5_Go, d, a, b, c, t[ 2], MD5C1d, MD5S11);
 737     MD5_STEP (MD5_Go, c, d, a, b, t[ 7], MD5C1e, MD5S12);
 738     MD5_STEP (MD5_Go, b, c, d, a, t[12], MD5C1f, MD5S13);
 739
 740     MD5_STEP (MD5_H , a, b, c, d, t[ 5], MD5C20, MD5S20);
 741     MD5_STEP (MD5_H , d, a, b, c, t[ 8], MD5C21, MD5S21);
 742     MD5_STEP (MD5_H , c, d, a, b, t[11], MD5C22, MD5S22);
 743     MD5_STEP (MD5_H , b, c, d, a, t[14], MD5C23, MD5S23);
 744     MD5_STEP (MD5_H , a, b, c, d, t[ 1], MD5C24, MD5S20);
 745     MD5_STEP (MD5_H , d, a, b, c, t[ 4], MD5C25, MD5S21);
 746     MD5_STEP (MD5_H , c, d, a, b, t[ 7], MD5C26, MD5S22);
 747     MD5_STEP (MD5_H , b, c, d, a, t[10], MD5C27, MD5S23);
 748     MD5_STEP (MD5_H , a, b, c, d, t[13], MD5C28, MD5S20);
 749     MD5_STEP (MD5_H , d, a, b, c, t[ 0], MD5C29, MD5S21);
 750     MD5_STEP (MD5_H , c, d, a, b, t[ 3], MD5C2a, MD5S22);
 751     MD5_STEP (MD5_H , b, c, d, a, t[ 6], MD5C2b, MD5S23);
 752     MD5_STEP (MD5_H , a, b, c, d, t[ 9], MD5C2c, MD5S20);
 753     MD5_STEP (MD5_H , d, a, b, c, t[12], MD5C2d, MD5S21);
 754     MD5_STEP (MD5_H , c, d, a, b, t[15], MD5C2e, MD5S22);
 755     MD5_STEP (MD5_H , b, c, d, a, t[ 2], MD5C2f, MD5S23);
 756
 757     MD5_STEP (MD5_I , a, b, c, d, t[ 0], MD5C30, MD5S30);
 758     MD5_STEP (MD5_I , d, a, b, c, t[ 7], MD5C31, MD5S31);
 759     MD5_STEP (MD5_I , c, d, a, b, t[14], MD5C32, MD5S32);
 760     MD5_STEP (MD5_I , b, c, d, a, t[ 5], MD5C33, MD5S33);
 761     MD5_STEP (MD5_I , a, b, c, d, t[12], MD5C34, MD5S30);
 762     MD5_STEP (MD5_I , d, a, b, c, t[ 3], MD5C35, MD5S31);
 763     MD5_STEP (MD5_I , c, d, a, b, t[10], MD5C36, MD5S32);
 764     MD5_STEP (MD5_I , b, c, d, a, t[ 1], MD5C37, MD5S33);
 765     MD5_STEP (MD5_I , a, b, c, d, t[ 8], MD5C38, MD5S30);
 766     MD5_STEP (MD5_I , d, a, b, c, t[15], MD5C39, MD5S31);
 767     MD5_STEP (MD5_I , c, d, a, b, t[ 6], MD5C3a, MD5S32);
 768     MD5_STEP (MD5_I , b, c, d, a, t[13], MD5C3b, MD5S33);
 769     MD5_STEP (MD5_I , a, b, c, d, t[ 4], MD5C3c, MD5S30);
 770     MD5_STEP (MD5_I , d, a, b, c, t[11], MD5C3d, MD5S31);
 771     MD5_STEP (MD5_I , c, d, a, b, t[ 2], MD5C3e, MD5S32);
 772     MD5_STEP (MD5_I , b, c, d, a, t[ 9], MD5C3f, MD5S33);
 773
 774     a += MD5M_A;
 775     b += MD5M_B;
 776     c += MD5M_C;
 777     d += MD5M_D;
 778
 779     a ^= c;
 780     b ^= d;
 781
 782     const u32x r0 = a;
 783     const u32x r1 = b;
 784     const u32x r2 = 0;
 785     const u32x r3 = 0;
 786
 787     #include VECT_COMPARE_S
 788   }
 789 }
 790
 791 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 792 {
 793   /**
 794    * base
 795    */
 796
 797   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 798
 799   if (gid >= gid_max) return;
 800
 801   u32x w0[4];
 802
 803   w0[0] = pws[gid].i[ 0];
 804   w0[1] = pws[gid].i[ 1];
 805   w0[2] = pws[gid].i[ 2];
 806   w0[3] = pws[gid].i[ 3];
 807
 808   u32x w1[4];
 809
 810   w1[0] = 0;
 811   w1[1] = 0;
 812   w1[2] = 0;
 813   w1[3] = 0;
 814
 815   u32x w2[4];
 816
 817   w2[0] = 0;
 818   w2[1] = 0;
 819   w2[2] = 0;
 820   w2[3] = 0;
 821
 822   u32x w3[4];
 823
 824   w3[0] = 0;
 825   w3[1] = 0;
 826   w3[2] = 0;
 827   w3[3] = 0;
 828
 829   const u32 pw_len = pws[gid].pw_len;
 830
 831   /**
 832    * main
 833    */
 834
 835   m07700m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 836 }
 837
 838 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 839 {
 840   /**
 841    * base
 842    */
 843
 844   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 845
 846   if (gid >= gid_max) return;
 847
 848   u32x w0[4];
 849
 850   w0[0] = pws[gid].i[ 0];
 851   w0[1] = pws[gid].i[ 1];
 852   w0[2] = pws[gid].i[ 2];
 853   w0[3] = pws[gid].i[ 3];
 854
 855   u32x w1[4];
 856
 857   w1[0] = pws[gid].i[ 4];
 858   w1[1] = pws[gid].i[ 5];
 859   w1[2] = pws[gid].i[ 6];
 860   w1[3] = pws[gid].i[ 7];
 861
 862   u32x w2[4];
 863
 864   w2[0] = 0;
 865   w2[1] = 0;
 866   w2[2] = 0;
 867   w2[3] = 0;
 868
 869   u32x w3[4];
 870
 871   w3[0] = 0;
 872   w3[1] = 0;
 873   w3[2] = 0;
 874   w3[3] = 0;
 875
 876   const u32 pw_len = pws[gid].pw_len;
 877
 878   /**
 879    * main
 880    */
 881
 882   m07700m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 883 }
 884
 885 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 886 {
 887 }
 888
 889 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 890 {
 891   /**
 892    * base
 893    */
 894
 895   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 896
 897   if (gid >= gid_max) return;
 898
 899   u32x w0[4];
 900
 901   w0[0] = pws[gid].i[ 0];
 902   w0[1] = pws[gid].i[ 1];
 903   w0[2] = pws[gid].i[ 2];
 904   w0[3] = pws[gid].i[ 3];
 905
 906   u32x w1[4];
 907
 908   w1[0] = 0;
 909   w1[1] = 0;
 910   w1[2] = 0;
 911   w1[3] = 0;
 912
 913   u32x w2[4];
 914
 915   w2[0] = 0;
 916   w2[1] = 0;
 917   w2[2] = 0;
 918   w2[3] = 0;
 919
 920   u32x w3[4];
 921
 922   w3[0] = 0;
 923   w3[1] = 0;
 924   w3[2] = 0;
 925   w3[3] = 0;
 926
 927   const u32 pw_len = pws[gid].pw_len;
 928
 929   /**
 930    * main
 931    */
 932
 933   m07700s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 934 }
 935
 936 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 937 {
 938   /**
 939    * base
 940    */
 941
 942   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 943
 944   if (gid >= gid_max) return;
 945
 946   u32x w0[4];
 947
 948   w0[0] = pws[gid].i[ 0];
 949   w0[1] = pws[gid].i[ 1];
 950   w0[2] = pws[gid].i[ 2];
 951   w0[3] = pws[gid].i[ 3];
 952
 953   u32x w1[4];
 954
 955   w1[0] = pws[gid].i[ 4];
 956   w1[1] = pws[gid].i[ 5];
 957   w1[2] = pws[gid].i[ 6];
 958   w1[3] = pws[gid].i[ 7];
 959
 960   u32x w2[4];
 961
 962   w2[0] = 0;
 963   w2[1] = 0;
 964   w2[2] = 0;
 965   w2[3] = 0;
 966
 967   u32x w3[4];
 968
 969   w3[0] = 0;
 970   w3[1] = 0;
 971   w3[2] = 0;
 972   w3[3] = 0;
 973
 974   const u32 pw_len = pws[gid].pw_len;
 975
 976   /**
 977    * main
 978    */
 979
 980   m07700s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 981 }
 982
 983 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 984 {
 985 }