nv/m09700_a0.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 #include "include/rp_gpu.h"
  28 #include "rp_nv.c"
  29
  30 #ifdef  VECT_SIZE1
  31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  33 #endif
  34
  35 #ifdef  VECT_SIZE2
  36 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  37 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  38 #endif
  39
  40 #ifdef  VECT_SIZE4
  41 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  42 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  43 #endif
  44
  45 typedef struct
  46 {
  47   u8 S[256];
  48
  49   u32 wtf_its_faster;
  50
  51 } RC4_KEY;
  52
  53 __device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
  54 {
  55   u8 tmp;
  56
  57   tmp           = rc4_key->S[i];
  58   rc4_key->S[i] = rc4_key->S[j];
  59   rc4_key->S[j] = tmp;
  60 }
  61
  62 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  63 {
  64   u32 v = 0x03020100;
  65   u32 a = 0x04040404;
  66
  67   u32 *ptr = (u32 *) rc4_key->S;
  68
  69   #pragma unroll 64
  70   for (u32 i = 0; i < 64; i++)
  71   {
  72     *ptr++ = v; v += a;
  73   }
  74
  75   u32 j = 0;
  76
  77   for (u32 i = 0; i < 16; i++)
  78   {
  79     u32 idx = i * 16;
  80
  81     u32 v;
  82
  83     v = data[0];
  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[1];
  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[2];
  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     v = data[3];
 105
 106     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
 107     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
 108     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 109     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 110   }
 111 }
 112
 113 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 114 {
 115   for (u32 k = 0; k < 4; k++)
 116   {
 117     u32 xor4 = 0;
 118
 119     u8 idx;
 120
 121     i += 1;
 122     j += rc4_key->S[i];
 123
 124     swap (rc4_key, i, j);
 125
 126     idx = rc4_key->S[i] + rc4_key->S[j];
 127
 128     xor4 |= rc4_key->S[idx] <<  0;
 129
 130     i += 1;
 131     j += rc4_key->S[i];
 132
 133     swap (rc4_key, i, j);
 134
 135     idx = rc4_key->S[i] + rc4_key->S[j];
 136
 137     xor4 |= rc4_key->S[idx] <<  8;
 138
 139     i += 1;
 140     j += rc4_key->S[i];
 141
 142     swap (rc4_key, i, j);
 143
 144     idx = rc4_key->S[i] + rc4_key->S[j];
 145
 146     xor4 |= rc4_key->S[idx] << 16;
 147
 148     i += 1;
 149     j += rc4_key->S[i];
 150
 151     swap (rc4_key, i, j);
 152
 153     idx = rc4_key->S[i] + rc4_key->S[j];
 154
 155     xor4 |= rc4_key->S[idx] << 24;
 156
 157     out[k] = in[k] ^ xor4;
 158   }
 159
 160   return j;
 161 }
 162
 163 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
 164 {
 165   u32x a = digest[0];
 166   u32x b = digest[1];
 167   u32x c = digest[2];
 168   u32x d = digest[3];
 169
 170   u32x w0_t = w0[0];
 171   u32x w1_t = w0[1];
 172   u32x w2_t = w0[2];
 173   u32x w3_t = w0[3];
 174   u32x w4_t = w1[0];
 175   u32x w5_t = w1[1];
 176   u32x w6_t = w1[2];
 177   u32x w7_t = w1[3];
 178   u32x w8_t = w2[0];
 179   u32x w9_t = w2[1];
 180   u32x wa_t = w2[2];
 181   u32x wb_t = w2[3];
 182   u32x wc_t = w3[0];
 183   u32x wd_t = w3[1];
 184   u32x we_t = w3[2];
 185   u32x wf_t = w3[3];
 186
 187   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
 188   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
 189   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
 190   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
 191   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
 192   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
 193   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
 194   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
 195   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
 196   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
 197   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
 198   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
 199   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
 200   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
 201   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
 202   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
 203
 204   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
 205   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
 206   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
 207   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
 208   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
 209   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
 210   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
 211   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
 212   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
 213   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
 214   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
 215   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
 216   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
 217   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
 218   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 219   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 220
 221   MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
 222   MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
 223   MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
 224   MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
 225   MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
 226   MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
 227   MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
 228   MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
 229   MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
 230   MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
 231   MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
 232   MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
 233   MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
 234   MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
 235   MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
 236   MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
 237
 238   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 239   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 240   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 241   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 242   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 243   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 244   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 245   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 246   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 247   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 248   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 249   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 250   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 251   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 252   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 253   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 254
 255   digest[0] += a;
 256   digest[1] += b;
 257   digest[2] += c;
 258   digest[3] += d;
 259 }
 260
 261 __device__ static void gen336 (u32x digest_pre[4], u32 salt_buf[4], u32x digest[4])
 262 {
 263   u32x digest_t0[2];
 264   u32x digest_t1[2];
 265   u32x digest_t2[2];
 266   u32x digest_t3[2];
 267
 268   digest_t0[0] = digest_pre[0];
 269   digest_t0[1] = digest_pre[1] & 0xff;
 270
 271   digest_t1[0] =                       digest_pre[0] <<  8;
 272   digest_t1[1] = digest_pre[0] >> 24 | digest_pre[1] <<  8;
 273
 274   digest_t2[0] =                       digest_pre[0] << 16;
 275   digest_t2[1] = digest_pre[0] >> 16 | digest_pre[1] << 16;
 276
 277   digest_t3[0] =                       digest_pre[0] << 24;
 278   digest_t3[1] = digest_pre[0] >>  8 | digest_pre[1] << 24;
 279
 280   u32 salt_buf_t0[4];
 281   u32 salt_buf_t1[5];
 282   u32 salt_buf_t2[5];
 283   u32 salt_buf_t3[5];
 284
 285   salt_buf_t0[0] = salt_buf[0];
 286   salt_buf_t0[1] = salt_buf[1];
 287   salt_buf_t0[2] = salt_buf[2];
 288   salt_buf_t0[3] = salt_buf[3];
 289
 290   salt_buf_t1[0] =                     salt_buf[0] <<  8;
 291   salt_buf_t1[1] = salt_buf[0] >> 24 | salt_buf[1] <<  8;
 292   salt_buf_t1[2] = salt_buf[1] >> 24 | salt_buf[2] <<  8;
 293   salt_buf_t1[3] = salt_buf[2] >> 24 | salt_buf[3] <<  8;
 294   salt_buf_t1[4] = salt_buf[3] >> 24;
 295
 296   salt_buf_t2[0] =                     salt_buf[0] << 16;
 297   salt_buf_t2[1] = salt_buf[0] >> 16 | salt_buf[1] << 16;
 298   salt_buf_t2[2] = salt_buf[1] >> 16 | salt_buf[2] << 16;
 299   salt_buf_t2[3] = salt_buf[2] >> 16 | salt_buf[3] << 16;
 300   salt_buf_t2[4] = salt_buf[3] >> 16;
 301
 302   salt_buf_t3[0] =                     salt_buf[0] << 24;
 303   salt_buf_t3[1] = salt_buf[0] >>  8 | salt_buf[1] << 24;
 304   salt_buf_t3[2] = salt_buf[1] >>  8 | salt_buf[2] << 24;
 305   salt_buf_t3[3] = salt_buf[2] >>  8 | salt_buf[3] << 24;
 306   salt_buf_t3[4] = salt_buf[3] >>  8;
 307
 308   u32x w0_t[4];
 309   u32x w1_t[4];
 310   u32x w2_t[4];
 311   u32x w3_t[4];
 312
 313   // generate the 16 * 21 buffer
 314
 315   w0_t[0] = 0;
 316   w0_t[1] = 0;
 317   w0_t[2] = 0;
 318   w0_t[3] = 0;
 319   w1_t[0] = 0;
 320   w1_t[1] = 0;
 321   w1_t[2] = 0;
 322   w1_t[3] = 0;
 323   w2_t[0] = 0;
 324   w2_t[1] = 0;
 325   w2_t[2] = 0;
 326   w2_t[3] = 0;
 327   w3_t[0] = 0;
 328   w3_t[1] = 0;
 329   w3_t[2] = 0;
 330   w3_t[3] = 0;
 331
 332   // 0..5
 333   w0_t[0]  = digest_t0[0];
 334   w0_t[1]  = digest_t0[1];
 335
 336   // 5..21
 337   w0_t[1] |= salt_buf_t1[0];
 338   w0_t[2]  = salt_buf_t1[1];
 339   w0_t[3]  = salt_buf_t1[2];
 340   w1_t[0]  = salt_buf_t1[3];
 341   w1_t[1]  = salt_buf_t1[4];
 342
 343   // 21..26
 344   w1_t[1] |= digest_t1[0];
 345   w1_t[2]  = digest_t1[1];
 346
 347   // 26..42
 348   w1_t[2] |= salt_buf_t2[0];
 349   w1_t[3]  = salt_buf_t2[1];
 350   w2_t[0]  = salt_buf_t2[2];
 351   w2_t[1]  = salt_buf_t2[3];
 352   w2_t[2]  = salt_buf_t2[4];
 353
 354   // 42..47
 355   w2_t[2] |= digest_t2[0];
 356   w2_t[3]  = digest_t2[1];
 357
 358   // 47..63
 359   w2_t[3] |= salt_buf_t3[0];
 360   w3_t[0]  = salt_buf_t3[1];
 361   w3_t[1]  = salt_buf_t3[2];
 362   w3_t[2]  = salt_buf_t3[3];
 363   w3_t[3]  = salt_buf_t3[4];
 364
 365   // 63..
 366
 367   w3_t[3] |= digest_t3[0];
 368
 369   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 370
 371   w0_t[0] = 0;
 372   w0_t[1] = 0;
 373   w0_t[2] = 0;
 374   w0_t[3] = 0;
 375   w1_t[0] = 0;
 376   w1_t[1] = 0;
 377   w1_t[2] = 0;
 378   w1_t[3] = 0;
 379   w2_t[0] = 0;
 380   w2_t[1] = 0;
 381   w2_t[2] = 0;
 382   w2_t[3] = 0;
 383   w3_t[0] = 0;
 384   w3_t[1] = 0;
 385   w3_t[2] = 0;
 386   w3_t[3] = 0;
 387
 388   // 0..4
 389   w0_t[0]  = digest_t3[1];
 390
 391   // 4..20
 392   w0_t[1]  = salt_buf_t0[0];
 393   w0_t[2]  = salt_buf_t0[1];
 394   w0_t[3]  = salt_buf_t0[2];
 395   w1_t[0]  = salt_buf_t0[3];
 396
 397   // 20..25
 398   w1_t[1]  = digest_t0[0];
 399   w1_t[2]  = digest_t0[1];
 400
 401   // 25..41
 402   w1_t[2] |= salt_buf_t1[0];
 403   w1_t[3]  = salt_buf_t1[1];
 404   w2_t[0]  = salt_buf_t1[2];
 405   w2_t[1]  = salt_buf_t1[3];
 406   w2_t[2]  = salt_buf_t1[4];
 407
 408   // 41..46
 409   w2_t[2] |= digest_t1[0];
 410   w2_t[3]  = digest_t1[1];
 411
 412   // 46..62
 413   w2_t[3] |= salt_buf_t2[0];
 414   w3_t[0]  = salt_buf_t2[1];
 415   w3_t[1]  = salt_buf_t2[2];
 416   w3_t[2]  = salt_buf_t2[3];
 417   w3_t[3]  = salt_buf_t2[4];
 418
 419   // 62..
 420   w3_t[3] |= digest_t2[0];
 421
 422   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 423
 424   w0_t[0] = 0;
 425   w0_t[1] = 0;
 426   w0_t[2] = 0;
 427   w0_t[3] = 0;
 428   w1_t[0] = 0;
 429   w1_t[1] = 0;
 430   w1_t[2] = 0;
 431   w1_t[3] = 0;
 432   w2_t[0] = 0;
 433   w2_t[1] = 0;
 434   w2_t[2] = 0;
 435   w2_t[3] = 0;
 436   w3_t[0] = 0;
 437   w3_t[1] = 0;
 438   w3_t[2] = 0;
 439   w3_t[3] = 0;
 440
 441   // 0..3
 442   w0_t[0]  = digest_t2[1];
 443
 444   // 3..19
 445   w0_t[0] |= salt_buf_t3[0];
 446   w0_t[1]  = salt_buf_t3[1];
 447   w0_t[2]  = salt_buf_t3[2];
 448   w0_t[3]  = salt_buf_t3[3];
 449   w1_t[0]  = salt_buf_t3[4];
 450
 451   // 19..24
 452   w1_t[0] |= digest_t3[0];
 453   w1_t[1]  = digest_t3[1];
 454
 455   // 24..40
 456   w1_t[2]  = salt_buf_t0[0];
 457   w1_t[3]  = salt_buf_t0[1];
 458   w2_t[0]  = salt_buf_t0[2];
 459   w2_t[1]  = salt_buf_t0[3];
 460
 461   // 40..45
 462   w2_t[2]  = digest_t0[0];
 463   w2_t[3]  = digest_t0[1];
 464
 465   // 45..61
 466   w2_t[3] |= salt_buf_t1[0];
 467   w3_t[0]  = salt_buf_t1[1];
 468   w3_t[1]  = salt_buf_t1[2];
 469   w3_t[2]  = salt_buf_t1[3];
 470   w3_t[3]  = salt_buf_t1[4];
 471
 472   // 61..
 473   w3_t[3] |= digest_t1[0];
 474
 475   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 476
 477   w0_t[0] = 0;
 478   w0_t[1] = 0;
 479   w0_t[2] = 0;
 480   w0_t[3] = 0;
 481   w1_t[0] = 0;
 482   w1_t[1] = 0;
 483   w1_t[2] = 0;
 484   w1_t[3] = 0;
 485   w2_t[0] = 0;
 486   w2_t[1] = 0;
 487   w2_t[2] = 0;
 488   w2_t[3] = 0;
 489   w3_t[0] = 0;
 490   w3_t[1] = 0;
 491   w3_t[2] = 0;
 492   w3_t[3] = 0;
 493
 494   // 0..2
 495   w0_t[0]  = digest_t1[1];
 496
 497   // 2..18
 498   w0_t[0] |= salt_buf_t2[0];
 499   w0_t[1]  = salt_buf_t2[1];
 500   w0_t[2]  = salt_buf_t2[2];
 501   w0_t[3]  = salt_buf_t2[3];
 502   w1_t[0]  = salt_buf_t2[4];
 503
 504   // 18..23
 505   w1_t[0] |= digest_t2[0];
 506   w1_t[1]  = digest_t2[1];
 507
 508   // 23..39
 509   w1_t[1] |= salt_buf_t3[0];
 510   w1_t[2]  = salt_buf_t3[1];
 511   w1_t[3]  = salt_buf_t3[2];
 512   w2_t[0]  = salt_buf_t3[3];
 513   w2_t[1]  = salt_buf_t3[4];
 514
 515   // 39..44
 516   w2_t[1] |= digest_t3[0];
 517   w2_t[2]  = digest_t3[1];
 518
 519   // 44..60
 520   w2_t[3]  = salt_buf_t0[0];
 521   w3_t[0]  = salt_buf_t0[1];
 522   w3_t[1]  = salt_buf_t0[2];
 523   w3_t[2]  = salt_buf_t0[3];
 524
 525   // 60..
 526   w3_t[3]  = digest_t0[0];
 527
 528   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 529
 530   w0_t[0] = 0;
 531   w0_t[1] = 0;
 532   w0_t[2] = 0;
 533   w0_t[3] = 0;
 534   w1_t[0] = 0;
 535   w1_t[1] = 0;
 536   w1_t[2] = 0;
 537   w1_t[3] = 0;
 538   w2_t[0] = 0;
 539   w2_t[1] = 0;
 540   w2_t[2] = 0;
 541   w2_t[3] = 0;
 542   w3_t[0] = 0;
 543   w3_t[1] = 0;
 544   w3_t[2] = 0;
 545   w3_t[3] = 0;
 546
 547   // 0..1
 548   w0_t[0]  = digest_t0[1];
 549
 550   // 1..17
 551   w0_t[0] |= salt_buf_t1[0];
 552   w0_t[1]  = salt_buf_t1[1];
 553   w0_t[2]  = salt_buf_t1[2];
 554   w0_t[3]  = salt_buf_t1[3];
 555   w1_t[0]  = salt_buf_t1[4];
 556
 557   // 17..22
 558   w1_t[0] |= digest_t1[0];
 559   w1_t[1]  = digest_t1[1];
 560
 561   // 22..38
 562   w1_t[1] |= salt_buf_t2[0];
 563   w1_t[2]  = salt_buf_t2[1];
 564   w1_t[3]  = salt_buf_t2[2];
 565   w2_t[0]  = salt_buf_t2[3];
 566   w2_t[1]  = salt_buf_t2[4];
 567
 568   // 38..43
 569   w2_t[1] |= digest_t2[0];
 570   w2_t[2]  = digest_t2[1];
 571
 572   // 43..59
 573   w2_t[2] |= salt_buf_t3[0];
 574   w2_t[3]  = salt_buf_t3[1];
 575   w3_t[0]  = salt_buf_t3[2];
 576   w3_t[1]  = salt_buf_t3[3];
 577   w3_t[2]  = salt_buf_t3[4];
 578
 579   // 59..
 580   w3_t[2] |= digest_t3[0];
 581   w3_t[3]  = digest_t3[1];
 582
 583   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 584
 585   w0_t[0]  = salt_buf_t0[0];
 586   w0_t[1]  = salt_buf_t0[1];
 587   w0_t[2]  = salt_buf_t0[2];
 588   w0_t[3]  = salt_buf_t0[3];
 589   w1_t[0]  = 0x80;
 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]  = 21 * 16 * 8;
 600   w3_t[3]  = 0;
 601
 602   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 603 }
 604
 605 __device__ __constant__ gpu_rule_t c_rules[1024];
 606
 607 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 608 {
 609   /**
 610    * modifier
 611    */
 612
 613   const u32 lid = threadIdx.x;
 614
 615   __shared__ RC4_KEY rc4_keys[64];
 616
 617   RC4_KEY *rc4_key = &rc4_keys[lid];
 618
 619   /**
 620    * base
 621    */
 622
 623   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 624
 625   if (gid >= gid_max) return;
 626
 627   u32x pw_buf0[4];
 628
 629   pw_buf0[0] = pws[gid].i[ 0];
 630   pw_buf0[1] = pws[gid].i[ 1];
 631   pw_buf0[2] = pws[gid].i[ 2];
 632   pw_buf0[3] = pws[gid].i[ 3];
 633
 634   u32x pw_buf1[4];
 635
 636   pw_buf1[0] = pws[gid].i[ 4];
 637   pw_buf1[1] = pws[gid].i[ 5];
 638   pw_buf1[2] = pws[gid].i[ 6];
 639   pw_buf1[3] = pws[gid].i[ 7];
 640
 641   const u32 pw_len = pws[gid].pw_len;
 642
 643   /**
 644    * salt
 645    */
 646
 647   u32 salt_buf[4];
 648
 649   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 650   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 651   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 652   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 653
 654   /**
 655    * esalt
 656    */
 657
 658   const u32 version = oldoffice01_bufs[salt_pos].version;
 659
 660   u32 encryptedVerifier[4];
 661
 662   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 663   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 664   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 665   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 666
 667   /**
 668    * loop
 669    */
 670
 671   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 672   {
 673     u32x w0[4];
 674
 675     w0[0] = pw_buf0[0];
 676     w0[1] = pw_buf0[1];
 677     w0[2] = pw_buf0[2];
 678     w0[3] = pw_buf0[3];
 679
 680     u32x w1[4];
 681
 682     w1[0] = pw_buf1[0];
 683     w1[1] = pw_buf1[1];
 684     w1[2] = pw_buf1[2];
 685     w1[3] = pw_buf1[3];
 686
 687     u32x w2[4];
 688
 689     w2[0] = 0;
 690     w2[1] = 0;
 691     w2[2] = 0;
 692     w2[3] = 0;
 693
 694     u32x w3[4];
 695
 696     w3[0] = 0;
 697     w3[1] = 0;
 698     w3[2] = 0;
 699     w3[3] = 0;
 700
 701     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 702
 703     append_0x80_2 (w0, w1, out_len);
 704
 705     u32x w0_t[4];
 706     u32x w1_t[4];
 707     u32x w2_t[4];
 708     u32x w3_t[4];
 709
 710     make_unicode (w0, w0_t, w1_t);
 711     make_unicode (w1, w2_t, w3_t);
 712
 713     w3_t[2] = out_len * 8 * 2;
 714
 715     u32x digest_pre[4];
 716
 717     digest_pre[0] = MD5M_A;
 718     digest_pre[1] = MD5M_B;
 719     digest_pre[2] = MD5M_C;
 720     digest_pre[3] = MD5M_D;
 721
 722     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_pre);
 723
 724     digest_pre[0] &= 0xffffffff;
 725     digest_pre[1] &= 0x000000ff;
 726     digest_pre[2] &= 0x00000000;
 727     digest_pre[3] &= 0x00000000;
 728
 729     u32x digest[4];
 730
 731     digest[0] = MD5M_A;
 732     digest[1] = MD5M_B;
 733     digest[2] = MD5M_C;
 734     digest[3] = MD5M_D;
 735
 736     gen336 (digest_pre, salt_buf, digest);
 737
 738     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
 739
 740     w0_t[0]  = digest[0];
 741     w0_t[1]  = digest[1] & 0xff;
 742     w0_t[2]  = 0x8000;
 743     w0_t[3]  = 0;
 744     w1_t[0]  = 0;
 745     w1_t[1]  = 0;
 746     w1_t[2]  = 0;
 747     w1_t[3]  = 0;
 748     w2_t[0]  = 0;
 749     w2_t[1]  = 0;
 750     w2_t[2]  = 0;
 751     w2_t[3]  = 0;
 752     w3_t[0]  = 0;
 753     w3_t[1]  = 0;
 754     w3_t[2]  = 9 * 8;
 755     w3_t[3]  = 0;
 756
 757     digest[0] = MD5M_A;
 758     digest[1] = MD5M_B;
 759     digest[2] = MD5M_C;
 760     digest[3] = MD5M_D;
 761
 762     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 763
 764     // now the RC4 part
 765
 766     u32x key[4];
 767
 768     key[0] = digest[0];
 769     key[1] = digest[1];
 770     key[2] = digest[2];
 771     key[3] = digest[3];
 772
 773     rc4_init_16 (rc4_key, key);
 774
 775     u32x out[4];
 776
 777     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 778
 779     w0_t[0] = out[0];
 780     w0_t[1] = out[1];
 781     w0_t[2] = out[2];
 782     w0_t[3] = out[3];
 783     w1_t[0] = 0x80;
 784     w1_t[1] = 0;
 785     w1_t[2] = 0;
 786     w1_t[3] = 0;
 787     w2_t[0] = 0;
 788     w2_t[1] = 0;
 789     w2_t[2] = 0;
 790     w2_t[3] = 0;
 791     w3_t[0] = 0;
 792     w3_t[1] = 0;
 793     w3_t[2] = 16 * 8;
 794     w3_t[3] = 0;
 795
 796     digest[0] = MD5M_A;
 797     digest[1] = MD5M_B;
 798     digest[2] = MD5M_C;
 799     digest[3] = MD5M_D;
 800
 801     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 802
 803     rc4_next_16 (rc4_key, 16, j, digest, out);
 804
 805     const u32x r0 = out[0];
 806     const u32x r1 = out[1];
 807     const u32x r2 = out[2];
 808     const u32x r3 = out[3];
 809
 810     #include VECT_COMPARE_M
 811   }
 812 }
 813
 814 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)
 815 {
 816 }
 817
 818 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)
 819 {
 820 }
 821
 822 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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 823 {
 824   /**
 825    * modifier
 826    */
 827
 828   const u32 lid = threadIdx.x;
 829
 830   __shared__ RC4_KEY rc4_keys[64];
 831
 832   RC4_KEY *rc4_key = &rc4_keys[lid];
 833
 834   /**
 835    * base
 836    */
 837
 838   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 839
 840   if (gid >= gid_max) return;
 841
 842   u32x pw_buf0[4];
 843
 844   pw_buf0[0] = pws[gid].i[ 0];
 845   pw_buf0[1] = pws[gid].i[ 1];
 846   pw_buf0[2] = pws[gid].i[ 2];
 847   pw_buf0[3] = pws[gid].i[ 3];
 848
 849   u32x pw_buf1[4];
 850
 851   pw_buf1[0] = pws[gid].i[ 4];
 852   pw_buf1[1] = pws[gid].i[ 5];
 853   pw_buf1[2] = pws[gid].i[ 6];
 854   pw_buf1[3] = pws[gid].i[ 7];
 855
 856   const u32 pw_len = pws[gid].pw_len;
 857
 858   /**
 859    * digest
 860    */
 861
 862   const u32 search[4] =
 863   {
 864     digests_buf[digests_offset].digest_buf[DGST_R0],
 865     digests_buf[digests_offset].digest_buf[DGST_R1],
 866     digests_buf[digests_offset].digest_buf[DGST_R2],
 867     digests_buf[digests_offset].digest_buf[DGST_R3]
 868   };
 869
 870   /**
 871    * salt
 872    */
 873
 874   u32 salt_buf[4];
 875
 876   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 877   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 878   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 879   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 880
 881   /**
 882    * esalt
 883    */
 884
 885   const u32 version = oldoffice01_bufs[salt_pos].version;
 886
 887   u32 encryptedVerifier[4];
 888
 889   encryptedVerifier[0] = oldoffice01_bufs[salt_pos].encryptedVerifier[0];
 890   encryptedVerifier[1] = oldoffice01_bufs[salt_pos].encryptedVerifier[1];
 891   encryptedVerifier[2] = oldoffice01_bufs[salt_pos].encryptedVerifier[2];
 892   encryptedVerifier[3] = oldoffice01_bufs[salt_pos].encryptedVerifier[3];
 893
 894   /**
 895    * loop
 896    */
 897
 898   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 899   {
 900     u32x w0[4];
 901
 902     w0[0] = pw_buf0[0];
 903     w0[1] = pw_buf0[1];
 904     w0[2] = pw_buf0[2];
 905     w0[3] = pw_buf0[3];
 906
 907     u32x w1[4];
 908
 909     w1[0] = pw_buf1[0];
 910     w1[1] = pw_buf1[1];
 911     w1[2] = pw_buf1[2];
 912     w1[3] = pw_buf1[3];
 913
 914     u32x w2[4];
 915
 916     w2[0] = 0;
 917     w2[1] = 0;
 918     w2[2] = 0;
 919     w2[3] = 0;
 920
 921     u32x w3[4];
 922
 923     w3[0] = 0;
 924     w3[1] = 0;
 925     w3[2] = 0;
 926     w3[3] = 0;
 927
 928     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 929
 930     append_0x80_2 (w0, w1, out_len);
 931
 932     u32x w0_t[4];
 933     u32x w1_t[4];
 934     u32x w2_t[4];
 935     u32x w3_t[4];
 936
 937     make_unicode (w0, w0_t, w1_t);
 938     make_unicode (w1, w2_t, w3_t);
 939
 940     w3_t[2] = out_len * 8 * 2;
 941
 942     u32x digest_pre[4];
 943
 944     digest_pre[0] = MD5M_A;
 945     digest_pre[1] = MD5M_B;
 946     digest_pre[2] = MD5M_C;
 947     digest_pre[3] = MD5M_D;
 948
 949     md5_transform (w0_t, w1_t, w2_t, w3_t, digest_pre);
 950
 951     digest_pre[0] &= 0xffffffff;
 952     digest_pre[1] &= 0x000000ff;
 953     digest_pre[2] &= 0x00000000;
 954     digest_pre[3] &= 0x00000000;
 955
 956     u32x digest[4];
 957
 958     digest[0] = MD5M_A;
 959     digest[1] = MD5M_B;
 960     digest[2] = MD5M_C;
 961     digest[3] = MD5M_D;
 962
 963     gen336 (digest_pre, salt_buf, digest);
 964
 965     // now the 40 bit input for the MD5 which then will generate the RC4 key, so it's precomputable!
 966
 967     w0_t[0]  = digest[0];
 968     w0_t[1]  = digest[1] & 0xff;
 969     w0_t[2]  = 0x8000;
 970     w0_t[3]  = 0;
 971     w1_t[0]  = 0;
 972     w1_t[1]  = 0;
 973     w1_t[2]  = 0;
 974     w1_t[3]  = 0;
 975     w2_t[0]  = 0;
 976     w2_t[1]  = 0;
 977     w2_t[2]  = 0;
 978     w2_t[3]  = 0;
 979     w3_t[0]  = 0;
 980     w3_t[1]  = 0;
 981     w3_t[2]  = 9 * 8;
 982     w3_t[3]  = 0;
 983
 984     digest[0] = MD5M_A;
 985     digest[1] = MD5M_B;
 986     digest[2] = MD5M_C;
 987     digest[3] = MD5M_D;
 988
 989     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 990
 991     // now the RC4 part
 992
 993     u32x key[4];
 994
 995     key[0] = digest[0];
 996     key[1] = digest[1];
 997     key[2] = digest[2];
 998     key[3] = digest[3];
 999
1000     rc4_init_16 (rc4_key, key);
1001
1002     u32x out[4];
1003
1004     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
1005
1006     w0_t[0] = out[0];
1007     w0_t[1] = out[1];
1008     w0_t[2] = out[2];
1009     w0_t[3] = out[3];
1010     w1_t[0] = 0x80;
1011     w1_t[1] = 0;
1012     w1_t[2] = 0;
1013     w1_t[3] = 0;
1014     w2_t[0] = 0;
1015     w2_t[1] = 0;
1016     w2_t[2] = 0;
1017     w2_t[3] = 0;
1018     w3_t[0] = 0;
1019     w3_t[1] = 0;
1020     w3_t[2] = 16 * 8;
1021     w3_t[3] = 0;
1022
1023     digest[0] = MD5M_A;
1024     digest[1] = MD5M_B;
1025     digest[2] = MD5M_C;
1026     digest[3] = MD5M_D;
1027
1028     md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
1029
1030     rc4_next_16 (rc4_key, 16, j, digest, out);
1031
1032     const u32x r0 = out[0];
1033     const u32x r1 = out[1];
1034     const u32x r2 = out[2];
1035     const u32x r3 = out[3];
1036
1037     #include VECT_COMPARE_S
1038   }
1039 }
1040
1041 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)
1042 {
1043 }
1044
1045 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)
1046 {
1047 }