nv/m10500.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MD5_
   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_M "check_multi_vect1_comp4.c"
  30 #endif
  31
  32 #ifdef  VECT_SIZE2
  33 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  34 #endif
  35
  36 __device__ __constant__ u32 padding[8] =
  37 {
  38   0x5e4ebf28,
  39   0x418a754e,
  40   0x564e0064,
  41   0x0801faff,
  42   0xb6002e2e,
  43   0x803e68d0,
  44   0xfea90c2f,
  45   0x7a695364
  46 };
  47
  48 typedef struct
  49 {
  50   u8 S[256];
  51
  52   u32 wtf_its_faster;
  53
  54 } RC4_KEY;
  55
  56 __device__ static void swap (RC4_KEY *rc4_key, const u8 i, const u8 j)
  57 {
  58   u8 tmp;
  59
  60   tmp           = rc4_key->S[i];
  61   rc4_key->S[i] = rc4_key->S[j];
  62   rc4_key->S[j] = tmp;
  63 }
  64
  65 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  66 {
  67   u32 v = 0x03020100;
  68   u32 a = 0x04040404;
  69
  70   u32 *ptr = (u32 *) rc4_key->S;
  71
  72   #pragma unroll 64
  73   for (u32 i = 0; i < 64; i++)
  74   {
  75     *ptr++ = v; v += a;
  76   }
  77
  78   u32 j = 0;
  79
  80   for (u32 i = 0; i < 16; i++)
  81   {
  82     u32 idx = i * 16;
  83
  84     u32 v;
  85
  86     v = data[0];
  87
  88     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  89     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  90     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  91     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  92
  93     v = data[1];
  94
  95     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
  96     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
  97     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
  98     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
  99
 100     v = data[2];
 101
 102     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
 103     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
 104     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 105     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 106
 107     v = data[3];
 108
 109     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j); idx++;
 110     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j); idx++;
 111     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j); idx++;
 112     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j); idx++;
 113   }
 114 }
 115
 116 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 117 {
 118   for (u32 k = 0; k < 4; k++)
 119   {
 120     u32 xor4 = 0;
 121
 122     u8 idx;
 123
 124     i += 1;
 125     j += rc4_key->S[i];
 126
 127     swap (rc4_key, i, j);
 128
 129     idx = rc4_key->S[i] + rc4_key->S[j];
 130
 131     xor4 |= rc4_key->S[idx] <<  0;
 132
 133     i += 1;
 134     j += rc4_key->S[i];
 135
 136     swap (rc4_key, i, j);
 137
 138     idx = rc4_key->S[i] + rc4_key->S[j];
 139
 140     xor4 |= rc4_key->S[idx] <<  8;
 141
 142     i += 1;
 143     j += rc4_key->S[i];
 144
 145     swap (rc4_key, i, j);
 146
 147     idx = rc4_key->S[i] + rc4_key->S[j];
 148
 149     xor4 |= rc4_key->S[idx] << 16;
 150
 151     i += 1;
 152     j += rc4_key->S[i];
 153
 154     swap (rc4_key, i, j);
 155
 156     idx = rc4_key->S[i] + rc4_key->S[j];
 157
 158     xor4 |= rc4_key->S[idx] << 24;
 159
 160     out[k] = in[k] ^ xor4;
 161   }
 162
 163   return j;
 164 }
 165
 166 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
 167 {
 168   u32x a = digest[0];
 169   u32x b = digest[1];
 170   u32x c = digest[2];
 171   u32x d = digest[3];
 172
 173   u32x w0_t = w0[0];
 174   u32x w1_t = w0[1];
 175   u32x w2_t = w0[2];
 176   u32x w3_t = w0[3];
 177   u32x w4_t = w1[0];
 178   u32x w5_t = w1[1];
 179   u32x w6_t = w1[2];
 180   u32x w7_t = w1[3];
 181   u32x w8_t = w2[0];
 182   u32x w9_t = w2[1];
 183   u32x wa_t = w2[2];
 184   u32x wb_t = w2[3];
 185   u32x wc_t = w3[0];
 186   u32x wd_t = w3[1];
 187   u32x we_t = w3[2];
 188   u32x wf_t = w3[3];
 189
 190   u32x tmp2;
 191
 192   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
 193   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
 194   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
 195   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
 196   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
 197   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
 198   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
 199   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
 200   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
 201   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
 202   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
 203   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
 204   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
 205   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
 206   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
 207   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
 208
 209   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
 210   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
 211   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
 212   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
 213   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
 214   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
 215   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
 216   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
 217   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
 218   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
 219   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
 220   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
 221   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
 222   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
 223   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 224   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 225
 226   MD5_STEP (MD5_H1, a, b, c, d, w5_t, MD5C20, MD5S20);
 227   MD5_STEP (MD5_H2, d, a, b, c, w8_t, MD5C21, MD5S21);
 228   MD5_STEP (MD5_H1, c, d, a, b, wb_t, MD5C22, MD5S22);
 229   MD5_STEP (MD5_H2, b, c, d, a, we_t, MD5C23, MD5S23);
 230   MD5_STEP (MD5_H1, a, b, c, d, w1_t, MD5C24, MD5S20);
 231   MD5_STEP (MD5_H2, d, a, b, c, w4_t, MD5C25, MD5S21);
 232   MD5_STEP (MD5_H1, c, d, a, b, w7_t, MD5C26, MD5S22);
 233   MD5_STEP (MD5_H2, b, c, d, a, wa_t, MD5C27, MD5S23);
 234   MD5_STEP (MD5_H1, a, b, c, d, wd_t, MD5C28, MD5S20);
 235   MD5_STEP (MD5_H2, d, a, b, c, w0_t, MD5C29, MD5S21);
 236   MD5_STEP (MD5_H1, c, d, a, b, w3_t, MD5C2a, MD5S22);
 237   MD5_STEP (MD5_H2, b, c, d, a, w6_t, MD5C2b, MD5S23);
 238   MD5_STEP (MD5_H1, a, b, c, d, w9_t, MD5C2c, MD5S20);
 239   MD5_STEP (MD5_H2, d, a, b, c, wc_t, MD5C2d, MD5S21);
 240   MD5_STEP (MD5_H1, c, d, a, b, wf_t, MD5C2e, MD5S22);
 241   MD5_STEP (MD5_H2, b, c, d, a, w2_t, MD5C2f, MD5S23);
 242
 243   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 244   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 245   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 246   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 247   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 248   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 249   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 250   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 251   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 252   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 253   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 254   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 255   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 256   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 257   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 258   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 259
 260   digest[0] += a;
 261   digest[1] += b;
 262   digest[2] += c;
 263   digest[3] += d;
 264 }
 265
 266 extern "C" __global__ void __launch_bounds__ (64, 1) m10500_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pdf14_tmp_t *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 pdf_t *pdf_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)
 267 {
 268   /**
 269    * base
 270    */
 271
 272   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 273
 274   if (gid >= gid_max) return;
 275
 276   u32x w0[4];
 277
 278   w0[0] = pws[gid].i[ 0];
 279   w0[1] = pws[gid].i[ 1];
 280   w0[2] = pws[gid].i[ 2];
 281   w0[3] = pws[gid].i[ 3];
 282
 283   u32x w1[4];
 284
 285   w1[0] = pws[gid].i[ 4];
 286   w1[1] = pws[gid].i[ 5];
 287   w1[2] = pws[gid].i[ 6];
 288   w1[3] = pws[gid].i[ 7];
 289
 290   u32x w2[4];
 291
 292   w2[0] = pws[gid].i[ 8];
 293   w2[1] = pws[gid].i[ 9];
 294   w2[2] = 0;
 295   w2[3] = 0;
 296
 297   const u32 pw_len = pws[gid].pw_len;
 298
 299   /**
 300    * U_buf
 301    */
 302
 303   u32 o_buf[8];
 304
 305   o_buf[0] = pdf_bufs[salt_pos].o_buf[0];
 306   o_buf[1] = pdf_bufs[salt_pos].o_buf[1];
 307   o_buf[2] = pdf_bufs[salt_pos].o_buf[2];
 308   o_buf[3] = pdf_bufs[salt_pos].o_buf[3];
 309   o_buf[4] = pdf_bufs[salt_pos].o_buf[4];
 310   o_buf[5] = pdf_bufs[salt_pos].o_buf[5];
 311   o_buf[6] = pdf_bufs[salt_pos].o_buf[6];
 312   o_buf[7] = pdf_bufs[salt_pos].o_buf[7];
 313
 314   u32 P = pdf_bufs[salt_pos].P;
 315
 316   u32 id_buf[12];
 317
 318   id_buf[ 0] = pdf_bufs[salt_pos].id_buf[0];
 319   id_buf[ 1] = pdf_bufs[salt_pos].id_buf[1];
 320   id_buf[ 2] = pdf_bufs[salt_pos].id_buf[2];
 321   id_buf[ 3] = pdf_bufs[salt_pos].id_buf[3];
 322
 323   id_buf[ 4] = pdf_bufs[salt_pos].id_buf[4];
 324   id_buf[ 5] = pdf_bufs[salt_pos].id_buf[5];
 325   id_buf[ 6] = pdf_bufs[salt_pos].id_buf[6];
 326   id_buf[ 7] = pdf_bufs[salt_pos].id_buf[7];
 327
 328   id_buf[ 8] = 0;
 329   id_buf[ 9] = 0;
 330   id_buf[10] = 0;
 331   id_buf[11] = 0;
 332
 333   u32 id_len  = pdf_bufs[salt_pos].id_len;
 334   u32 id_len4 = id_len / 4;
 335
 336   u32 rc4data[2];
 337
 338   rc4data[0] = pdf_bufs[salt_pos].rc4data[0];
 339   rc4data[1] = pdf_bufs[salt_pos].rc4data[1];
 340
 341   u32 final_length = 68 + id_len;
 342
 343   u32x w11 = 0x80;
 344   u32x w12 = 0;
 345
 346   if (pdf_bufs[salt_pos].enc_md != 1)
 347   {
 348     w11 = 0xffffffff;
 349     w12 = 0x80;
 350
 351     final_length += 4;
 352   }
 353
 354   id_buf[id_len4 + 0] = w11;
 355   id_buf[id_len4 + 1] = w12;
 356
 357   /**
 358    * main init
 359    */
 360
 361   u32x w0_t[4];
 362   u32x w1_t[4];
 363   u32x w2_t[4];
 364   u32x w3_t[4];
 365
 366   // max length supported by pdf11 is 32
 367
 368   w0_t[0] = padding[0];
 369   w0_t[1] = padding[1];
 370   w0_t[2] = padding[2];
 371   w0_t[3] = padding[3];
 372   w1_t[0] = padding[4];
 373   w1_t[1] = padding[5];
 374   w1_t[2] = padding[6];
 375   w1_t[3] = padding[7];
 376   w2_t[0] = 0;
 377   w2_t[1] = 0;
 378   w2_t[2] = 0;
 379   w2_t[3] = 0;
 380   w3_t[0] = 0;
 381   w3_t[1] = 0;
 382   w3_t[2] = 0;
 383   w3_t[3] = 0;
 384
 385   switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, pw_len);
 386
 387   // add password
 388   // truncate at 32 is wanted, not a bug!
 389   // add o_buf
 390
 391   w0_t[0] |= w0[0];
 392   w0_t[1] |= w0[1];
 393   w0_t[2] |= w0[2];
 394   w0_t[3] |= w0[3];
 395   w1_t[0] |= w1[0];
 396   w1_t[1] |= w1[1];
 397   w1_t[2] |= w1[2];
 398   w1_t[3] |= w1[3];
 399   w2_t[0]  = o_buf[0];
 400   w2_t[1]  = o_buf[1];
 401   w2_t[2]  = o_buf[2];
 402   w2_t[3]  = o_buf[3];
 403   w3_t[0]  = o_buf[4];
 404   w3_t[1]  = o_buf[5];
 405   w3_t[2]  = o_buf[6];
 406   w3_t[3]  = o_buf[7];
 407
 408   u32x digest[4];
 409
 410   digest[0] = MD5M_A;
 411   digest[1] = MD5M_B;
 412   digest[2] = MD5M_C;
 413   digest[3] = MD5M_D;
 414
 415   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 416
 417   w0_t[0] = P;
 418   w0_t[1] = id_buf[ 0];
 419   w0_t[2] = id_buf[ 1];
 420   w0_t[3] = id_buf[ 2];
 421   w1_t[0] = id_buf[ 3];
 422   w1_t[1] = id_buf[ 4];
 423   w1_t[2] = id_buf[ 5];
 424   w1_t[3] = id_buf[ 6];
 425   w2_t[0] = id_buf[ 7];
 426   w2_t[1] = id_buf[ 8];
 427   w2_t[2] = id_buf[ 9];
 428   w2_t[3] = id_buf[10];
 429   w3_t[0] = id_buf[11];
 430   w3_t[1] = 0;
 431   w3_t[2] = final_length * 8;
 432   w3_t[3] = 0;
 433
 434   md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 435
 436   tmps[gid].digest[0] = digest[0];
 437   tmps[gid].digest[1] = digest[1];
 438   tmps[gid].digest[2] = digest[2];
 439   tmps[gid].digest[3] = digest[3];
 440
 441   tmps[gid].out[0] = rc4data[0];
 442   tmps[gid].out[1] = rc4data[1];
 443   tmps[gid].out[2] = 0;
 444   tmps[gid].out[3] = 0;
 445 }
 446
 447 extern "C" __global__ void __launch_bounds__ (64, 1) m10500_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pdf14_tmp_t *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 pdf_t *pdf_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)
 448 {
 449   /**
 450    * base
 451    */
 452
 453   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 454   const u32 lid = threadIdx.x;
 455
 456   if (gid >= gid_max) return;
 457
 458   /**
 459    * shared
 460    */
 461
 462   __shared__ RC4_KEY rc4_keys[64];
 463
 464   RC4_KEY *rc4_key = &rc4_keys[lid];
 465
 466   /**
 467    * loop
 468    */
 469
 470   u32x digest[4];
 471
 472   digest[0] = tmps[gid].digest[0];
 473   digest[1] = tmps[gid].digest[1];
 474   digest[2] = tmps[gid].digest[2];
 475   digest[3] = tmps[gid].digest[3];
 476
 477   u32x out[4];
 478
 479   out[0] = tmps[gid].out[0];
 480   out[1] = tmps[gid].out[1];
 481   out[2] = tmps[gid].out[2];
 482   out[3] = tmps[gid].out[3];
 483
 484   for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
 485   {
 486     if (j < 50)
 487     {
 488       u32x w0_t[4];
 489       u32x w1_t[4];
 490       u32x w2_t[4];
 491       u32x w3_t[4];
 492
 493       w0_t[0] = digest[0];
 494       w0_t[1] = digest[1];
 495       w0_t[2] = digest[2];
 496       w0_t[3] = digest[3];
 497       w1_t[0] = 0x80;
 498       w1_t[1] = 0;
 499       w1_t[2] = 0;
 500       w1_t[3] = 0;
 501       w2_t[0] = 0;
 502       w2_t[1] = 0;
 503       w2_t[2] = 0;
 504       w2_t[3] = 0;
 505       w3_t[0] = 0;
 506       w3_t[1] = 0;
 507       w3_t[2] = 16 * 8;
 508       w3_t[3] = 0;
 509
 510       digest[0] = MD5M_A;
 511       digest[1] = MD5M_B;
 512       digest[2] = MD5M_C;
 513       digest[3] = MD5M_D;
 514
 515       md5_transform (w0_t, w1_t, w2_t, w3_t, digest);
 516     }
 517     else
 518     {
 519       const u32 x = j - 50;
 520
 521       const u32 xv = x <<  0
 522                     | x <<  8
 523                     | x << 16
 524                     | x << 24;
 525
 526       u32x tmp[4];
 527
 528       tmp[0] = digest[0] ^ xv;
 529       tmp[1] = digest[1] ^ xv;
 530       tmp[2] = digest[2] ^ xv;
 531       tmp[3] = digest[3] ^ xv;
 532
 533       rc4_init_16 (rc4_key, tmp);
 534
 535       rc4_next_16 (rc4_key, 0, 0, out, out);
 536     }
 537   }
 538
 539   tmps[gid].digest[0] = digest[0];
 540   tmps[gid].digest[1] = digest[1];
 541   tmps[gid].digest[2] = digest[2];
 542   tmps[gid].digest[3] = digest[3];
 543
 544   tmps[gid].out[0] = out[0];
 545   tmps[gid].out[1] = out[1];
 546   tmps[gid].out[2] = out[2];
 547   tmps[gid].out[3] = out[3];
 548 }
 549
 550 extern "C" __global__ void __launch_bounds__ (64, 1) m10500_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, pdf14_tmp_t *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 pdf_t *pdf_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)
 551 {
 552   /**
 553    * modifier
 554    */
 555
 556   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 557
 558   if (gid >= gid_max) return;
 559
 560   const u32 lid = threadIdx.x;
 561
 562   /**
 563    * digest
 564    */
 565
 566   const u32x r0 = tmps[gid].out[0];
 567   const u32x r1 = tmps[gid].out[1];
 568   const u32x r2 = 0;
 569   const u32x r3 = 0;
 570
 571   #define il_pos 0
 572
 573   #include VECT_COMPARE_M
 574 }