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