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