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