nv/m09710_a1.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_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 #ifdef  VECT_SIZE4
  39 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  40 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  41 #endif
  42
  43 typedef struct
  44 {
  45   u8 S[256];
  46
  47   u32 wtf_its_faster;
  48
  49 } RC4_KEY;
  50
  51 __device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
  52 {
  53   u8 tmp;
  54
  55   tmp           = rc4_key->S[i];
  56   rc4_key->S[i] = rc4_key->S[j];
  57   rc4_key->S[j] = tmp;
  58 }
  59
  60 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  61 {
  62   u32 v = 0x03020100;
  63   u32 a = 0x04040404;
  64
  65   u32 *ptr = (u32 *) rc4_key->S;
  66
  67   #pragma unroll 64
  68   for (u32 i = 0; i < 64; i++)
  69   {
  70     *ptr++ = v; v += a;
  71   }
  72
  73   u32 j = 0;
  74
  75   for (u32 i = 0; i < 16; i++)
  76   {
  77     u32 idx = i * 16;
  78
  79     u32 v;
  80
  81     v = data[0];
  82
  83     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  84     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  85     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  86     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  87
  88     v = data[1];
  89
  90     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  91     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  92     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  93     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  94
  95     v = data[2];
  96
  97     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  98     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  99     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 100     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 101
 102     v = data[3];
 103
 104     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
 105     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
 106     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 107     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 108   }
 109 }
 110
 111 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 112 {
 113   for (u32 k = 0; k < 4; k++)
 114   {
 115     u32 xor4 = 0;
 116
 117     u8 idx;
 118
 119     i += 1;
 120     j += rc4_key->S[i];
 121
 122     swap (rc4_key, i, j);
 123
 124     idx = rc4_key->S[i] + rc4_key->S[j];
 125
 126     xor4 |= rc4_key->S[idx] <<  0;
 127
 128     i += 1;
 129     j += rc4_key->S[i];
 130
 131     swap (rc4_key, i, j);
 132
 133     idx = rc4_key->S[i] + rc4_key->S[j];
 134
 135     xor4 |= rc4_key->S[idx] <<  8;
 136
 137     i += 1;
 138     j += rc4_key->S[i];
 139
 140     swap (rc4_key, i, j);
 141
 142     idx = rc4_key->S[i] + rc4_key->S[j];
 143
 144     xor4 |= rc4_key->S[idx] << 16;
 145
 146     i += 1;
 147     j += rc4_key->S[i];
 148
 149     swap (rc4_key, i, j);
 150
 151     idx = rc4_key->S[i] + rc4_key->S[j];
 152
 153     xor4 |= rc4_key->S[idx] << 24;
 154
 155     out[k] = in[k] ^ xor4;
 156   }
 157
 158   return j;
 159 }
 160
 161 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
 162 {
 163   u32x a = digest[0];
 164   u32x b = digest[1];
 165   u32x c = digest[2];
 166   u32x d = digest[3];
 167
 168   u32x w0_t = w0[0];
 169   u32x w1_t = w0[1];
 170   u32x w2_t = w0[2];
 171   u32x w3_t = w0[3];
 172   u32x w4_t = w1[0];
 173   u32x w5_t = w1[1];
 174   u32x w6_t = w1[2];
 175   u32x w7_t = w1[3];
 176   u32x w8_t = w2[0];
 177   u32x w9_t = w2[1];
 178   u32x wa_t = w2[2];
 179   u32x wb_t = w2[3];
 180   u32x wc_t = w3[0];
 181   u32x wd_t = w3[1];
 182   u32x we_t = w3[2];
 183   u32x wf_t = w3[3];
 184
 185   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
 186   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
 187   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
 188   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
 189   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
 190   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
 191   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
 192   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
 193   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
 194   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
 195   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
 196   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
 197   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
 198   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
 199   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
 200   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
 201
 202   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
 203   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
 204   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
 205   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
 206   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
 207   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
 208   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
 209   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
 210   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
 211   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
 212   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
 213   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
 214   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
 215   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
 216   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 217   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 218
 219   MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
 220   MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
 221   MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
 222   MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
 223   MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
 224   MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
 225   MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
 226   MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
 227   MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
 228   MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
 229   MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
 230   MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
 231   MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
 232   MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
 233   MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
 234   MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
 235
 236   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 237   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 238   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 239   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 240   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 241   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 242   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 243   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 244   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 245   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 246   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 247   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 248   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 249   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 250   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 251   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 252
 253   digest[0] += a;
 254   digest[1] += b;
 255   digest[2] += c;
 256   digest[3] += d;
 257 }
 258
 259 __device__ __constant__ comb_t c_combs[1024];
 260
 261 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 262 {
 263   /**
 264    * modifier
 265    */
 266
 267   const u32 lid = threadIdx.x;
 268
 269   __shared__ RC4_KEY rc4_keys[64];
 270
 271   RC4_KEY *rc4_key = &rc4_keys[lid];
 272
 273   /**
 274    * base
 275    */
 276
 277   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 278
 279   if (gid >= gid_max) return;
 280
 281   u32x wordl0[4];
 282
 283   wordl0[0] = pws[gid].i[ 0];
 284   wordl0[1] = pws[gid].i[ 1];
 285   wordl0[2] = pws[gid].i[ 2];
 286   wordl0[3] = pws[gid].i[ 3];
 287
 288   u32x wordl1[4];
 289
 290   wordl1[0] = pws[gid].i[ 4];
 291   wordl1[1] = pws[gid].i[ 5];
 292   wordl1[2] = pws[gid].i[ 6];
 293   wordl1[3] = pws[gid].i[ 7];
 294
 295   u32x wordl2[4];
 296
 297   wordl2[0] = 0;
 298   wordl2[1] = 0;
 299   wordl2[2] = 0;
 300   wordl2[3] = 0;
 301
 302   u32x wordl3[4];
 303
 304   wordl3[0] = 0;
 305   wordl3[1] = 0;
 306   wordl3[2] = 0;
 307   wordl3[3] = 0;
 308
 309   const u32 pw_l_len = pws[gid].pw_len;
 310
 311   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 312   {
 313     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 314   }
 315
 316   /**
 317    * esalt
 318    */
 319
 320   const u32 version = oldoffice01_bufs[salt_pos].version;
 321
 322   u32 encryptedVerifier[4];
 323
 324   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 325   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 326   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 327   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 328
 329   /**
 330    * loop
 331    */
 332
 333   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 334   {
 335     u32 wordr0[4];
 336
 337     wordr0[0] = c_combs[il_pos].i[0];
 338     wordr0[1] = c_combs[il_pos].i[1];
 339     wordr0[2] = c_combs[il_pos].i[2];
 340     wordr0[3] = c_combs[il_pos].i[3];
 341
 342     u32 wordr1[4];
 343
 344     wordr1[0] = c_combs[il_pos].i[4];
 345     wordr1[1] = c_combs[il_pos].i[5];
 346     wordr1[2] = c_combs[il_pos].i[6];
 347     wordr1[3] = c_combs[il_pos].i[7];
 348
 349     u32 wordr2[4];
 350
 351     wordr2[0] = 0;
 352     wordr2[1] = 0;
 353     wordr2[2] = 0;
 354     wordr2[3] = 0;
 355
 356     u32 wordr3[4];
 357
 358     wordr3[0] = 0;
 359     wordr3[1] = 0;
 360     wordr3[2] = 0;
 361     wordr3[3] = 0;
 362
 363     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 364     {
 365       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 366     }
 367
 368     // first md5 to generate RC4 128 bit key
 369
 370     u32x w0_t[4];
 371     u32x w1_t[4];
 372     u32x w2_t[4];
 373     u32x w3_t[4];
 374
 375     w0_t[0]  =  wordl0[0] | wordr0[0];
 376     w0_t[1]  = (wordl0[1] | wordr0[1]) & 0xff;
 377     w0_t[2]  = 0x8000;
 378     w0_t[3]  = 0;
 379     w1_t[0]  = 0;
 380     w1_t[1]  = 0;
 381     w1_t[2]  = 0;
 382     w1_t[3]  = 0;
 383     w2_t[0]  = 0;
 384     w2_t[1]  = 0;
 385     w2_t[2]  = 0;
 386     w2_t[3]  = 0;
 387     w3_t[0]  = 0;
 388     w3_t[1]  = 0;
 389     w3_t[2]  = 9 * 8;
 390     w3_t[3]  = 0;
 391
 392     u32x digest[4];
 393
 394     digest[0] = MD5M_A;
 395     digest[1] = MD5M_B;
 396     digest[2] = MD5M_C;
 397     digest[3] = MD5M_D;
 398
 399     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 400
 401     // now the RC4 part
 402
 403     u32x key[4];
 404
 405     key[0] = digest[0];
 406     key[1] = digest[1];
 407     key[2] = digest[2];
 408     key[3] = digest[3];
 409
 410     rc4_init_16 (rc4_key, key);
 411
 412     u32x out[4];
 413
 414     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 415
 416     w0_t[0] = out[0];
 417     w0_t[1] = out[1];
 418     w0_t[2] = out[2];
 419     w0_t[3] = out[3];
 420     w1_t[0] = 0x80;
 421     w1_t[1] = 0;
 422     w1_t[2] = 0;
 423     w1_t[3] = 0;
 424     w2_t[0] = 0;
 425     w2_t[1] = 0;
 426     w2_t[2] = 0;
 427     w2_t[3] = 0;
 428     w3_t[0] = 0;
 429     w3_t[1] = 0;
 430     w3_t[2] = 16 * 8;
 431     w3_t[3] = 0;
 432
 433     digest[0] = MD5M_A;
 434     digest[1] = MD5M_B;
 435     digest[2] = MD5M_C;
 436     digest[3] = MD5M_D;
 437
 438     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 439
 440     rc4_next_16 (rc4_key, 16, j, digest, out);
 441
 442     const u32x r0 = out[0];
 443     const u32x r1 = out[1];
 444     const u32x r2 = out[2];
 445     const u32x r3 = out[3];
 446
 447     #include VECT_COMPARE_M
 448   }
 449 }
 450
 451 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 452 {
 453 }
 454
 455 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 456 {
 457 }
 458
 459 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 460 {
 461   /**
 462    * modifier
 463    */
 464
 465   const u32 lid = threadIdx.x;
 466
 467   __shared__ RC4_KEY rc4_keys[64];
 468
 469   RC4_KEY *rc4_key = &rc4_keys[lid];
 470
 471   /**
 472    * base
 473    */
 474
 475   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 476
 477   if (gid >= gid_max) return;
 478
 479   u32x wordl0[4];
 480
 481   wordl0[0] = pws[gid].i[ 0];
 482   wordl0[1] = pws[gid].i[ 1];
 483   wordl0[2] = pws[gid].i[ 2];
 484   wordl0[3] = pws[gid].i[ 3];
 485
 486   u32x wordl1[4];
 487
 488   wordl1[0] = pws[gid].i[ 4];
 489   wordl1[1] = pws[gid].i[ 5];
 490   wordl1[2] = pws[gid].i[ 6];
 491   wordl1[3] = pws[gid].i[ 7];
 492
 493   u32x wordl2[4];
 494
 495   wordl2[0] = 0;
 496   wordl2[1] = 0;
 497   wordl2[2] = 0;
 498   wordl2[3] = 0;
 499
 500   u32x wordl3[4];
 501
 502   wordl3[0] = 0;
 503   wordl3[1] = 0;
 504   wordl3[2] = 0;
 505   wordl3[3] = 0;
 506
 507   const u32 pw_l_len = pws[gid].pw_len;
 508
 509   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 510   {
 511     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 512   }
 513
 514   /**
 515    * digest
 516    */
 517
 518   const u32 search[4] =
 519   {
 520     digests_buf[digests_offset].digest_buf[DGST_R0],
 521     digests_buf[digests_offset].digest_buf[DGST_R1],
 522     digests_buf[digests_offset].digest_buf[DGST_R2],
 523     digests_buf[digests_offset].digest_buf[DGST_R3]
 524   };
 525
 526   /**
 527    * esalt
 528    */
 529
 530   const u32 version = oldoffice01_bufs[salt_pos].version;
 531
 532   u32 encryptedVerifier[4];
 533
 534   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 535   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 536   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 537   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 538
 539   /**
 540    * loop
 541    */
 542
 543   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 544   {
 545     u32 wordr0[4];
 546
 547     wordr0[0] = c_combs[il_pos].i[0];
 548     wordr0[1] = c_combs[il_pos].i[1];
 549     wordr0[2] = c_combs[il_pos].i[2];
 550     wordr0[3] = c_combs[il_pos].i[3];
 551
 552     u32 wordr1[4];
 553
 554     wordr1[0] = c_combs[il_pos].i[4];
 555     wordr1[1] = c_combs[il_pos].i[5];
 556     wordr1[2] = c_combs[il_pos].i[6];
 557     wordr1[3] = c_combs[il_pos].i[7];
 558
 559     u32 wordr2[4];
 560
 561     wordr2[0] = 0;
 562     wordr2[1] = 0;
 563     wordr2[2] = 0;
 564     wordr2[3] = 0;
 565
 566     u32 wordr3[4];
 567
 568     wordr3[0] = 0;
 569     wordr3[1] = 0;
 570     wordr3[2] = 0;
 571     wordr3[3] = 0;
 572
 573     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 574     {
 575       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 576     }
 577
 578     // first md5 to generate RC4 128 bit key
 579
 580     u32x w0_t[4];
 581     u32x w1_t[4];
 582     u32x w2_t[4];
 583     u32x w3_t[4];
 584
 585     w0_t[0]  =  wordl0[0] | wordr0[0];
 586     w0_t[1]  = (wordl0[1] | wordr0[1]) & 0xff;
 587     w0_t[2]  = 0x8000;
 588     w0_t[3]  = 0;
 589     w1_t[0]  = 0;
 590     w1_t[1]  = 0;
 591     w1_t[2]  = 0;
 592     w1_t[3]  = 0;
 593     w2_t[0]  = 0;
 594     w2_t[1]  = 0;
 595     w2_t[2]  = 0;
 596     w2_t[3]  = 0;
 597     w3_t[0]  = 0;
 598     w3_t[1]  = 0;
 599     w3_t[2]  = 9 * 8;
 600     w3_t[3]  = 0;
 601
 602     u32x digest[4];
 603
 604     digest[0] = MD5M_A;
 605     digest[1] = MD5M_B;
 606     digest[2] = MD5M_C;
 607     digest[3] = MD5M_D;
 608
 609     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 610
 611     // now the RC4 part
 612
 613     u32x key[4];
 614
 615     key[0] = digest[0];
 616     key[1] = digest[1];
 617     key[2] = digest[2];
 618     key[3] = digest[3];
 619
 620     rc4_init_16 (rc4_key, key);
 621
 622     u32x out[4];
 623
 624     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 625
 626     w0_t[0] = out[0];
 627     w0_t[1] = out[1];
 628     w0_t[2] = out[2];
 629     w0_t[3] = out[3];
 630     w1_t[0] = 0x80;
 631     w1_t[1] = 0;
 632     w1_t[2] = 0;
 633     w1_t[3] = 0;
 634     w2_t[0] = 0;
 635     w2_t[1] = 0;
 636     w2_t[2] = 0;
 637     w2_t[3] = 0;
 638     w3_t[0] = 0;
 639     w3_t[1] = 0;
 640     w3_t[2] = 16 * 8;
 641     w3_t[3] = 0;
 642
 643     digest[0] = MD5M_A;
 644     digest[1] = MD5M_B;
 645     digest[2] = MD5M_C;
 646     digest[3] = MD5M_D;
 647
 648     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 649
 650     rc4_next_16 (rc4_key, 16, j, digest, out);
 651
 652     const u32x r0 = out[0];
 653     const u32x r1 = out[1];
 654     const u32x r2 = out[2];
 655     const u32x r3 = out[3];
 656
 657     #include VECT_COMPARE_S
 658   }
 659 }
 660
 661 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 662 {
 663 }
 664
 665 extern "C" __global__ void __launch_bounds__ (64, 1) m09710_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 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 666 {
 667 }