nv/m09700_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__ static void gen336 (u32x digest_pre[4], u32 salt_buf[4], u32x digest[4])
 260 {
 261   u32x digest_t0[2];
 262   u32x digest_t1[2];
 263   u32x digest_t2[2];
 264   u32x digest_t3[2];
 265
 266   digest_t0[0] = digest_pre[0];
 267   digest_t0[1] = digest_pre[1] & 0xff;
 268
 269   digest_t1[0] =                       digest_pre[0] <<  8;
 270   digest_t1[1] = digest_pre[0] >> 24 | digest_pre[1] <<  8;
 271
 272   digest_t2[0] =                       digest_pre[0] << 16;
 273   digest_t2[1] = digest_pre[0] >> 16 | digest_pre[1] << 16;
 274
 275   digest_t3[0] =                       digest_pre[0] << 24;
 276   digest_t3[1] = digest_pre[0] >>  8 | digest_pre[1] << 24;
 277
 278   u32 salt_buf_t0[4];
 279   u32 salt_buf_t1[5];
 280   u32 salt_buf_t2[5];
 281   u32 salt_buf_t3[5];
 282
 283   salt_buf_t0[0] = salt_buf[0];
 284   salt_buf_t0[1] = salt_buf[1];
 285   salt_buf_t0[2] = salt_buf[2];
 286   salt_buf_t0[3] = salt_buf[3];
 287
 288   salt_buf_t1[0] =                     salt_buf[0] <<  8;
 289   salt_buf_t1[1] = salt_buf[0] >> 24 | salt_buf[1] <<  8;
 290   salt_buf_t1[2] = salt_buf[1] >> 24 | salt_buf[2] <<  8;
 291   salt_buf_t1[3] = salt_buf[2] >> 24 | salt_buf[3] <<  8;
 292   salt_buf_t1[4] = salt_buf[3] >> 24;
 293
 294   salt_buf_t2[0] =                     salt_buf[0] << 16;
 295   salt_buf_t2[1] = salt_buf[0] >> 16 | salt_buf[1] << 16;
 296   salt_buf_t2[2] = salt_buf[1] >> 16 | salt_buf[2] << 16;
 297   salt_buf_t2[3] = salt_buf[2] >> 16 | salt_buf[3] << 16;
 298   salt_buf_t2[4] = salt_buf[3] >> 16;
 299
 300   salt_buf_t3[0] =                     salt_buf[0] << 24;
 301   salt_buf_t3[1] = salt_buf[0] >>  8 | salt_buf[1] << 24;
 302   salt_buf_t3[2] = salt_buf[1] >>  8 | salt_buf[2] << 24;
 303   salt_buf_t3[3] = salt_buf[2] >>  8 | salt_buf[3] << 24;
 304   salt_buf_t3[4] = salt_buf[3] >>  8;
 305
 306   u32x w0_t[4];
 307   u32x w1_t[4];
 308   u32x w2_t[4];
 309   u32x w3_t[4];
 310
 311   // generate the 16 * 21 buffer
 312
 313   w0_t[0] = 0;
 314   w0_t[1] = 0;
 315   w0_t[2] = 0;
 316   w0_t[3] = 0;
 317   w1_t[0] = 0;
 318   w1_t[1] = 0;
 319   w1_t[2] = 0;
 320   w1_t[3] = 0;
 321   w2_t[0] = 0;
 322   w2_t[1] = 0;
 323   w2_t[2] = 0;
 324   w2_t[3] = 0;
 325   w3_t[0] = 0;
 326   w3_t[1] = 0;
 327   w3_t[2] = 0;
 328   w3_t[3] = 0;
 329
 330   // 0..5
 331   w0_t[0]  = digest_t0[0];
 332   w0_t[1]  = digest_t0[1];
 333
 334   // 5..21
 335   w0_t[1] |= salt_buf_t1[0];
 336   w0_t[2]  = salt_buf_t1[1];
 337   w0_t[3]  = salt_buf_t1[2];
 338   w1_t[0]  = salt_buf_t1[3];
 339   w1_t[1]  = salt_buf_t1[4];
 340
 341   // 21..26
 342   w1_t[1] |= digest_t1[0];
 343   w1_t[2]  = digest_t1[1];
 344
 345   // 26..42
 346   w1_t[2] |= salt_buf_t2[0];
 347   w1_t[3]  = salt_buf_t2[1];
 348   w2_t[0]  = salt_buf_t2[2];
 349   w2_t[1]  = salt_buf_t2[3];
 350   w2_t[2]  = salt_buf_t2[4];
 351
 352   // 42..47
 353   w2_t[2] |= digest_t2[0];
 354   w2_t[3]  = digest_t2[1];
 355
 356   // 47..63
 357   w2_t[3] |= salt_buf_t3[0];
 358   w3_t[0]  = salt_buf_t3[1];
 359   w3_t[1]  = salt_buf_t3[2];
 360   w3_t[2]  = salt_buf_t3[3];
 361   w3_t[3]  = salt_buf_t3[4];
 362
 363   // 63..
 364
 365   w3_t[3] |= digest_t3[0];
 366
 367   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 368
 369   w0_t[0] = 0;
 370   w0_t[1] = 0;
 371   w0_t[2] = 0;
 372   w0_t[3] = 0;
 373   w1_t[0] = 0;
 374   w1_t[1] = 0;
 375   w1_t[2] = 0;
 376   w1_t[3] = 0;
 377   w2_t[0] = 0;
 378   w2_t[1] = 0;
 379   w2_t[2] = 0;
 380   w2_t[3] = 0;
 381   w3_t[0] = 0;
 382   w3_t[1] = 0;
 383   w3_t[2] = 0;
 384   w3_t[3] = 0;
 385
 386   // 0..4
 387   w0_t[0]  = digest_t3[1];
 388
 389   // 4..20
 390   w0_t[1]  = salt_buf_t0[0];
 391   w0_t[2]  = salt_buf_t0[1];
 392   w0_t[3]  = salt_buf_t0[2];
 393   w1_t[0]  = salt_buf_t0[3];
 394
 395   // 20..25
 396   w1_t[1]  = digest_t0[0];
 397   w1_t[2]  = digest_t0[1];
 398
 399   // 25..41
 400   w1_t[2] |= salt_buf_t1[0];
 401   w1_t[3]  = salt_buf_t1[1];
 402   w2_t[0]  = salt_buf_t1[2];
 403   w2_t[1]  = salt_buf_t1[3];
 404   w2_t[2]  = salt_buf_t1[4];
 405
 406   // 41..46
 407   w2_t[2] |= digest_t1[0];
 408   w2_t[3]  = digest_t1[1];
 409
 410   // 46..62
 411   w2_t[3] |= salt_buf_t2[0];
 412   w3_t[0]  = salt_buf_t2[1];
 413   w3_t[1]  = salt_buf_t2[2];
 414   w3_t[2]  = salt_buf_t2[3];
 415   w3_t[3]  = salt_buf_t2[4];
 416
 417   // 62..
 418   w3_t[3] |= digest_t2[0];
 419
 420   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 421
 422   w0_t[0] = 0;
 423   w0_t[1] = 0;
 424   w0_t[2] = 0;
 425   w0_t[3] = 0;
 426   w1_t[0] = 0;
 427   w1_t[1] = 0;
 428   w1_t[2] = 0;
 429   w1_t[3] = 0;
 430   w2_t[0] = 0;
 431   w2_t[1] = 0;
 432   w2_t[2] = 0;
 433   w2_t[3] = 0;
 434   w3_t[0] = 0;
 435   w3_t[1] = 0;
 436   w3_t[2] = 0;
 437   w3_t[3] = 0;
 438
 439   // 0..3
 440   w0_t[0]  = digest_t2[1];
 441
 442   // 3..19
 443   w0_t[0] |= salt_buf_t3[0];
 444   w0_t[1]  = salt_buf_t3[1];
 445   w0_t[2]  = salt_buf_t3[2];
 446   w0_t[3]  = salt_buf_t3[3];
 447   w1_t[0]  = salt_buf_t3[4];
 448
 449   // 19..24
 450   w1_t[0] |= digest_t3[0];
 451   w1_t[1]  = digest_t3[1];
 452
 453   // 24..40
 454   w1_t[2]  = salt_buf_t0[0];
 455   w1_t[3]  = salt_buf_t0[1];
 456   w2_t[0]  = salt_buf_t0[2];
 457   w2_t[1]  = salt_buf_t0[3];
 458
 459   // 40..45
 460   w2_t[2]  = digest_t0[0];
 461   w2_t[3]  = digest_t0[1];
 462
 463   // 45..61
 464   w2_t[3] |= salt_buf_t1[0];
 465   w3_t[0]  = salt_buf_t1[1];
 466   w3_t[1]  = salt_buf_t1[2];
 467   w3_t[2]  = salt_buf_t1[3];
 468   w3_t[3]  = salt_buf_t1[4];
 469
 470   // 61..
 471   w3_t[3] |= digest_t1[0];
 472
 473   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 474
 475   w0_t[0] = 0;
 476   w0_t[1] = 0;
 477   w0_t[2] = 0;
 478   w0_t[3] = 0;
 479   w1_t[0] = 0;
 480   w1_t[1] = 0;
 481   w1_t[2] = 0;
 482   w1_t[3] = 0;
 483   w2_t[0] = 0;
 484   w2_t[1] = 0;
 485   w2_t[2] = 0;
 486   w2_t[3] = 0;
 487   w3_t[0] = 0;
 488   w3_t[1] = 0;
 489   w3_t[2] = 0;
 490   w3_t[3] = 0;
 491
 492   // 0..2
 493   w0_t[0]  = digest_t1[1];
 494
 495   // 2..18
 496   w0_t[0] |= salt_buf_t2[0];
 497   w0_t[1]  = salt_buf_t2[1];
 498   w0_t[2]  = salt_buf_t2[2];
 499   w0_t[3]  = salt_buf_t2[3];
 500   w1_t[0]  = salt_buf_t2[4];
 501
 502   // 18..23
 503   w1_t[0] |= digest_t2[0];
 504   w1_t[1]  = digest_t2[1];
 505
 506   // 23..39
 507   w1_t[1] |= salt_buf_t3[0];
 508   w1_t[2]  = salt_buf_t3[1];
 509   w1_t[3]  = salt_buf_t3[2];
 510   w2_t[0]  = salt_buf_t3[3];
 511   w2_t[1]  = salt_buf_t3[4];
 512
 513   // 39..44
 514   w2_t[1] |= digest_t3[0];
 515   w2_t[2]  = digest_t3[1];
 516
 517   // 44..60
 518   w2_t[3]  = salt_buf_t0[0];
 519   w3_t[0]  = salt_buf_t0[1];
 520   w3_t[1]  = salt_buf_t0[2];
 521   w3_t[2]  = salt_buf_t0[3];
 522
 523   // 60..
 524   w3_t[3]  = digest_t0[0];
 525
 526   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 527
 528   w0_t[0] = 0;
 529   w0_t[1] = 0;
 530   w0_t[2] = 0;
 531   w0_t[3] = 0;
 532   w1_t[0] = 0;
 533   w1_t[1] = 0;
 534   w1_t[2] = 0;
 535   w1_t[3] = 0;
 536   w2_t[0] = 0;
 537   w2_t[1] = 0;
 538   w2_t[2] = 0;
 539   w2_t[3] = 0;
 540   w3_t[0] = 0;
 541   w3_t[1] = 0;
 542   w3_t[2] = 0;
 543   w3_t[3] = 0;
 544
 545   // 0..1
 546   w0_t[0]  = digest_t0[1];
 547
 548   // 1..17
 549   w0_t[0] |= salt_buf_t1[0];
 550   w0_t[1]  = salt_buf_t1[1];
 551   w0_t[2]  = salt_buf_t1[2];
 552   w0_t[3]  = salt_buf_t1[3];
 553   w1_t[0]  = salt_buf_t1[4];
 554
 555   // 17..22
 556   w1_t[0] |= digest_t1[0];
 557   w1_t[1]  = digest_t1[1];
 558
 559   // 22..38
 560   w1_t[1] |= salt_buf_t2[0];
 561   w1_t[2]  = salt_buf_t2[1];
 562   w1_t[3]  = salt_buf_t2[2];
 563   w2_t[0]  = salt_buf_t2[3];
 564   w2_t[1]  = salt_buf_t2[4];
 565
 566   // 38..43
 567   w2_t[1] |= digest_t2[0];
 568   w2_t[2]  = digest_t2[1];
 569
 570   // 43..59
 571   w2_t[2] |= salt_buf_t3[0];
 572   w2_t[3]  = salt_buf_t3[1];
 573   w3_t[0]  = salt_buf_t3[2];
 574   w3_t[1]  = salt_buf_t3[3];
 575   w3_t[2]  = salt_buf_t3[4];
 576
 577   // 59..
 578   w3_t[2] |= digest_t3[0];
 579   w3_t[3]  = digest_t3[1];
 580
 581   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 582
 583   w0_t[0]  = salt_buf_t0[0];
 584   w0_t[1]  = salt_buf_t0[1];
 585   w0_t[2]  = salt_buf_t0[2];
 586   w0_t[3]  = salt_buf_t0[3];
 587   w1_t[0]  = 0x80;
 588   w1_t[1]  = 0;
 589   w1_t[2]  = 0;
 590   w1_t[3]  = 0;
 591   w2_t[0]  = 0;
 592   w2_t[1]  = 0;
 593   w2_t[2]  = 0;
 594   w2_t[3]  = 0;
 595   w3_t[0]  = 0;
 596   w3_t[1]  = 0;
 597   w3_t[2]  = 21 * 16 * 8;
 598   w3_t[3]  = 0;
 599
 600   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 601 }
 602
 603 __device__ __constant__ comb_t c_combs[1024];
 604
 605 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 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)
 606 {
 607   /**
 608    * modifier
 609    */
 610
 611   const u32 lid = threadIdx.x;
 612
 613   __shared__ RC4_KEY rc4_keys[64];
 614
 615   RC4_KEY *rc4_key = &rc4_keys[lid];
 616
 617   /**
 618    * base
 619    */
 620
 621   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 622
 623   if (gid >= gid_max) return;
 624
 625   u32x wordl0[4];
 626
 627   wordl0[0] = pws[gid].i[ 0];
 628   wordl0[1] = pws[gid].i[ 1];
 629   wordl0[2] = pws[gid].i[ 2];
 630   wordl0[3] = pws[gid].i[ 3];
 631
 632   u32x wordl1[4];
 633
 634   wordl1[0] = pws[gid].i[ 4];
 635   wordl1[1] = pws[gid].i[ 5];
 636   wordl1[2] = pws[gid].i[ 6];
 637   wordl1[3] = pws[gid].i[ 7];
 638
 639   u32x wordl2[4];
 640
 641   wordl2[0] = 0;
 642   wordl2[1] = 0;
 643   wordl2[2] = 0;
 644   wordl2[3] = 0;
 645
 646   u32x wordl3[4];
 647
 648   wordl3[0] = 0;
 649   wordl3[1] = 0;
 650   wordl3[2] = 0;
 651   wordl3[3] = 0;
 652
 653   const u32 pw_l_len = pws[gid].pw_len;
 654
 655   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 656   {
 657     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 658   }
 659
 660   /**
 661    * digest
 662    */
 663
 664   const u32 search[4] =
 665   {
 666     digests_buf[digests_offset].digest_buf[DGST_R0],
 667     digests_buf[digests_offset].digest_buf[DGST_R1],
 668     digests_buf[digests_offset].digest_buf[DGST_R2],
 669     digests_buf[digests_offset].digest_buf[DGST_R3]
 670   };
 671
 672   /**
 673    * salt
 674    */
 675
 676   u32 salt_buf[4];
 677
 678   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 679   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 680   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 681   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 682
 683   /**
 684    * esalt
 685    */
 686
 687   const u32 version = oldoffice01_bufs[salt_pos].version;
 688
 689   u32 encryptedVerifier[4];
 690
 691   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 692   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 693   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 694   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 695
 696   /**
 697    * loop
 698    */
 699
 700   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 701   {
 702     const u32 pw_r_len = c_combs[il_pos].pw_len;
 703
 704     const u32 pw_len = pw_l_len + pw_r_len;
 705
 706     u32 wordr0[4];
 707
 708     wordr0[0] = c_combs[il_pos].i[0];
 709     wordr0[1] = c_combs[il_pos].i[1];
 710     wordr0[2] = c_combs[il_pos].i[2];
 711     wordr0[3] = c_combs[il_pos].i[3];
 712
 713     u32 wordr1[4];
 714
 715     wordr1[0] = c_combs[il_pos].i[4];
 716     wordr1[1] = c_combs[il_pos].i[5];
 717     wordr1[2] = c_combs[il_pos].i[6];
 718     wordr1[3] = c_combs[il_pos].i[7];
 719
 720     u32 wordr2[4];
 721
 722     wordr2[0] = 0;
 723     wordr2[1] = 0;
 724     wordr2[2] = 0;
 725     wordr2[3] = 0;
 726
 727     u32 wordr3[4];
 728
 729     wordr3[0] = 0;
 730     wordr3[1] = 0;
 731     wordr3[2] = 0;
 732     wordr3[3] = 0;
 733
 734     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 735     {
 736       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 737     }
 738
 739     u32x w0[4];
 740
 741     w0[0] = wordl0[0] | wordr0[0];
 742     w0[1] = wordl0[1] | wordr0[1];
 743     w0[2] = wordl0[2] | wordr0[2];
 744     w0[3] = wordl0[3] | wordr0[3];
 745
 746     u32x w1[4];
 747
 748     w1[0] = wordl1[0] | wordr1[0];
 749     w1[1] = wordl1[1] | wordr1[1];
 750     w1[2] = wordl1[2] | wordr1[2];
 751     w1[3] = wordl1[3] | wordr1[3];
 752
 753     u32x w2[4];
 754
 755     w2[0] = wordl2[0] | wordr2[0];
 756     w2[1] = wordl2[1] | wordr2[1];
 757     w2[2] = wordl2[2] | wordr2[2];
 758     w2[3] = wordl2[3] | wordr2[3];
 759
 760     u32x w3[4];
 761
 762     w3[0] = wordl3[0] | wordr3[0];
 763     w3[1] = wordl3[1] | wordr3[1];
 764     w3[2] = 0;
 765     w3[3] = 0;
 766
 767     append_0x80_2 (w0, w1, pw_len);
 768
 769     u32x w0_t[4];
 770     u32x w1_t[4];
 771     u32x w2_t[4];
 772     u32x w3_t[4];
 773
 774     make_unicode (w0, w0_t, w1_t);
 775     make_unicode (w1, w2_t, w3_t);
 776
 777     w3_t[2] = pw_len * 8 * 2;
 778
 779     u32x digest_pre[4];
 780
 781     digest_pre[0] = MD5M_A;
 782     digest_pre[1] = MD5M_B;
 783     digest_pre[2] = MD5M_C;
 784     digest_pre[3] = MD5M_D;
 785
 786     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_pre);
 787
 788     digest_pre[0] &= 0xffffffff;
 789     digest_pre[1] &= 0x000000ff;
 790     digest_pre[2] &= 0x00000000;
 791     digest_pre[3] &= 0x00000000;
 792
 793     u32x digest[4];
 794
 795     digest[0] = MD5M_A;
 796     digest[1] = MD5M_B;
 797     digest[2] = MD5M_C;
 798     digest[3] = MD5M_D;
 799
 800     gen336 (digest_pre, salt_buf, digest);
 801
 802     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
 803
 804     w0_t[0]  = digest[0];
 805     w0_t[1]  = digest[1] & 0xff;
 806     w0_t[2]  = 0x8000;
 807     w0_t[3]  = 0;
 808     w1_t[0]  = 0;
 809     w1_t[1]  = 0;
 810     w1_t[2]  = 0;
 811     w1_t[3]  = 0;
 812     w2_t[0]  = 0;
 813     w2_t[1]  = 0;
 814     w2_t[2]  = 0;
 815     w2_t[3]  = 0;
 816     w3_t[0]  = 0;
 817     w3_t[1]  = 0;
 818     w3_t[2]  = 9 * 8;
 819     w3_t[3]  = 0;
 820
 821     digest[0] = MD5M_A;
 822     digest[1] = MD5M_B;
 823     digest[2] = MD5M_C;
 824     digest[3] = MD5M_D;
 825
 826     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 827
 828     // now the RC4 part
 829
 830     u32x key[4];
 831
 832     key[0] = digest[0];
 833     key[1] = digest[1];
 834     key[2] = digest[2];
 835     key[3] = digest[3];
 836
 837     rc4_init_16 (rc4_key, key);
 838
 839     u32x out[4];
 840
 841     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 842
 843     w0_t[0] = out[0];
 844     w0_t[1] = out[1];
 845     w0_t[2] = out[2];
 846     w0_t[3] = out[3];
 847     w1_t[0] = 0x80;
 848     w1_t[1] = 0;
 849     w1_t[2] = 0;
 850     w1_t[3] = 0;
 851     w2_t[0] = 0;
 852     w2_t[1] = 0;
 853     w2_t[2] = 0;
 854     w2_t[3] = 0;
 855     w3_t[0] = 0;
 856     w3_t[1] = 0;
 857     w3_t[2] = 16 * 8;
 858     w3_t[3] = 0;
 859
 860     digest[0] = MD5M_A;
 861     digest[1] = MD5M_B;
 862     digest[2] = MD5M_C;
 863     digest[3] = MD5M_D;
 864
 865     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 866
 867     rc4_next_16 (rc4_key, 16, j, digest, out);
 868
 869     const u32x r0 = out[0];
 870     const u32x r1 = out[1];
 871     const u32x r2 = out[2];
 872     const u32x r3 = out[3];
 873
 874     #include VECT_COMPARE_M
 875   }
 876 }
 877
 878 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 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)
 879 {
 880 }
 881
 882 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 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)
 883 {
 884 }
 885
 886 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 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)
 887 {
 888   /**
 889    * modifier
 890    */
 891
 892   const u32 lid = threadIdx.x;
 893
 894   __shared__ RC4_KEY rc4_keys[64];
 895
 896   RC4_KEY *rc4_key = &rc4_keys[lid];
 897
 898   /**
 899    * base
 900    */
 901
 902   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 903
 904   if (gid >= gid_max) return;
 905
 906   u32x wordl0[4];
 907
 908   wordl0[0] = pws[gid].i[ 0];
 909   wordl0[1] = pws[gid].i[ 1];
 910   wordl0[2] = pws[gid].i[ 2];
 911   wordl0[3] = pws[gid].i[ 3];
 912
 913   u32x wordl1[4];
 914
 915   wordl1[0] = pws[gid].i[ 4];
 916   wordl1[1] = pws[gid].i[ 5];
 917   wordl1[2] = pws[gid].i[ 6];
 918   wordl1[3] = pws[gid].i[ 7];
 919
 920   u32x wordl2[4];
 921
 922   wordl2[0] = 0;
 923   wordl2[1] = 0;
 924   wordl2[2] = 0;
 925   wordl2[3] = 0;
 926
 927   u32x wordl3[4];
 928
 929   wordl3[0] = 0;
 930   wordl3[1] = 0;
 931   wordl3[2] = 0;
 932   wordl3[3] = 0;
 933
 934   const u32 pw_l_len = pws[gid].pw_len;
 935
 936   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 937   {
 938     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 939   }
 940
 941   /**
 942    * digest
 943    */
 944
 945   const u32 search[4] =
 946   {
 947     digests_buf[digests_offset].digest_buf[DGST_R0],
 948     digests_buf[digests_offset].digest_buf[DGST_R1],
 949     digests_buf[digests_offset].digest_buf[DGST_R2],
 950     digests_buf[digests_offset].digest_buf[DGST_R3]
 951   };
 952
 953   /**
 954    * salt
 955    */
 956
 957   u32 salt_buf[4];
 958
 959   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 960   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 961   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 962   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 963
 964   /**
 965    * esalt
 966    */
 967
 968   const u32 version = oldoffice01_bufs[salt_pos].version;
 969
 970   u32 encryptedVerifier[4];
 971
 972   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 973   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 974   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 975   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 976
 977   /**
 978    * loop
 979    */
 980
 981   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 982   {
 983     const u32 pw_r_len = c_combs[il_pos].pw_len;
 984
 985     const u32 pw_len = pw_l_len + pw_r_len;
 986
 987     u32 wordr0[4];
 988
 989     wordr0[0] = c_combs[il_pos].i[0];
 990     wordr0[1] = c_combs[il_pos].i[1];
 991     wordr0[2] = c_combs[il_pos].i[2];
 992     wordr0[3] = c_combs[il_pos].i[3];
 993
 994     u32 wordr1[4];
 995
 996     wordr1[0] = c_combs[il_pos].i[4];
 997     wordr1[1] = c_combs[il_pos].i[5];
 998     wordr1[2] = c_combs[il_pos].i[6];
 999     wordr1[3] = c_combs[il_pos].i[7];
1000
1001     u32 wordr2[4];
1002
1003     wordr2[0] = 0;
1004     wordr2[1] = 0;
1005     wordr2[2] = 0;
1006     wordr2[3] = 0;
1007
1008     u32 wordr3[4];
1009
1010     wordr3[0] = 0;
1011     wordr3[1] = 0;
1012     wordr3[2] = 0;
1013     wordr3[3] = 0;
1014
1015     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
1016     {
1017       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
1018     }
1019
1020     u32x w0[4];
1021
1022     w0[0] = wordl0[0] | wordr0[0];
1023     w0[1] = wordl0[1] | wordr0[1];
1024     w0[2] = wordl0[2] | wordr0[2];
1025     w0[3] = wordl0[3] | wordr0[3];
1026
1027     u32x w1[4];
1028
1029     w1[0] = wordl1[0] | wordr1[0];
1030     w1[1] = wordl1[1] | wordr1[1];
1031     w1[2] = wordl1[2] | wordr1[2];
1032     w1[3] = wordl1[3] | wordr1[3];
1033
1034     u32x w2[4];
1035
1036     w2[0] = wordl2[0] | wordr2[0];
1037     w2[1] = wordl2[1] | wordr2[1];
1038     w2[2] = wordl2[2] | wordr2[2];
1039     w2[3] = wordl2[3] | wordr2[3];
1040
1041     u32x w3[4];
1042
1043     w3[0] = wordl3[0] | wordr3[0];
1044     w3[1] = wordl3[1] | wordr3[1];
1045     w3[2] = 0;
1046     w3[3] = 0;
1047
1048     append_0x80_2 (w0, w1, pw_len);
1049
1050     u32x w0_t[4];
1051     u32x w1_t[4];
1052     u32x w2_t[4];
1053     u32x w3_t[4];
1054
1055     make_unicode (w0, w0_t, w1_t);
1056     make_unicode (w1, w2_t, w3_t);
1057
1058     w3_t[2] = pw_len * 8 * 2;
1059
1060     u32x digest_pre[4];
1061
1062     digest_pre[0] = MD5M_A;
1063     digest_pre[1] = MD5M_B;
1064     digest_pre[2] = MD5M_C;
1065     digest_pre[3] = MD5M_D;
1066
1067     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_pre);
1068
1069     digest_pre[0] &= 0xffffffff;
1070     digest_pre[1] &= 0x000000ff;
1071     digest_pre[2] &= 0x00000000;
1072     digest_pre[3] &= 0x00000000;
1073
1074     u32x digest[4];
1075
1076     digest[0] = MD5M_A;
1077     digest[1] = MD5M_B;
1078     digest[2] = MD5M_C;
1079     digest[3] = MD5M_D;
1080
1081     gen336 (digest_pre, salt_buf, digest);
1082
1083     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
1084
1085     w0_t[0]  = digest[0];
1086     w0_t[1]  = digest[1] & 0xff;
1087     w0_t[2]  = 0x8000;
1088     w0_t[3]  = 0;
1089     w1_t[0]  = 0;
1090     w1_t[1]  = 0;
1091     w1_t[2]  = 0;
1092     w1_t[3]  = 0;
1093     w2_t[0]  = 0;
1094     w2_t[1]  = 0;
1095     w2_t[2]  = 0;
1096     w2_t[3]  = 0;
1097     w3_t[0]  = 0;
1098     w3_t[1]  = 0;
1099     w3_t[2]  = 9 * 8;
1100     w3_t[3]  = 0;
1101
1102     digest[0] = MD5M_A;
1103     digest[1] = MD5M_B;
1104     digest[2] = MD5M_C;
1105     digest[3] = MD5M_D;
1106
1107     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1108
1109     // now the RC4 part
1110
1111     u32x key[4];
1112
1113     key[0] = digest[0];
1114     key[1] = digest[1];
1115     key[2] = digest[2];
1116     key[3] = digest[3];
1117
1118     rc4_init_16 (rc4_key, key);
1119
1120     u32x out[4];
1121
1122     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
1123
1124     w0_t[0] = out[0];
1125     w0_t[1] = out[1];
1126     w0_t[2] = out[2];
1127     w0_t[3] = out[3];
1128     w1_t[0] = 0x80;
1129     w1_t[1] = 0;
1130     w1_t[2] = 0;
1131     w1_t[3] = 0;
1132     w2_t[0] = 0;
1133     w2_t[1] = 0;
1134     w2_t[2] = 0;
1135     w2_t[3] = 0;
1136     w3_t[0] = 0;
1137     w3_t[1] = 0;
1138     w3_t[2] = 16 * 8;
1139     w3_t[3] = 0;
1140
1141     digest[0] = MD5M_A;
1142     digest[1] = MD5M_B;
1143     digest[2] = MD5M_C;
1144     digest[3] = MD5M_D;
1145
1146     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1147
1148     rc4_next_16 (rc4_key, 16, j, digest, out);
1149
1150     const u32x r0 = out[0];
1151     const u32x r1 = out[1];
1152     const u32x r2 = out[2];
1153     const u32x r3 = out[3];
1154
1155     #include VECT_COMPARE_S
1156   }
1157 }
1158
1159 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 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)
1160 {
1161 }
1162
1163 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 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)
1164 {
1165 }