nv/m09700_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _OLDOFFICE01_
   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 #ifdef  VECT_SIZE4
  34 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  36 #endif
  37
  38 typedef struct
  39 {
  40   u8 S[256];
  41
  42   u32 wtf_its_faster;
  43
  44 } RC4_KEY;
  45
  46 __device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
  47 {
  48   u8 tmp;
  49
  50   tmp           = rc4_key->S[i];
  51   rc4_key->S[i] = rc4_key->S[j];
  52   rc4_key->S[j] = tmp;
  53 }
  54
  55 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  56 {
  57   u32 v = 0x03020100;
  58   u32 a = 0x04040404;
  59
  60   u32 *ptr = (u32 *) rc4_key->S;
  61
  62   #pragma unroll 64
  63   for (u32 i = 0; i < 64; i++)
  64   {
  65     *ptr++ = v; v += a;
  66   }
  67
  68   u32 j = 0;
  69
  70   for (u32 i = 0; i < 16; i++)
  71   {
  72     u32 idx = i * 16;
  73
  74     u32 v;
  75
  76     v = data[0];
  77
  78     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  79     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  80     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  81     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  82
  83     v = data[1];
  84
  85     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  86     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  87     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  88     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  89
  90     v = data[2];
  91
  92     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  93     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  94     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  95     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  96
  97     v = data[3];
  98
  99     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
 100     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
 101     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 102     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 103   }
 104 }
 105
 106 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 107 {
 108   for (u32 k = 0; k < 4; k++)
 109   {
 110     u32 xor4 = 0;
 111
 112     u8 idx;
 113
 114     i += 1;
 115     j += rc4_key->S[i];
 116
 117     swap (rc4_key, i, j);
 118
 119     idx = rc4_key->S[i] + rc4_key->S[j];
 120
 121     xor4 |= rc4_key->S[idx] <<  0;
 122
 123     i += 1;
 124     j += rc4_key->S[i];
 125
 126     swap (rc4_key, i, j);
 127
 128     idx = rc4_key->S[i] + rc4_key->S[j];
 129
 130     xor4 |= rc4_key->S[idx] <<  8;
 131
 132     i += 1;
 133     j += rc4_key->S[i];
 134
 135     swap (rc4_key, i, j);
 136
 137     idx = rc4_key->S[i] + rc4_key->S[j];
 138
 139     xor4 |= rc4_key->S[idx] << 16;
 140
 141     i += 1;
 142     j += rc4_key->S[i];
 143
 144     swap (rc4_key, i, j);
 145
 146     idx = rc4_key->S[i] + rc4_key->S[j];
 147
 148     xor4 |= rc4_key->S[idx] << 24;
 149
 150     out[k] = in[k] ^ xor4;
 151   }
 152
 153   return j;
 154 }
 155
 156 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
 157 {
 158   u32x a = digest[0];
 159   u32x b = digest[1];
 160   u32x c = digest[2];
 161   u32x d = digest[3];
 162
 163   u32x w0_t = w0[0];
 164   u32x w1_t = w0[1];
 165   u32x w2_t = w0[2];
 166   u32x w3_t = w0[3];
 167   u32x w4_t = w1[0];
 168   u32x w5_t = w1[1];
 169   u32x w6_t = w1[2];
 170   u32x w7_t = w1[3];
 171   u32x w8_t = w2[0];
 172   u32x w9_t = w2[1];
 173   u32x wa_t = w2[2];
 174   u32x wb_t = w2[3];
 175   u32x wc_t = w3[0];
 176   u32x wd_t = w3[1];
 177   u32x we_t = w3[2];
 178   u32x wf_t = w3[3];
 179
 180   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
 181   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
 182   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
 183   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
 184   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
 185   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
 186   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
 187   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
 188   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
 189   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
 190   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
 191   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
 192   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
 193   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
 194   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
 195   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
 196
 197   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
 198   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
 199   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
 200   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
 201   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
 202   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
 203   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
 204   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
 205   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
 206   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
 207   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
 208   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
 209   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
 210   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
 211   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 212   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 213
 214   MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
 215   MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
 216   MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
 217   MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
 218   MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
 219   MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
 220   MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
 221   MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
 222   MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
 223   MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
 224   MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
 225   MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
 226   MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
 227   MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
 228   MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
 229   MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
 230
 231   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 232   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 233   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 234   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 235   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 236   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 237   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 238   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 239   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 240   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 241   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 242   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 243   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 244   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 245   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 246   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 247
 248   digest[0] += a;
 249   digest[1] += b;
 250   digest[2] += c;
 251   digest[3] += d;
 252 }
 253
 254 __device__ __constant__ bf_t c_bfs[1024];
 255
 256 __device__ static void m09700m (RC4_KEY rc4_keys[64], 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 u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
 257 {
 258   /**
 259    * modifier
 260    */
 261
 262   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 263   const u32 lid = threadIdx.x;
 264
 265   RC4_KEY *rc4_key = &rc4_keys[lid];
 266
 267   /**
 268    * salt
 269    */
 270
 271   u32 salt_buf_t0[4];
 272   u32 salt_buf_t1[5];
 273   u32 salt_buf_t2[5];
 274   u32 salt_buf_t3[5];
 275
 276   salt_buf_t0[0] = salt_bufs[salt_pos].salt_buf[0];
 277   salt_buf_t0[1] = salt_bufs[salt_pos].salt_buf[1];
 278   salt_buf_t0[2] = salt_bufs[salt_pos].salt_buf[2];
 279   salt_buf_t0[3] = salt_bufs[salt_pos].salt_buf[3];
 280
 281   salt_buf_t1[0] =                        salt_buf_t0[0] <<  8;
 282   salt_buf_t1[1] = salt_buf_t0[0] >> 24 | salt_buf_t0[1] <<  8;
 283   salt_buf_t1[2] = salt_buf_t0[1] >> 24 | salt_buf_t0[2] <<  8;
 284   salt_buf_t1[3] = salt_buf_t0[2] >> 24 | salt_buf_t0[3] <<  8;
 285   salt_buf_t1[4] = salt_buf_t0[3] >> 24;
 286
 287   salt_buf_t2[0] =                        salt_buf_t0[0] << 16;
 288   salt_buf_t2[1] = salt_buf_t0[0] >> 16 | salt_buf_t0[1] << 16;
 289   salt_buf_t2[2] = salt_buf_t0[1] >> 16 | salt_buf_t0[2] << 16;
 290   salt_buf_t2[3] = salt_buf_t0[2] >> 16 | salt_buf_t0[3] << 16;
 291   salt_buf_t2[4] = salt_buf_t0[3] >> 16;
 292
 293   salt_buf_t3[0] =                        salt_buf_t0[0] << 24;
 294   salt_buf_t3[1] = salt_buf_t0[0] >>  8 | salt_buf_t0[1] << 24;
 295   salt_buf_t3[2] = salt_buf_t0[1] >>  8 | salt_buf_t0[2] << 24;
 296   salt_buf_t3[3] = salt_buf_t0[2] >>  8 | salt_buf_t0[3] << 24;
 297   salt_buf_t3[4] = salt_buf_t0[3] >>  8;
 298
 299   const u32 salt_len = 16;
 300
 301   /**
 302    * esalt
 303    */
 304
 305   const u32 version = oldoffice01_bufs[salt_pos].version;
 306
 307   u32 encryptedVerifier[4];
 308
 309   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 310   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 311   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 312   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 313
 314   /**
 315    * loop
 316    */
 317
 318   u32x w0l = w0[0];
 319
 320   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 321   {
 322     const u32 w0r = c_bfs[il_pos].i;
 323
 324     w0[0] = w0l | w0r;
 325
 326     u32x w0_t[4];
 327     u32x w1_t[4];
 328     u32x w2_t[4];
 329     u32x w3_t[4];
 330
 331     w0_t[0] = w0[0];
 332     w0_t[1] = w0[1];
 333     w0_t[2] = w0[2];
 334     w0_t[3] = w0[3];
 335     w1_t[0] = w1[0];
 336     w1_t[1] = w1[1];
 337     w1_t[2] = w1[2];
 338     w1_t[3] = w1[3];
 339     w2_t[0] = w2[0];
 340     w2_t[1] = w2[1];
 341     w2_t[2] = w2[2];
 342     w2_t[3] = w2[3];
 343     w3_t[0] = w3[0];
 344     w3_t[1] = w3[1];
 345     w3_t[2] = pw_len * 8;
 346     w3_t[3] = 0;
 347
 348     u32x digest_t0[4];
 349     u32x digest_t1[2]; // need only first 5 byte
 350     u32x digest_t2[2];
 351     u32x digest_t3[2];
 352
 353     digest_t0[0] = MD5M_A;
 354     digest_t0[1] = MD5M_B;
 355     digest_t0[2] = MD5M_C;
 356     digest_t0[3] = MD5M_D;
 357
 358     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_t0);
 359
 360     // prepare 16 * 21 buffer stuff
 361
 362     u32x digest[4];
 363
 364     digest[0] = MD5M_A;
 365     digest[1] = MD5M_B;
 366     digest[2] = MD5M_C;
 367     digest[3] = MD5M_D;
 368
 369     // offsets
 370
 371     digest_t0[0] &= 0xffffffff;
 372     digest_t0[1] &= 0x000000ff;
 373     digest_t0[2] &= 0x00000000;
 374     digest_t0[3] &= 0x00000000;
 375
 376     digest_t1[0] =                      digest_t0[0] <<  8;
 377     digest_t1[1] = digest_t0[0] >> 24 | digest_t0[1] <<  8;
 378
 379     digest_t2[0] =                      digest_t0[0] << 16;
 380     digest_t2[1] = digest_t0[0] >> 16 | digest_t0[1] << 16;
 381
 382     digest_t3[0] =                      digest_t0[0] << 24;
 383     digest_t3[1] = digest_t0[0] >>  8 | digest_t0[1] << 24;
 384
 385     // generate the 16 * 21 buffer
 386
 387     w0_t[0] = 0;
 388     w0_t[1] = 0;
 389     w0_t[2] = 0;
 390     w0_t[3] = 0;
 391     w1_t[0] = 0;
 392     w1_t[1] = 0;
 393     w1_t[2] = 0;
 394     w1_t[3] = 0;
 395     w2_t[0] = 0;
 396     w2_t[1] = 0;
 397     w2_t[2] = 0;
 398     w2_t[3] = 0;
 399     w3_t[0] = 0;
 400     w3_t[1] = 0;
 401     w3_t[2] = 0;
 402     w3_t[3] = 0;
 403
 404     // 0..5
 405     w0_t[0]  = digest_t0[0];
 406     w0_t[1]  = digest_t0[1];
 407
 408     // 5..21
 409     w0_t[1] |= salt_buf_t1[0];
 410     w0_t[2]  = salt_buf_t1[1];
 411     w0_t[3]  = salt_buf_t1[2];
 412     w1_t[0]  = salt_buf_t1[3];
 413     w1_t[1]  = salt_buf_t1[4];
 414
 415     // 21..26
 416     w1_t[1] |= digest_t1[0];
 417     w1_t[2]  = digest_t1[1];
 418
 419     // 26..42
 420     w1_t[2] |= salt_buf_t2[0];
 421     w1_t[3]  = salt_buf_t2[1];
 422     w2_t[0]  = salt_buf_t2[2];
 423     w2_t[1]  = salt_buf_t2[3];
 424     w2_t[2]  = salt_buf_t2[4];
 425
 426     // 42..47
 427     w2_t[2] |= digest_t2[0];
 428     w2_t[3]  = digest_t2[1];
 429
 430     // 47..63
 431     w2_t[3] |= salt_buf_t3[0];
 432     w3_t[0]  = salt_buf_t3[1];
 433     w3_t[1]  = salt_buf_t3[2];
 434     w3_t[2]  = salt_buf_t3[3];
 435     w3_t[3]  = salt_buf_t3[4];
 436
 437     // 63..
 438
 439     w3_t[3] |= digest_t3[0];
 440
 441     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 442
 443     w0_t[0] = 0;
 444     w0_t[1] = 0;
 445     w0_t[2] = 0;
 446     w0_t[3] = 0;
 447     w1_t[0] = 0;
 448     w1_t[1] = 0;
 449     w1_t[2] = 0;
 450     w1_t[3] = 0;
 451     w2_t[0] = 0;
 452     w2_t[1] = 0;
 453     w2_t[2] = 0;
 454     w2_t[3] = 0;
 455     w3_t[0] = 0;
 456     w3_t[1] = 0;
 457     w3_t[2] = 0;
 458     w3_t[3] = 0;
 459
 460     // 0..4
 461     w0_t[0]  = digest_t3[1];
 462
 463     // 4..20
 464     w0_t[1]  = salt_buf_t0[0];
 465     w0_t[2]  = salt_buf_t0[1];
 466     w0_t[3]  = salt_buf_t0[2];
 467     w1_t[0]  = salt_buf_t0[3];
 468
 469     // 20..25
 470     w1_t[1]  = digest_t0[0];
 471     w1_t[2]  = digest_t0[1];
 472
 473     // 25..41
 474     w1_t[2] |= salt_buf_t1[0];
 475     w1_t[3]  = salt_buf_t1[1];
 476     w2_t[0]  = salt_buf_t1[2];
 477     w2_t[1]  = salt_buf_t1[3];
 478     w2_t[2]  = salt_buf_t1[4];
 479
 480     // 41..46
 481     w2_t[2] |= digest_t1[0];
 482     w2_t[3]  = digest_t1[1];
 483
 484     // 46..62
 485     w2_t[3] |= salt_buf_t2[0];
 486     w3_t[0]  = salt_buf_t2[1];
 487     w3_t[1]  = salt_buf_t2[2];
 488     w3_t[2]  = salt_buf_t2[3];
 489     w3_t[3]  = salt_buf_t2[4];
 490
 491     // 62..
 492     w3_t[3] |= digest_t2[0];
 493
 494     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 495
 496     w0_t[0] = 0;
 497     w0_t[1] = 0;
 498     w0_t[2] = 0;
 499     w0_t[3] = 0;
 500     w1_t[0] = 0;
 501     w1_t[1] = 0;
 502     w1_t[2] = 0;
 503     w1_t[3] = 0;
 504     w2_t[0] = 0;
 505     w2_t[1] = 0;
 506     w2_t[2] = 0;
 507     w2_t[3] = 0;
 508     w3_t[0] = 0;
 509     w3_t[1] = 0;
 510     w3_t[2] = 0;
 511     w3_t[3] = 0;
 512
 513     // 0..3
 514     w0_t[0]  = digest_t2[1];
 515
 516     // 3..19
 517     w0_t[0] |= salt_buf_t3[0];
 518     w0_t[1]  = salt_buf_t3[1];
 519     w0_t[2]  = salt_buf_t3[2];
 520     w0_t[3]  = salt_buf_t3[3];
 521     w1_t[0]  = salt_buf_t3[4];
 522
 523     // 19..24
 524     w1_t[0] |= digest_t3[0];
 525     w1_t[1]  = digest_t3[1];
 526
 527     // 24..40
 528     w1_t[2]  = salt_buf_t0[0];
 529     w1_t[3]  = salt_buf_t0[1];
 530     w2_t[0]  = salt_buf_t0[2];
 531     w2_t[1]  = salt_buf_t0[3];
 532
 533     // 40..45
 534     w2_t[2]  = digest_t0[0];
 535     w2_t[3]  = digest_t0[1];
 536
 537     // 45..61
 538     w2_t[3] |= salt_buf_t1[0];
 539     w3_t[0]  = salt_buf_t1[1];
 540     w3_t[1]  = salt_buf_t1[2];
 541     w3_t[2]  = salt_buf_t1[3];
 542     w3_t[3]  = salt_buf_t1[4];
 543
 544     // 61..
 545     w3_t[3] |= digest_t1[0];
 546
 547     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 548
 549     w0_t[0] = 0;
 550     w0_t[1] = 0;
 551     w0_t[2] = 0;
 552     w0_t[3] = 0;
 553     w1_t[0] = 0;
 554     w1_t[1] = 0;
 555     w1_t[2] = 0;
 556     w1_t[3] = 0;
 557     w2_t[0] = 0;
 558     w2_t[1] = 0;
 559     w2_t[2] = 0;
 560     w2_t[3] = 0;
 561     w3_t[0] = 0;
 562     w3_t[1] = 0;
 563     w3_t[2] = 0;
 564     w3_t[3] = 0;
 565
 566     // 0..2
 567     w0_t[0]  = digest_t1[1];
 568
 569     // 2..18
 570     w0_t[0] |= salt_buf_t2[0];
 571     w0_t[1]  = salt_buf_t2[1];
 572     w0_t[2]  = salt_buf_t2[2];
 573     w0_t[3]  = salt_buf_t2[3];
 574     w1_t[0]  = salt_buf_t2[4];
 575
 576     // 18..23
 577     w1_t[0] |= digest_t2[0];
 578     w1_t[1]  = digest_t2[1];
 579
 580     // 23..39
 581     w1_t[1] |= salt_buf_t3[0];
 582     w1_t[2]  = salt_buf_t3[1];
 583     w1_t[3]  = salt_buf_t3[2];
 584     w2_t[0]  = salt_buf_t3[3];
 585     w2_t[1]  = salt_buf_t3[4];
 586
 587     // 39..44
 588     w2_t[1] |= digest_t3[0];
 589     w2_t[2]  = digest_t3[1];
 590
 591     // 44..60
 592     w2_t[3]  = salt_buf_t0[0];
 593     w3_t[0]  = salt_buf_t0[1];
 594     w3_t[1]  = salt_buf_t0[2];
 595     w3_t[2]  = salt_buf_t0[3];
 596
 597     // 60..
 598     w3_t[3]  = digest_t0[0];
 599
 600     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 601
 602     w0_t[0] = 0;
 603     w0_t[1] = 0;
 604     w0_t[2] = 0;
 605     w0_t[3] = 0;
 606     w1_t[0] = 0;
 607     w1_t[1] = 0;
 608     w1_t[2] = 0;
 609     w1_t[3] = 0;
 610     w2_t[0] = 0;
 611     w2_t[1] = 0;
 612     w2_t[2] = 0;
 613     w2_t[3] = 0;
 614     w3_t[0] = 0;
 615     w3_t[1] = 0;
 616     w3_t[2] = 0;
 617     w3_t[3] = 0;
 618
 619     // 0..1
 620     w0_t[0]  = digest_t0[1];
 621
 622     // 1..17
 623     w0_t[0] |= salt_buf_t1[0];
 624     w0_t[1]  = salt_buf_t1[1];
 625     w0_t[2]  = salt_buf_t1[2];
 626     w0_t[3]  = salt_buf_t1[3];
 627     w1_t[0]  = salt_buf_t1[4];
 628
 629     // 17..22
 630     w1_t[0] |= digest_t1[0];
 631     w1_t[1]  = digest_t1[1];
 632
 633     // 22..38
 634     w1_t[1] |= salt_buf_t2[0];
 635     w1_t[2]  = salt_buf_t2[1];
 636     w1_t[3]  = salt_buf_t2[2];
 637     w2_t[0]  = salt_buf_t2[3];
 638     w2_t[1]  = salt_buf_t2[4];
 639
 640     // 38..43
 641     w2_t[1] |= digest_t2[0];
 642     w2_t[2]  = digest_t2[1];
 643
 644     // 43..59
 645     w2_t[2] |= salt_buf_t3[0];
 646     w2_t[3]  = salt_buf_t3[1];
 647     w3_t[0]  = salt_buf_t3[2];
 648     w3_t[1]  = salt_buf_t3[3];
 649     w3_t[2]  = salt_buf_t3[4];
 650
 651     // 59..
 652     w3_t[2] |= digest_t3[0];
 653     w3_t[3]  = digest_t3[1];
 654
 655     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 656
 657     w0_t[0]  = salt_buf_t0[0];
 658     w0_t[1]  = salt_buf_t0[1];
 659     w0_t[2]  = salt_buf_t0[2];
 660     w0_t[3]  = salt_buf_t0[3];
 661     w1_t[0]  = 0x80;
 662     w1_t[1]  = 0;
 663     w1_t[2]  = 0;
 664     w1_t[3]  = 0;
 665     w2_t[0]  = 0;
 666     w2_t[1]  = 0;
 667     w2_t[2]  = 0;
 668     w2_t[3]  = 0;
 669     w3_t[0]  = 0;
 670     w3_t[1]  = 0;
 671     w3_t[2]  = 21 * 16 * 8;
 672     w3_t[3]  = 0;
 673
 674     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 675
 676     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
 677
 678     w0_t[0]  = digest[0];
 679     w0_t[1]  = digest[1] & 0xff;
 680     w0_t[2]  = 0x8000;
 681     w0_t[3]  = 0;
 682     w1_t[0]  = 0;
 683     w1_t[1]  = 0;
 684     w1_t[2]  = 0;
 685     w1_t[3]  = 0;
 686     w2_t[0]  = 0;
 687     w2_t[1]  = 0;
 688     w2_t[2]  = 0;
 689     w2_t[3]  = 0;
 690     w3_t[0]  = 0;
 691     w3_t[1]  = 0;
 692     w3_t[2]  = 9 * 8;
 693     w3_t[3]  = 0;
 694
 695     digest[0] = MD5M_A;
 696     digest[1] = MD5M_B;
 697     digest[2] = MD5M_C;
 698     digest[3] = MD5M_D;
 699
 700     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 701
 702     // now the RC4 part
 703
 704     u32x key[4];
 705
 706     key[0] = digest[0];
 707     key[1] = digest[1];
 708     key[2] = digest[2];
 709     key[3] = digest[3];
 710
 711     rc4_init_16 (rc4_key, key);
 712
 713     u32x out[4];
 714
 715     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 716
 717     w0_t[0] = out[0];
 718     w0_t[1] = out[1];
 719     w0_t[2] = out[2];
 720     w0_t[3] = out[3];
 721     w1_t[0] = 0x80;
 722     w1_t[1] = 0;
 723     w1_t[2] = 0;
 724     w1_t[3] = 0;
 725     w2_t[0] = 0;
 726     w2_t[1] = 0;
 727     w2_t[2] = 0;
 728     w2_t[3] = 0;
 729     w3_t[0] = 0;
 730     w3_t[1] = 0;
 731     w3_t[2] = 16 * 8;
 732     w3_t[3] = 0;
 733
 734     digest[0] = MD5M_A;
 735     digest[1] = MD5M_B;
 736     digest[2] = MD5M_C;
 737     digest[3] = MD5M_D;
 738
 739     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 740
 741     rc4_next_16 (rc4_key, 16, j, digest, out);
 742
 743     const u32x r0 = out[0];
 744     const u32x r1 = out[1];
 745     const u32x r2 = out[2];
 746     const u32x r3 = out[3];
 747
 748     #include VECT_COMPARE_M
 749   }
 750 }
 751
 752 __device__ static void m09700s (RC4_KEY rc4_keys[64], 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 u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
 753 {
 754   /**
 755    * modifier
 756    */
 757
 758   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 759   const u32 lid = threadIdx.x;
 760
 761   RC4_KEY *rc4_key = &rc4_keys[lid];
 762
 763   /**
 764    * digest
 765    */
 766
 767   const u32 search[4] =
 768   {
 769     digests_buf[digests_offset].digest_buf[DGST_R0],
 770     digests_buf[digests_offset].digest_buf[DGST_R1],
 771     digests_buf[digests_offset].digest_buf[DGST_R2],
 772     digests_buf[digests_offset].digest_buf[DGST_R3]
 773   };
 774
 775   /**
 776    * salt
 777    */
 778
 779   u32 salt_buf_t0[4];
 780   u32 salt_buf_t1[5];
 781   u32 salt_buf_t2[5];
 782   u32 salt_buf_t3[5];
 783
 784   salt_buf_t0[0] = salt_bufs[salt_pos].salt_buf[0];
 785   salt_buf_t0[1] = salt_bufs[salt_pos].salt_buf[1];
 786   salt_buf_t0[2] = salt_bufs[salt_pos].salt_buf[2];
 787   salt_buf_t0[3] = salt_bufs[salt_pos].salt_buf[3];
 788
 789   salt_buf_t1[0] =                        salt_buf_t0[0] <<  8;
 790   salt_buf_t1[1] = salt_buf_t0[0] >> 24 | salt_buf_t0[1] <<  8;
 791   salt_buf_t1[2] = salt_buf_t0[1] >> 24 | salt_buf_t0[2] <<  8;
 792   salt_buf_t1[3] = salt_buf_t0[2] >> 24 | salt_buf_t0[3] <<  8;
 793   salt_buf_t1[4] = salt_buf_t0[3] >> 24;
 794
 795   salt_buf_t2[0] =                        salt_buf_t0[0] << 16;
 796   salt_buf_t2[1] = salt_buf_t0[0] >> 16 | salt_buf_t0[1] << 16;
 797   salt_buf_t2[2] = salt_buf_t0[1] >> 16 | salt_buf_t0[2] << 16;
 798   salt_buf_t2[3] = salt_buf_t0[2] >> 16 | salt_buf_t0[3] << 16;
 799   salt_buf_t2[4] = salt_buf_t0[3] >> 16;
 800
 801   salt_buf_t3[0] =                        salt_buf_t0[0] << 24;
 802   salt_buf_t3[1] = salt_buf_t0[0] >>  8 | salt_buf_t0[1] << 24;
 803   salt_buf_t3[2] = salt_buf_t0[1] >>  8 | salt_buf_t0[2] << 24;
 804   salt_buf_t3[3] = salt_buf_t0[2] >>  8 | salt_buf_t0[3] << 24;
 805   salt_buf_t3[4] = salt_buf_t0[3] >>  8;
 806
 807   const u32 salt_len = 16;
 808
 809   /**
 810    * esalt
 811    */
 812
 813   const u32 version = oldoffice01_bufs[salt_pos].version;
 814
 815   u32 encryptedVerifier[4];
 816
 817   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 818   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 819   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 820   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 821
 822   /**
 823    * loop
 824    */
 825
 826   u32x w0l = w0[0];
 827
 828   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 829   {
 830     const u32 w0r = c_bfs[il_pos].i;
 831
 832     w0[0] = w0l | w0r;
 833
 834     u32x w0_t[4];
 835     u32x w1_t[4];
 836     u32x w2_t[4];
 837     u32x w3_t[4];
 838
 839     w0_t[0] = w0[0];
 840     w0_t[1] = w0[1];
 841     w0_t[2] = w0[2];
 842     w0_t[3] = w0[3];
 843     w1_t[0] = w1[0];
 844     w1_t[1] = w1[1];
 845     w1_t[2] = w1[2];
 846     w1_t[3] = w1[3];
 847     w2_t[0] = w2[0];
 848     w2_t[1] = w2[1];
 849     w2_t[2] = w2[2];
 850     w2_t[3] = w2[3];
 851     w3_t[0] = w3[0];
 852     w3_t[1] = w3[1];
 853     w3_t[2] = pw_len * 8;
 854     w3_t[3] = 0;
 855
 856     u32x digest_t0[4];
 857     u32x digest_t1[2]; // need only first 5 byte
 858     u32x digest_t2[2];
 859     u32x digest_t3[2];
 860
 861     digest_t0[0] = MD5M_A;
 862     digest_t0[1] = MD5M_B;
 863     digest_t0[2] = MD5M_C;
 864     digest_t0[3] = MD5M_D;
 865
 866     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_t0);
 867
 868     // prepare 16 * 21 buffer stuff
 869
 870     u32x digest[4];
 871
 872     digest[0] = MD5M_A;
 873     digest[1] = MD5M_B;
 874     digest[2] = MD5M_C;
 875     digest[3] = MD5M_D;
 876
 877     // offsets
 878
 879     digest_t0[0] &= 0xffffffff;
 880     digest_t0[1] &= 0x000000ff;
 881     digest_t0[2] &= 0x00000000;
 882     digest_t0[3] &= 0x00000000;
 883
 884     digest_t1[0] =                      digest_t0[0] <<  8;
 885     digest_t1[1] = digest_t0[0] >> 24 | digest_t0[1] <<  8;
 886
 887     digest_t2[0] =                      digest_t0[0] << 16;
 888     digest_t2[1] = digest_t0[0] >> 16 | digest_t0[1] << 16;
 889
 890     digest_t3[0] =                      digest_t0[0] << 24;
 891     digest_t3[1] = digest_t0[0] >>  8 | digest_t0[1] << 24;
 892
 893     // generate the 16 * 21 buffer
 894
 895     w0_t[0] = 0;
 896     w0_t[1] = 0;
 897     w0_t[2] = 0;
 898     w0_t[3] = 0;
 899     w1_t[0] = 0;
 900     w1_t[1] = 0;
 901     w1_t[2] = 0;
 902     w1_t[3] = 0;
 903     w2_t[0] = 0;
 904     w2_t[1] = 0;
 905     w2_t[2] = 0;
 906     w2_t[3] = 0;
 907     w3_t[0] = 0;
 908     w3_t[1] = 0;
 909     w3_t[2] = 0;
 910     w3_t[3] = 0;
 911
 912     // 0..5
 913     w0_t[0]  = digest_t0[0];
 914     w0_t[1]  = digest_t0[1];
 915
 916     // 5..21
 917     w0_t[1] |= salt_buf_t1[0];
 918     w0_t[2]  = salt_buf_t1[1];
 919     w0_t[3]  = salt_buf_t1[2];
 920     w1_t[0]  = salt_buf_t1[3];
 921     w1_t[1]  = salt_buf_t1[4];
 922
 923     // 21..26
 924     w1_t[1] |= digest_t1[0];
 925     w1_t[2]  = digest_t1[1];
 926
 927     // 26..42
 928     w1_t[2] |= salt_buf_t2[0];
 929     w1_t[3]  = salt_buf_t2[1];
 930     w2_t[0]  = salt_buf_t2[2];
 931     w2_t[1]  = salt_buf_t2[3];
 932     w2_t[2]  = salt_buf_t2[4];
 933
 934     // 42..47
 935     w2_t[2] |= digest_t2[0];
 936     w2_t[3]  = digest_t2[1];
 937
 938     // 47..63
 939     w2_t[3] |= salt_buf_t3[0];
 940     w3_t[0]  = salt_buf_t3[1];
 941     w3_t[1]  = salt_buf_t3[2];
 942     w3_t[2]  = salt_buf_t3[3];
 943     w3_t[3]  = salt_buf_t3[4];
 944
 945     // 63..
 946
 947     w3_t[3] |= digest_t3[0];
 948
 949     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 950
 951     w0_t[0] = 0;
 952     w0_t[1] = 0;
 953     w0_t[2] = 0;
 954     w0_t[3] = 0;
 955     w1_t[0] = 0;
 956     w1_t[1] = 0;
 957     w1_t[2] = 0;
 958     w1_t[3] = 0;
 959     w2_t[0] = 0;
 960     w2_t[1] = 0;
 961     w2_t[2] = 0;
 962     w2_t[3] = 0;
 963     w3_t[0] = 0;
 964     w3_t[1] = 0;
 965     w3_t[2] = 0;
 966     w3_t[3] = 0;
 967
 968     // 0..4
 969     w0_t[0]  = digest_t3[1];
 970
 971     // 4..20
 972     w0_t[1]  = salt_buf_t0[0];
 973     w0_t[2]  = salt_buf_t0[1];
 974     w0_t[3]  = salt_buf_t0[2];
 975     w1_t[0]  = salt_buf_t0[3];
 976
 977     // 20..25
 978     w1_t[1]  = digest_t0[0];
 979     w1_t[2]  = digest_t0[1];
 980
 981     // 25..41
 982     w1_t[2] |= salt_buf_t1[0];
 983     w1_t[3]  = salt_buf_t1[1];
 984     w2_t[0]  = salt_buf_t1[2];
 985     w2_t[1]  = salt_buf_t1[3];
 986     w2_t[2]  = salt_buf_t1[4];
 987
 988     // 41..46
 989     w2_t[2] |= digest_t1[0];
 990     w2_t[3]  = digest_t1[1];
 991
 992     // 46..62
 993     w2_t[3] |= salt_buf_t2[0];
 994     w3_t[0]  = salt_buf_t2[1];
 995     w3_t[1]  = salt_buf_t2[2];
 996     w3_t[2]  = salt_buf_t2[3];
 997     w3_t[3]  = salt_buf_t2[4];
 998
 999     // 62..
1000     w3_t[3] |= digest_t2[0];
1001
1002     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1003
1004     w0_t[0] = 0;
1005     w0_t[1] = 0;
1006     w0_t[2] = 0;
1007     w0_t[3] = 0;
1008     w1_t[0] = 0;
1009     w1_t[1] = 0;
1010     w1_t[2] = 0;
1011     w1_t[3] = 0;
1012     w2_t[0] = 0;
1013     w2_t[1] = 0;
1014     w2_t[2] = 0;
1015     w2_t[3] = 0;
1016     w3_t[0] = 0;
1017     w3_t[1] = 0;
1018     w3_t[2] = 0;
1019     w3_t[3] = 0;
1020
1021     // 0..3
1022     w0_t[0]  = digest_t2[1];
1023
1024     // 3..19
1025     w0_t[0] |= salt_buf_t3[0];
1026     w0_t[1]  = salt_buf_t3[1];
1027     w0_t[2]  = salt_buf_t3[2];
1028     w0_t[3]  = salt_buf_t3[3];
1029     w1_t[0]  = salt_buf_t3[4];
1030
1031     // 19..24
1032     w1_t[0] |= digest_t3[0];
1033     w1_t[1]  = digest_t3[1];
1034
1035     // 24..40
1036     w1_t[2]  = salt_buf_t0[0];
1037     w1_t[3]  = salt_buf_t0[1];
1038     w2_t[0]  = salt_buf_t0[2];
1039     w2_t[1]  = salt_buf_t0[3];
1040
1041     // 40..45
1042     w2_t[2]  = digest_t0[0];
1043     w2_t[3]  = digest_t0[1];
1044
1045     // 45..61
1046     w2_t[3] |= salt_buf_t1[0];
1047     w3_t[0]  = salt_buf_t1[1];
1048     w3_t[1]  = salt_buf_t1[2];
1049     w3_t[2]  = salt_buf_t1[3];
1050     w3_t[3]  = salt_buf_t1[4];
1051
1052     // 61..
1053     w3_t[3] |= digest_t1[0];
1054
1055     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1056
1057     w0_t[0] = 0;
1058     w0_t[1] = 0;
1059     w0_t[2] = 0;
1060     w0_t[3] = 0;
1061     w1_t[0] = 0;
1062     w1_t[1] = 0;
1063     w1_t[2] = 0;
1064     w1_t[3] = 0;
1065     w2_t[0] = 0;
1066     w2_t[1] = 0;
1067     w2_t[2] = 0;
1068     w2_t[3] = 0;
1069     w3_t[0] = 0;
1070     w3_t[1] = 0;
1071     w3_t[2] = 0;
1072     w3_t[3] = 0;
1073
1074     // 0..2
1075     w0_t[0]  = digest_t1[1];
1076
1077     // 2..18
1078     w0_t[0] |= salt_buf_t2[0];
1079     w0_t[1]  = salt_buf_t2[1];
1080     w0_t[2]  = salt_buf_t2[2];
1081     w0_t[3]  = salt_buf_t2[3];
1082     w1_t[0]  = salt_buf_t2[4];
1083
1084     // 18..23
1085     w1_t[0] |= digest_t2[0];
1086     w1_t[1]  = digest_t2[1];
1087
1088     // 23..39
1089     w1_t[1] |= salt_buf_t3[0];
1090     w1_t[2]  = salt_buf_t3[1];
1091     w1_t[3]  = salt_buf_t3[2];
1092     w2_t[0]  = salt_buf_t3[3];
1093     w2_t[1]  = salt_buf_t3[4];
1094
1095     // 39..44
1096     w2_t[1] |= digest_t3[0];
1097     w2_t[2]  = digest_t3[1];
1098
1099     // 44..60
1100     w2_t[3]  = salt_buf_t0[0];
1101     w3_t[0]  = salt_buf_t0[1];
1102     w3_t[1]  = salt_buf_t0[2];
1103     w3_t[2]  = salt_buf_t0[3];
1104
1105     // 60..
1106     w3_t[3]  = digest_t0[0];
1107
1108     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1109
1110     w0_t[0] = 0;
1111     w0_t[1] = 0;
1112     w0_t[2] = 0;
1113     w0_t[3] = 0;
1114     w1_t[0] = 0;
1115     w1_t[1] = 0;
1116     w1_t[2] = 0;
1117     w1_t[3] = 0;
1118     w2_t[0] = 0;
1119     w2_t[1] = 0;
1120     w2_t[2] = 0;
1121     w2_t[3] = 0;
1122     w3_t[0] = 0;
1123     w3_t[1] = 0;
1124     w3_t[2] = 0;
1125     w3_t[3] = 0;
1126
1127     // 0..1
1128     w0_t[0]  = digest_t0[1];
1129
1130     // 1..17
1131     w0_t[0] |= salt_buf_t1[0];
1132     w0_t[1]  = salt_buf_t1[1];
1133     w0_t[2]  = salt_buf_t1[2];
1134     w0_t[3]  = salt_buf_t1[3];
1135     w1_t[0]  = salt_buf_t1[4];
1136
1137     // 17..22
1138     w1_t[0] |= digest_t1[0];
1139     w1_t[1]  = digest_t1[1];
1140
1141     // 22..38
1142     w1_t[1] |= salt_buf_t2[0];
1143     w1_t[2]  = salt_buf_t2[1];
1144     w1_t[3]  = salt_buf_t2[2];
1145     w2_t[0]  = salt_buf_t2[3];
1146     w2_t[1]  = salt_buf_t2[4];
1147
1148     // 38..43
1149     w2_t[1] |= digest_t2[0];
1150     w2_t[2]  = digest_t2[1];
1151
1152     // 43..59
1153     w2_t[2] |= salt_buf_t3[0];
1154     w2_t[3]  = salt_buf_t3[1];
1155     w3_t[0]  = salt_buf_t3[2];
1156     w3_t[1]  = salt_buf_t3[3];
1157     w3_t[2]  = salt_buf_t3[4];
1158
1159     // 59..
1160     w3_t[2] |= digest_t3[0];
1161     w3_t[3]  = digest_t3[1];
1162
1163     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1164
1165     w0_t[0]  = salt_buf_t0[0];
1166     w0_t[1]  = salt_buf_t0[1];
1167     w0_t[2]  = salt_buf_t0[2];
1168     w0_t[3]  = salt_buf_t0[3];
1169     w1_t[0]  = 0x80;
1170     w1_t[1]  = 0;
1171     w1_t[2]  = 0;
1172     w1_t[3]  = 0;
1173     w2_t[0]  = 0;
1174     w2_t[1]  = 0;
1175     w2_t[2]  = 0;
1176     w2_t[3]  = 0;
1177     w3_t[0]  = 0;
1178     w3_t[1]  = 0;
1179     w3_t[2]  = 21 * 16 * 8;
1180     w3_t[3]  = 0;
1181
1182     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1183
1184     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
1185
1186     w0_t[0]  = digest[0];
1187     w0_t[1]  = digest[1] & 0xff;
1188     w0_t[2]  = 0x8000;
1189     w0_t[3]  = 0;
1190     w1_t[0]  = 0;
1191     w1_t[1]  = 0;
1192     w1_t[2]  = 0;
1193     w1_t[3]  = 0;
1194     w2_t[0]  = 0;
1195     w2_t[1]  = 0;
1196     w2_t[2]  = 0;
1197     w2_t[3]  = 0;
1198     w3_t[0]  = 0;
1199     w3_t[1]  = 0;
1200     w3_t[2]  = 9 * 8;
1201     w3_t[3]  = 0;
1202
1203     digest[0] = MD5M_A;
1204     digest[1] = MD5M_B;
1205     digest[2] = MD5M_C;
1206     digest[3] = MD5M_D;
1207
1208     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1209
1210     // now the RC4 part
1211
1212     u32x key[4];
1213
1214     key[0] = digest[0];
1215     key[1] = digest[1];
1216     key[2] = digest[2];
1217     key[3] = digest[3];
1218
1219     rc4_init_16 (rc4_key, key);
1220
1221     u32x out[4];
1222
1223     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
1224
1225     w0_t[0] = out[0];
1226     w0_t[1] = out[1];
1227     w0_t[2] = out[2];
1228     w0_t[3] = out[3];
1229     w1_t[0] = 0x80;
1230     w1_t[1] = 0;
1231     w1_t[2] = 0;
1232     w1_t[3] = 0;
1233     w2_t[0] = 0;
1234     w2_t[1] = 0;
1235     w2_t[2] = 0;
1236     w2_t[3] = 0;
1237     w3_t[0] = 0;
1238     w3_t[1] = 0;
1239     w3_t[2] = 16 * 8;
1240     w3_t[3] = 0;
1241
1242     digest[0] = MD5M_A;
1243     digest[1] = MD5M_B;
1244     digest[2] = MD5M_C;
1245     digest[3] = MD5M_D;
1246
1247     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1248
1249     rc4_next_16 (rc4_key, 16, j, digest, out);
1250
1251     const u32x r0 = out[0];
1252     const u32x r1 = out[1];
1253     const u32x r2 = out[2];
1254     const u32x r3 = out[3];
1255
1256     #include VECT_COMPARE_S
1257   }
1258 }
1259
1260 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1261 {
1262   /**
1263    * base
1264    */
1265
1266   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1267
1268   if (gid >= gid_max) return;
1269
1270   u32x w0[4];
1271
1272   w0[0] = pws[gid].i[ 0];
1273   w0[1] = pws[gid].i[ 1];
1274   w0[2] = pws[gid].i[ 2];
1275   w0[3] = pws[gid].i[ 3];
1276
1277   u32x w1[4];
1278
1279   w1[0] = 0;
1280   w1[1] = 0;
1281   w1[2] = 0;
1282   w1[3] = 0;
1283
1284   u32x w2[4];
1285
1286   w2[0] = 0;
1287   w2[1] = 0;
1288   w2[2] = 0;
1289   w2[3] = 0;
1290
1291   u32x w3[4];
1292
1293   w3[0] = 0;
1294   w3[1] = 0;
1295   w3[2] = 0;
1296   w3[3] = 0;
1297
1298   const u32 pw_len = pws[gid].pw_len;
1299
1300   /**
1301    * main
1302    */
1303
1304   __shared__ RC4_KEY rc4_keys[64];
1305
1306   m09700m (rc4_keys, 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, oldoffice01_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);
1307 }
1308
1309 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1310 {
1311   /**
1312    * base
1313    */
1314
1315   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1316
1317   if (gid >= gid_max) return;
1318
1319   u32x w0[4];
1320
1321   w0[0] = pws[gid].i[ 0];
1322   w0[1] = pws[gid].i[ 1];
1323   w0[2] = pws[gid].i[ 2];
1324   w0[3] = pws[gid].i[ 3];
1325
1326   u32x w1[4];
1327
1328   w1[0] = pws[gid].i[ 4];
1329   w1[1] = pws[gid].i[ 5];
1330   w1[2] = pws[gid].i[ 6];
1331   w1[3] = pws[gid].i[ 7];
1332
1333   u32x w2[4];
1334
1335   w2[0] = 0;
1336   w2[1] = 0;
1337   w2[2] = 0;
1338   w2[3] = 0;
1339
1340   u32x w3[4];
1341
1342   w3[0] = 0;
1343   w3[1] = 0;
1344   w3[2] = 0;
1345   w3[3] = 0;
1346
1347   const u32 pw_len = pws[gid].pw_len;
1348
1349   /**
1350    * main
1351    */
1352
1353   __shared__ RC4_KEY rc4_keys[64];
1354
1355   m09700m (rc4_keys, 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, oldoffice01_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);
1356 }
1357
1358 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1359 {
1360   /**
1361    * base
1362    */
1363
1364   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1365
1366   if (gid >= gid_max) return;
1367
1368   u32x w0[4];
1369
1370   w0[0] = pws[gid].i[ 0];
1371   w0[1] = pws[gid].i[ 1];
1372   w0[2] = pws[gid].i[ 2];
1373   w0[3] = pws[gid].i[ 3];
1374
1375   u32x w1[4];
1376
1377   w1[0] = pws[gid].i[ 4];
1378   w1[1] = pws[gid].i[ 5];
1379   w1[2] = pws[gid].i[ 6];
1380   w1[3] = pws[gid].i[ 7];
1381
1382   u32x w2[4];
1383
1384   w2[0] = pws[gid].i[ 8];
1385   w2[1] = pws[gid].i[ 9];
1386   w2[2] = pws[gid].i[10];
1387   w2[3] = pws[gid].i[11];
1388
1389   u32x w3[4];
1390
1391   w3[0] = pws[gid].i[12];
1392   w3[1] = pws[gid].i[13];
1393   w3[2] = 0;
1394   w3[3] = 0;
1395
1396   const u32 pw_len = pws[gid].pw_len;
1397
1398   /**
1399    * main
1400    */
1401
1402   __shared__ RC4_KEY rc4_keys[64];
1403
1404   m09700m (rc4_keys, 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, oldoffice01_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);
1405 }
1406
1407 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1408 {
1409   /**
1410    * base
1411    */
1412
1413   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1414
1415   if (gid >= gid_max) return;
1416
1417   u32x w0[4];
1418
1419   w0[0] = pws[gid].i[ 0];
1420   w0[1] = pws[gid].i[ 1];
1421   w0[2] = pws[gid].i[ 2];
1422   w0[3] = pws[gid].i[ 3];
1423
1424   u32x w1[4];
1425
1426   w1[0] = 0;
1427   w1[1] = 0;
1428   w1[2] = 0;
1429   w1[3] = 0;
1430
1431   u32x w2[4];
1432
1433   w2[0] = 0;
1434   w2[1] = 0;
1435   w2[2] = 0;
1436   w2[3] = 0;
1437
1438   u32x w3[4];
1439
1440   w3[0] = 0;
1441   w3[1] = 0;
1442   w3[2] = 0;
1443   w3[3] = 0;
1444
1445   const u32 pw_len = pws[gid].pw_len;
1446
1447   /**
1448    * main
1449    */
1450
1451   __shared__ RC4_KEY rc4_keys[64];
1452
1453   m09700s (rc4_keys, 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, oldoffice01_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);
1454 }
1455
1456 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1457 {
1458   /**
1459    * base
1460    */
1461
1462   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1463
1464   if (gid >= gid_max) return;
1465
1466   u32x w0[4];
1467
1468   w0[0] = pws[gid].i[ 0];
1469   w0[1] = pws[gid].i[ 1];
1470   w0[2] = pws[gid].i[ 2];
1471   w0[3] = pws[gid].i[ 3];
1472
1473   u32x w1[4];
1474
1475   w1[0] = pws[gid].i[ 4];
1476   w1[1] = pws[gid].i[ 5];
1477   w1[2] = pws[gid].i[ 6];
1478   w1[3] = pws[gid].i[ 7];
1479
1480   u32x w2[4];
1481
1482   w2[0] = 0;
1483   w2[1] = 0;
1484   w2[2] = 0;
1485   w2[3] = 0;
1486
1487   u32x w3[4];
1488
1489   w3[0] = 0;
1490   w3[1] = 0;
1491   w3[2] = 0;
1492   w3[3] = 0;
1493
1494   const u32 pw_len = pws[gid].pw_len;
1495
1496   /**
1497    * main
1498    */
1499
1500   __shared__ RC4_KEY rc4_keys[64];
1501
1502   m09700s (rc4_keys, 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, oldoffice01_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);
1503 }
1504
1505 extern "C" __global__ void __launch_bounds__ (64, 1) m09700_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 oldoffice01_t *oldoffice01_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)
1506 {
1507   /**
1508    * base
1509    */
1510
1511   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1512
1513   if (gid >= gid_max) return;
1514
1515   u32x w0[4];
1516
1517   w0[0] = pws[gid].i[ 0];
1518   w0[1] = pws[gid].i[ 1];
1519   w0[2] = pws[gid].i[ 2];
1520   w0[3] = pws[gid].i[ 3];
1521
1522   u32x w1[4];
1523
1524   w1[0] = pws[gid].i[ 4];
1525   w1[1] = pws[gid].i[ 5];
1526   w1[2] = pws[gid].i[ 6];
1527   w1[3] = pws[gid].i[ 7];
1528
1529   u32x w2[4];
1530
1531   w2[0] = pws[gid].i[ 8];
1532   w2[1] = pws[gid].i[ 9];
1533   w2[2] = pws[gid].i[10];
1534   w2[3] = pws[gid].i[11];
1535
1536   u32x w3[4];
1537
1538   w3[0] = pws[gid].i[12];
1539   w3[1] = pws[gid].i[13];
1540   w3[2] = 0;
1541   w3[3] = 0;
1542
1543   const u32 pw_len = pws[gid].pw_len;
1544
1545   /**
1546    * main
1547    */
1548
1549   __shared__ RC4_KEY rc4_keys[64];
1550
1551   m09700s (rc4_keys, 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, oldoffice01_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);
1552 }