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