nv/m11000_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) m11000_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_buf0[4];
 101
 102   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 103   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
 104   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
 105   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
 106
 107   u32 salt_buf1[4];
 108
 109   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
 110   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
 111   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
 112   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
 113
 114   u32 salt_buf2[4];
 115
 116   salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
 117   salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
 118   salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
 119   salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
 120
 121   u32 salt_buf3[2];
 122
 123   salt_buf3[0] = salt_bufs[salt_pos].salt_buf[12];
 124   salt_buf3[1] = salt_bufs[salt_pos].salt_buf[13];
 125
 126   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 127
 128   /**
 129    * loop
 130    */
 131
 132   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 133   {
 134     const u32 pw_r_len = c_combs[il_pos].pw_len;
 135
 136     const u32 pw_len = pw_l_len + pw_r_len;
 137
 138     u32 wordr0[4];
 139     u32 wordr1[4];
 140     u32 wordr2[4];
 141     u32 wordr3[4];
 142
 143     wordr0[0] = c_combs[il_pos].i[0];
 144     wordr0[1] = c_combs[il_pos].i[1];
 145     wordr0[2] = c_combs[il_pos].i[2];
 146     wordr0[3] = c_combs[il_pos].i[3];
 147     wordr1[0] = c_combs[il_pos].i[4];
 148     wordr1[1] = c_combs[il_pos].i[5];
 149     wordr1[2] = c_combs[il_pos].i[6];
 150     wordr1[3] = c_combs[il_pos].i[7];
 151     wordr2[0] = 0;
 152     wordr2[1] = 0;
 153     wordr2[2] = 0;
 154     wordr2[3] = 0;
 155     wordr3[0] = 0;
 156     wordr3[1] = 0;
 157     wordr3[2] = 0;
 158     wordr3[3] = 0;
 159
 160     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 161     {
 162       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 163     }
 164
 165     u32x w0[4];
 166     u32x w1[4];
 167     u32x w2[4];
 168     u32x w3[4];
 169
 170     w0[0] = wordl0[0] | wordr0[0];
 171     w0[1] = wordl0[1] | wordr0[1];
 172     w0[2] = wordl0[2] | wordr0[2];
 173     w0[3] = wordl0[3] | wordr0[3];
 174     w1[0] = wordl1[0] | wordr1[0];
 175     w1[1] = wordl1[1] | wordr1[1];
 176     w1[2] = wordl1[2] | wordr1[2];
 177     w1[3] = wordl1[3] | wordr1[3];
 178     w2[0] = wordl2[0] | wordr2[0];
 179     w2[1] = wordl2[1] | wordr2[1];
 180     w2[2] = wordl2[2] | wordr2[2];
 181     w2[3] = wordl2[3] | wordr2[3];
 182     w3[0] = wordl3[0] | wordr3[0];
 183     w3[1] = wordl3[1] | wordr3[1];
 184     w3[2] = wordl3[2] | wordr3[2];
 185     w3[3] = wordl3[3] | wordr3[3];
 186
 187     append_0x80_4 (w0, w1, w2, w3, pw_len);
 188
 189     const u32 pw_salt_len = pw_len + salt_len;
 190
 191     /**
 192      * prepend salt
 193      */
 194
 195     // first step fixed 56 bytes of salt
 196
 197     u32x w0_t[4];
 198     u32x w1_t[4];
 199     u32x w2_t[4];
 200     u32x w3_t[4];
 201
 202     w0_t[0] = salt_buf0[0];
 203     w0_t[1] = salt_buf0[1];
 204     w0_t[2] = salt_buf0[2];
 205     w0_t[3] = salt_buf0[3];
 206     w1_t[0] = salt_buf1[0];
 207     w1_t[1] = salt_buf1[1];
 208     w1_t[2] = salt_buf1[2];
 209     w1_t[3] = salt_buf1[3];
 210     w2_t[0] = salt_buf2[0];
 211     w2_t[1] = salt_buf2[1];
 212     w2_t[2] = salt_buf2[2];
 213     w2_t[3] = salt_buf2[3];
 214     w3_t[0] = salt_buf3[0];
 215     w3_t[1] = salt_buf3[1];
 216
 217     // after 56 byte salt, we have beginning of the password
 218
 219     w3_t[2] = w0[0];
 220     w3_t[3] = w0[1];
 221
 222     /**
 223      * md5
 224      */
 225
 226     // first transform
 227
 228     u32x a = MD5M_A;
 229     u32x b = MD5M_B;
 230     u32x c = MD5M_C;
 231     u32x d = MD5M_D;
 232
 233     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
 234     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
 235     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
 236     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
 237     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
 238     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
 239     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
 240     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
 241     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
 242     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
 243     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
 244     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
 245     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
 246     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
 247     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
 248     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
 249
 250     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
 251     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
 252     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
 253     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
 254     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
 255     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
 256     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
 257     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
 258     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
 259     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
 260     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
 261     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
 262     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
 263     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
 264     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
 265     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
 266
 267     MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
 268     MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
 269     MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
 270     MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
 271     MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
 272     MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
 273     MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
 274     MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
 275     MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
 276     MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
 277     MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
 278     MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
 279     MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
 280     MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
 281     MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
 282     MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
 283
 284     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
 285     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
 286     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
 287     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
 288     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
 289     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
 290     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
 291     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
 292     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
 293     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
 294     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
 295     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
 296     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
 297     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
 298     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
 299     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
 300
 301     a += MD5M_A;
 302     b += MD5M_B;
 303     c += MD5M_C;
 304     d += MD5M_D;
 305
 306     u32x r_a = a;
 307     u32x r_b = b;
 308     u32x r_c = c;
 309     u32x r_d = d;
 310
 311     // 2nd transform
 312
 313     w0_t[0] = w0[2];
 314     w0_t[1] = w0[3];
 315     w0_t[2] = w1[0];
 316     w0_t[3] = w1[1];
 317     w1_t[0] = w1[2];
 318     w1_t[1] = w1[3];
 319     w1_t[2] = w2[0];
 320     w1_t[3] = w2[1];
 321     w2_t[0] = w2[2];
 322     w2_t[1] = w2[3];
 323     w2_t[2] = w3[0];
 324     w2_t[3] = w3[1];
 325     w3_t[0] = w3[2];
 326     w3_t[1] = w3[3];
 327     w3_t[2] = pw_salt_len * 8;
 328     w3_t[3] = 0;
 329
 330     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
 331     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
 332     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
 333     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
 334     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
 335     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
 336     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
 337     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
 338     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
 339     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
 340     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
 341     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
 342     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
 343     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
 344     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
 345     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
 346
 347     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
 348     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
 349     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
 350     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
 351     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
 352     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
 353     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
 354     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
 355     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
 356     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
 357     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
 358     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
 359     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
 360     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
 361     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
 362     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
 363
 364     MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
 365     MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
 366     MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
 367     MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
 368     MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
 369     MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
 370     MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
 371     MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
 372     MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
 373     MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
 374     MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
 375     MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
 376     MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
 377     MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
 378     MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
 379     MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
 380
 381     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
 382     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
 383     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
 384     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
 385     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
 386     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
 387     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
 388     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
 389     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
 390     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
 391     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
 392     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
 393     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
 394     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
 395     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
 396     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
 397
 398     a += r_a;
 399     b += r_b;
 400     c += r_c;
 401     d += r_d;
 402
 403     const u32x r0 = a;
 404     const u32x r1 = d;
 405     const u32x r2 = c;
 406     const u32x r3 = b;
 407
 408     #include VECT_COMPARE_M
 409   }
 410 }
 411
 412 extern "C" __global__ void __launch_bounds__ (256, 1) m11000_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)
 413 {
 414 }
 415
 416 extern "C" __global__ void __launch_bounds__ (256, 1) m11000_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)
 417 {
 418 }
 419
 420 extern "C" __global__ void __launch_bounds__ (256, 1) m11000_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)
 421 {
 422   /**
 423    * modifier
 424    */
 425
 426   const u32 lid = threadIdx.x;
 427
 428   /**
 429    * base
 430    */
 431
 432   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 433
 434   if (gid >= gid_max) return;
 435
 436   u32x wordl0[4];
 437
 438   wordl0[0] = pws[gid].i[ 0];
 439   wordl0[1] = pws[gid].i[ 1];
 440   wordl0[2] = pws[gid].i[ 2];
 441   wordl0[3] = pws[gid].i[ 3];
 442
 443   u32x wordl1[4];
 444
 445   wordl1[0] = pws[gid].i[ 4];
 446   wordl1[1] = pws[gid].i[ 5];
 447   wordl1[2] = pws[gid].i[ 6];
 448   wordl1[3] = pws[gid].i[ 7];
 449
 450   u32x wordl2[4];
 451
 452   wordl2[0] = 0;
 453   wordl2[1] = 0;
 454   wordl2[2] = 0;
 455   wordl2[3] = 0;
 456
 457   u32x wordl3[4];
 458
 459   wordl3[0] = 0;
 460   wordl3[1] = 0;
 461   wordl3[2] = 0;
 462   wordl3[3] = 0;
 463
 464   const u32 pw_l_len = pws[gid].pw_len;
 465
 466   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 467   {
 468     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 469   }
 470
 471   /**
 472    * salt
 473    */
 474
 475   u32 salt_buf0[4];
 476
 477   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 478   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
 479   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
 480   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
 481
 482   u32 salt_buf1[4];
 483
 484   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
 485   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
 486   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
 487   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
 488
 489   u32 salt_buf2[4];
 490
 491   salt_buf2[0] = salt_bufs[salt_pos].salt_buf[ 8];
 492   salt_buf2[1] = salt_bufs[salt_pos].salt_buf[ 9];
 493   salt_buf2[2] = salt_bufs[salt_pos].salt_buf[10];
 494   salt_buf2[3] = salt_bufs[salt_pos].salt_buf[11];
 495
 496   u32 salt_buf3[2];
 497
 498   salt_buf3[0] = salt_bufs[salt_pos].salt_buf[12];
 499   salt_buf3[1] = salt_bufs[salt_pos].salt_buf[13];
 500
 501   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 502
 503   /**
 504    * digest
 505    */
 506
 507   const u32 search[4] =
 508   {
 509     digests_buf[digests_offset].digest_buf[DGST_R0],
 510     digests_buf[digests_offset].digest_buf[DGST_R1],
 511     digests_buf[digests_offset].digest_buf[DGST_R2],
 512     digests_buf[digests_offset].digest_buf[DGST_R3]
 513   };
 514
 515   /**
 516    * loop
 517    */
 518
 519   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 520   {
 521     const u32 pw_r_len = c_combs[il_pos].pw_len;
 522
 523     const u32 pw_len = pw_l_len + pw_r_len;
 524
 525     u32 wordr0[4];
 526     u32 wordr1[4];
 527     u32 wordr2[4];
 528     u32 wordr3[4];
 529
 530     wordr0[0] = c_combs[il_pos].i[0];
 531     wordr0[1] = c_combs[il_pos].i[1];
 532     wordr0[2] = c_combs[il_pos].i[2];
 533     wordr0[3] = c_combs[il_pos].i[3];
 534     wordr1[0] = c_combs[il_pos].i[4];
 535     wordr1[1] = c_combs[il_pos].i[5];
 536     wordr1[2] = c_combs[il_pos].i[6];
 537     wordr1[3] = c_combs[il_pos].i[7];
 538     wordr2[0] = 0;
 539     wordr2[1] = 0;
 540     wordr2[2] = 0;
 541     wordr2[3] = 0;
 542     wordr3[0] = 0;
 543     wordr3[1] = 0;
 544     wordr3[2] = 0;
 545     wordr3[3] = 0;
 546
 547     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 548     {
 549       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 550     }
 551
 552     u32x w0[4];
 553     u32x w1[4];
 554     u32x w2[4];
 555     u32x w3[4];
 556
 557     w0[0] = wordl0[0] | wordr0[0];
 558     w0[1] = wordl0[1] | wordr0[1];
 559     w0[2] = wordl0[2] | wordr0[2];
 560     w0[3] = wordl0[3] | wordr0[3];
 561     w1[0] = wordl1[0] | wordr1[0];
 562     w1[1] = wordl1[1] | wordr1[1];
 563     w1[2] = wordl1[2] | wordr1[2];
 564     w1[3] = wordl1[3] | wordr1[3];
 565     w2[0] = wordl2[0] | wordr2[0];
 566     w2[1] = wordl2[1] | wordr2[1];
 567     w2[2] = wordl2[2] | wordr2[2];
 568     w2[3] = wordl2[3] | wordr2[3];
 569     w3[0] = wordl3[0] | wordr3[0];
 570     w3[1] = wordl3[1] | wordr3[1];
 571     w3[2] = wordl3[2] | wordr3[2];
 572     w3[3] = wordl3[3] | wordr3[3];
 573
 574     append_0x80_4 (w0, w1, w2, w3, pw_len);
 575
 576     const u32 pw_salt_len = pw_len + salt_len;
 577
 578     /**
 579      * prepend salt
 580      */
 581
 582     // first step fixed 56 bytes of salt
 583
 584     u32x w0_t[4];
 585     u32x w1_t[4];
 586     u32x w2_t[4];
 587     u32x w3_t[4];
 588
 589     w0_t[0] = salt_buf0[0];
 590     w0_t[1] = salt_buf0[1];
 591     w0_t[2] = salt_buf0[2];
 592     w0_t[3] = salt_buf0[3];
 593     w1_t[0] = salt_buf1[0];
 594     w1_t[1] = salt_buf1[1];
 595     w1_t[2] = salt_buf1[2];
 596     w1_t[3] = salt_buf1[3];
 597     w2_t[0] = salt_buf2[0];
 598     w2_t[1] = salt_buf2[1];
 599     w2_t[2] = salt_buf2[2];
 600     w2_t[3] = salt_buf2[3];
 601     w3_t[0] = salt_buf3[0];
 602     w3_t[1] = salt_buf3[1];
 603
 604     // after 56 byte salt, we have beginning of the password
 605
 606     w3_t[2] = w0[0];
 607     w3_t[3] = w0[1];
 608
 609     /**
 610      * md5
 611      */
 612
 613     // first transform
 614
 615     u32x a = MD5M_A;
 616     u32x b = MD5M_B;
 617     u32x c = MD5M_C;
 618     u32x d = MD5M_D;
 619
 620     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
 621     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
 622     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
 623     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
 624     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
 625     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
 626     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
 627     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
 628     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
 629     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
 630     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
 631     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
 632     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
 633     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
 634     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
 635     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
 636
 637     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
 638     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
 639     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
 640     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
 641     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
 642     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
 643     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
 644     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
 645     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
 646     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
 647     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
 648     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
 649     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
 650     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
 651     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
 652     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
 653
 654     MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
 655     MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
 656     MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
 657     MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
 658     MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
 659     MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
 660     MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
 661     MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
 662     MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
 663     MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
 664     MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
 665     MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
 666     MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
 667     MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
 668     MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
 669     MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
 670
 671     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
 672     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
 673     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
 674     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
 675     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
 676     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
 677     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
 678     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
 679     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
 680     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
 681     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
 682     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
 683     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
 684     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
 685     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
 686     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
 687
 688     a += MD5M_A;
 689     b += MD5M_B;
 690     c += MD5M_C;
 691     d += MD5M_D;
 692
 693     u32x r_a = a;
 694     u32x r_b = b;
 695     u32x r_c = c;
 696     u32x r_d = d;
 697
 698     // 2nd transform
 699
 700     w0_t[0] = w0[2];
 701     w0_t[1] = w0[3];
 702     w0_t[2] = w1[0];
 703     w0_t[3] = w1[1];
 704     w1_t[0] = w1[2];
 705     w1_t[1] = w1[3];
 706     w1_t[2] = w2[0];
 707     w1_t[3] = w2[1];
 708     w2_t[0] = w2[2];
 709     w2_t[1] = w2[3];
 710     w2_t[2] = w3[0];
 711     w2_t[3] = w3[1];
 712     w3_t[0] = w3[2];
 713     w3_t[1] = w3[3];
 714     w3_t[2] = pw_salt_len * 8;
 715     w3_t[3] = 0;
 716
 717     MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00);
 718     MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01);
 719     MD5_STEP (MD5_Fo, c, d, a, b, w0_t[2], MD5C02, MD5S02);
 720     MD5_STEP (MD5_Fo, b, c, d, a, w0_t[3], MD5C03, MD5S03);
 721     MD5_STEP (MD5_Fo, a, b, c, d, w1_t[0], MD5C04, MD5S00);
 722     MD5_STEP (MD5_Fo, d, a, b, c, w1_t[1], MD5C05, MD5S01);
 723     MD5_STEP (MD5_Fo, c, d, a, b, w1_t[2], MD5C06, MD5S02);
 724     MD5_STEP (MD5_Fo, b, c, d, a, w1_t[3], MD5C07, MD5S03);
 725     MD5_STEP (MD5_Fo, a, b, c, d, w2_t[0], MD5C08, MD5S00);
 726     MD5_STEP (MD5_Fo, d, a, b, c, w2_t[1], MD5C09, MD5S01);
 727     MD5_STEP (MD5_Fo, c, d, a, b, w2_t[2], MD5C0a, MD5S02);
 728     MD5_STEP (MD5_Fo, b, c, d, a, w2_t[3], MD5C0b, MD5S03);
 729     MD5_STEP (MD5_Fo, a, b, c, d, w3_t[0], MD5C0c, MD5S00);
 730     MD5_STEP (MD5_Fo, d, a, b, c, w3_t[1], MD5C0d, MD5S01);
 731     MD5_STEP (MD5_Fo, c, d, a, b, w3_t[2], MD5C0e, MD5S02);
 732     MD5_STEP (MD5_Fo, b, c, d, a, w3_t[3], MD5C0f, MD5S03);
 733
 734     MD5_STEP (MD5_Go, a, b, c, d, w0_t[1], MD5C10, MD5S10);
 735     MD5_STEP (MD5_Go, d, a, b, c, w1_t[2], MD5C11, MD5S11);
 736     MD5_STEP (MD5_Go, c, d, a, b, w2_t[3], MD5C12, MD5S12);
 737     MD5_STEP (MD5_Go, b, c, d, a, w0_t[0], MD5C13, MD5S13);
 738     MD5_STEP (MD5_Go, a, b, c, d, w1_t[1], MD5C14, MD5S10);
 739     MD5_STEP (MD5_Go, d, a, b, c, w2_t[2], MD5C15, MD5S11);
 740     MD5_STEP (MD5_Go, c, d, a, b, w3_t[3], MD5C16, MD5S12);
 741     MD5_STEP (MD5_Go, b, c, d, a, w1_t[0], MD5C17, MD5S13);
 742     MD5_STEP (MD5_Go, a, b, c, d, w2_t[1], MD5C18, MD5S10);
 743     MD5_STEP (MD5_Go, d, a, b, c, w3_t[2], MD5C19, MD5S11);
 744     MD5_STEP (MD5_Go, c, d, a, b, w0_t[3], MD5C1a, MD5S12);
 745     MD5_STEP (MD5_Go, b, c, d, a, w2_t[0], MD5C1b, MD5S13);
 746     MD5_STEP (MD5_Go, a, b, c, d, w3_t[1], MD5C1c, MD5S10);
 747     MD5_STEP (MD5_Go, d, a, b, c, w0_t[2], MD5C1d, MD5S11);
 748     MD5_STEP (MD5_Go, c, d, a, b, w1_t[3], MD5C1e, MD5S12);
 749     MD5_STEP (MD5_Go, b, c, d, a, w3_t[0], MD5C1f, MD5S13);
 750
 751     MD5_STEP (MD5_H , a, b, c, d, w1_t[1], MD5C20, MD5S20);
 752     MD5_STEP (MD5_H , d, a, b, c, w2_t[0], MD5C21, MD5S21);
 753     MD5_STEP (MD5_H , c, d, a, b, w2_t[3], MD5C22, MD5S22);
 754     MD5_STEP (MD5_H , b, c, d, a, w3_t[2], MD5C23, MD5S23);
 755     MD5_STEP (MD5_H , a, b, c, d, w0_t[1], MD5C24, MD5S20);
 756     MD5_STEP (MD5_H , d, a, b, c, w1_t[0], MD5C25, MD5S21);
 757     MD5_STEP (MD5_H , c, d, a, b, w1_t[3], MD5C26, MD5S22);
 758     MD5_STEP (MD5_H , b, c, d, a, w2_t[2], MD5C27, MD5S23);
 759     MD5_STEP (MD5_H , a, b, c, d, w3_t[1], MD5C28, MD5S20);
 760     MD5_STEP (MD5_H , d, a, b, c, w0_t[0], MD5C29, MD5S21);
 761     MD5_STEP (MD5_H , c, d, a, b, w0_t[3], MD5C2a, MD5S22);
 762     MD5_STEP (MD5_H , b, c, d, a, w1_t[2], MD5C2b, MD5S23);
 763     MD5_STEP (MD5_H , a, b, c, d, w2_t[1], MD5C2c, MD5S20);
 764     MD5_STEP (MD5_H , d, a, b, c, w3_t[0], MD5C2d, MD5S21);
 765     MD5_STEP (MD5_H , c, d, a, b, w3_t[3], MD5C2e, MD5S22);
 766     MD5_STEP (MD5_H , b, c, d, a, w0_t[2], MD5C2f, MD5S23);
 767
 768     MD5_STEP (MD5_I , a, b, c, d, w0_t[0], MD5C30, MD5S30);
 769     MD5_STEP (MD5_I , d, a, b, c, w1_t[3], MD5C31, MD5S31);
 770     MD5_STEP (MD5_I , c, d, a, b, w3_t[2], MD5C32, MD5S32);
 771     MD5_STEP (MD5_I , b, c, d, a, w1_t[1], MD5C33, MD5S33);
 772     MD5_STEP (MD5_I , a, b, c, d, w3_t[0], MD5C34, MD5S30);
 773     MD5_STEP (MD5_I , d, a, b, c, w0_t[3], MD5C35, MD5S31);
 774     MD5_STEP (MD5_I , c, d, a, b, w2_t[2], MD5C36, MD5S32);
 775     MD5_STEP (MD5_I , b, c, d, a, w0_t[1], MD5C37, MD5S33);
 776     MD5_STEP (MD5_I , a, b, c, d, w2_t[0], MD5C38, MD5S30);
 777     MD5_STEP (MD5_I , d, a, b, c, w3_t[3], MD5C39, MD5S31);
 778     MD5_STEP (MD5_I , c, d, a, b, w1_t[2], MD5C3a, MD5S32);
 779     MD5_STEP (MD5_I , b, c, d, a, w3_t[1], MD5C3b, MD5S33);
 780     MD5_STEP (MD5_I , a, b, c, d, w1_t[0], MD5C3c, MD5S30);
 781     MD5_STEP (MD5_I , d, a, b, c, w2_t[3], MD5C3d, MD5S31);
 782     MD5_STEP (MD5_I , c, d, a, b, w0_t[2], MD5C3e, MD5S32);
 783     MD5_STEP (MD5_I , b, c, d, a, w2_t[1], MD5C3f, MD5S33);
 784
 785     a += r_a;
 786     b += r_b;
 787     c += r_c;
 788     d += r_d;
 789
 790     const u32x r0 = a;
 791     const u32x r1 = d;
 792     const u32x r2 = c;
 793     const u32x r3 = b;
 794
 795     #include VECT_COMPARE_S
 796   }
 797 }
 798
 799 extern "C" __global__ void __launch_bounds__ (256, 1) m11000_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)
 800 {
 801 }
 802
 803 extern "C" __global__ void __launch_bounds__ (256, 1) m11000_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)
 804 {
 805 }