nv/m02410_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MD5_
   7 #define _SCALAR_
   8
   9 #include "include/constants.h"
  10 #include "include/kernel_vendor.h"
  11
  12 #ifdef  VLIW1
  13 #define VECT_SIZE1
  14 #endif
  15
  16 #ifdef  VLIW2
  17 #define VECT_SIZE4
  18 #endif
  19
  20 #define DGST_R0 0
  21 #define DGST_R1 3
  22 #define DGST_R2 2
  23 #define DGST_R3 1
  24
  25 #include "include/kernel_functions.c"
  26 #include "types_nv.c"
  27 #include "common_nv.c"
  28
  29 #ifdef  VECT_SIZE1
  30 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
  31 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
  32 #endif
  33
  34 #ifdef  VECT_SIZE2
  35 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
  36 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
  37 #endif
  38
  39 #ifdef  VECT_SIZE4
  40 #define VECT_COMPARE_S "check_single_vect4_comp4_warp.c"
  41 #define VECT_COMPARE_M "check_multi_vect4_comp4_warp.c"
  42 #endif
  43
  44 __device__ __constant__ u32x c_bfs[1024];
  45
  46 __device__ static void m02410m (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
  47 {
  48   /**
  49    * modifier
  50    */
  51
  52   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  53   const u32 lid = threadIdx.x;
  54
  55   /**
  56    * salt
  57    */
  58
  59   u32 salt_buf0[4];
  60
  61   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
  62   salt_buf0[1] = 0;
  63   salt_buf0[2] = 0;
  64   salt_buf0[3] = 0;
  65
  66   u32 salt_buf1[4];
  67
  68   salt_buf1[0] = 0;
  69   salt_buf1[1] = 0;
  70   salt_buf1[2] = 0;
  71   salt_buf1[3] = 0;
  72
  73   u32 salt_buf2[4];
  74
  75   salt_buf2[0] = 0;
  76   salt_buf2[1] = 0;
  77   salt_buf2[2] = 0;
  78   salt_buf2[3] = 0;
  79
  80   u32 salt_buf3[4];
  81
  82   salt_buf3[0] = 0;
  83   salt_buf3[1] = 0;
  84   salt_buf3[2] = 0;
  85   salt_buf3[3] = 0;
  86
  87   switch_buffer_by_offset (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
  88
  89   w[0] |= salt_buf0[0];
  90   w[1] |= salt_buf0[1];
  91   w[2] |= salt_buf0[2];
  92   w[3] |= salt_buf0[3];
  93
  94   const u32 salt_len = (salt_bufs[salt_pos].salt_len < 4) ? salt_bufs[salt_pos].salt_len : 4;
  95
  96   const u32 pw_salt_len = pw_len + salt_len;
  97
  98   truncate_block (w, pw_salt_len);
  99
 100   /**
 101    * algorithm specific
 102    */
 103
 104   w[ 4] = 0x80;
 105   w[14] = 16 * 8;
 106
 107   /**
 108    * base
 109    */
 110
 111   const u32 F_w0c00 =    0  + MD5C00;
 112   const u32 F_w1c01 = w[ 1] + MD5C01;
 113   const u32 F_w2c02 = w[ 2] + MD5C02;
 114   const u32 F_w3c03 = w[ 3] + MD5C03;
 115   const u32 F_w4c04 = w[ 4] + MD5C04;
 116   const u32 F_w5c05 = w[ 5] + MD5C05;
 117   const u32 F_w6c06 = w[ 6] + MD5C06;
 118   const u32 F_w7c07 = w[ 7] + MD5C07;
 119   const u32 F_w8c08 = w[ 8] + MD5C08;
 120   const u32 F_w9c09 = w[ 9] + MD5C09;
 121   const u32 F_wac0a = w[10] + MD5C0a;
 122   const u32 F_wbc0b = w[11] + MD5C0b;
 123   const u32 F_wcc0c = w[12] + MD5C0c;
 124   const u32 F_wdc0d = w[13] + MD5C0d;
 125   const u32 F_wec0e = w[14] + MD5C0e;
 126   const u32 F_wfc0f = w[15] + MD5C0f;
 127
 128   const u32 G_w1c10 = w[ 1] + MD5C10;
 129   const u32 G_w6c11 = w[ 6] + MD5C11;
 130   const u32 G_wbc12 = w[11] + MD5C12;
 131   const u32 G_w0c13 =    0  + MD5C13;
 132   const u32 G_w5c14 = w[ 5] + MD5C14;
 133   const u32 G_wac15 = w[10] + MD5C15;
 134   const u32 G_wfc16 = w[15] + MD5C16;
 135   const u32 G_w4c17 = w[ 4] + MD5C17;
 136   const u32 G_w9c18 = w[ 9] + MD5C18;
 137   const u32 G_wec19 = w[14] + MD5C19;
 138   const u32 G_w3c1a = w[ 3] + MD5C1a;
 139   const u32 G_w8c1b = w[ 8] + MD5C1b;
 140   const u32 G_wdc1c = w[13] + MD5C1c;
 141   const u32 G_w2c1d = w[ 2] + MD5C1d;
 142   const u32 G_w7c1e = w[ 7] + MD5C1e;
 143   const u32 G_wcc1f = w[12] + MD5C1f;
 144
 145   const u32 H_w5c20 = w[ 5] + MD5C20;
 146   const u32 H_w8c21 = w[ 8] + MD5C21;
 147   const u32 H_wbc22 = w[11] + MD5C22;
 148   const u32 H_wec23 = w[14] + MD5C23;
 149   const u32 H_w1c24 = w[ 1] + MD5C24;
 150   const u32 H_w4c25 = w[ 4] + MD5C25;
 151   const u32 H_w7c26 = w[ 7] + MD5C26;
 152   const u32 H_wac27 = w[10] + MD5C27;
 153   const u32 H_wdc28 = w[13] + MD5C28;
 154   const u32 H_w0c29 =    0  + MD5C29;
 155   const u32 H_w3c2a = w[ 3] + MD5C2a;
 156   const u32 H_w6c2b = w[ 6] + MD5C2b;
 157   const u32 H_w9c2c = w[ 9] + MD5C2c;
 158   const u32 H_wcc2d = w[12] + MD5C2d;
 159   const u32 H_wfc2e = w[15] + MD5C2e;
 160   const u32 H_w2c2f = w[ 2] + MD5C2f;
 161
 162   const u32 I_w0c30 =    0  + MD5C30;
 163   const u32 I_w7c31 = w[ 7] + MD5C31;
 164   const u32 I_wec32 = w[14] + MD5C32;
 165   const u32 I_w5c33 = w[ 5] + MD5C33;
 166   const u32 I_wcc34 = w[12] + MD5C34;
 167   const u32 I_w3c35 = w[ 3] + MD5C35;
 168   const u32 I_wac36 = w[10] + MD5C36;
 169   const u32 I_w1c37 = w[ 1] + MD5C37;
 170   const u32 I_w8c38 = w[ 8] + MD5C38;
 171   const u32 I_wfc39 = w[15] + MD5C39;
 172   const u32 I_w6c3a = w[ 6] + MD5C3a;
 173   const u32 I_wdc3b = w[13] + MD5C3b;
 174   const u32 I_w4c3c = w[ 4] + MD5C3c;
 175   const u32 I_wbc3d = w[11] + MD5C3d;
 176   const u32 I_w2c3e = w[ 2] + MD5C3e;
 177   const u32 I_w9c3f = w[ 9] + MD5C3f;
 178
 179   /**
 180    * loop
 181    */
 182
 183   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 184
 185   u32x w0l = w[0];
 186
 187   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 188   {
 189     const u32x w0r = c_bfs[il_pos];
 190
 191     const u32x w0 = w0l | w0r;
 192
 193     u32x tmp2;
 194
 195     u32x a = MD5M_A;
 196     u32x b = MD5M_B;
 197     u32x c = MD5M_C;
 198     u32x d = MD5M_D;
 199
 200     MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00);
 201     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w1c01, MD5S01);
 202     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w2c02, MD5S02);
 203     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w3c03, MD5S03);
 204     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w4c04, MD5S00);
 205     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w5c05, MD5S01);
 206     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w6c06, MD5S02);
 207     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w7c07, MD5S03);
 208     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w8c08, MD5S00);
 209     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w9c09, MD5S01);
 210     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wac0a, MD5S02);
 211     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wbc0b, MD5S03);
 212     MD5_STEP0(MD5_Fo, a, b, c, d,     F_wcc0c, MD5S00);
 213     MD5_STEP0(MD5_Fo, d, a, b, c,     F_wdc0d, MD5S01);
 214     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wec0e, MD5S02);
 215     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wfc0f, MD5S03);
 216
 217     MD5_STEP0(MD5_Go, a, b, c, d,     G_w1c10, MD5S10);
 218     MD5_STEP0(MD5_Go, d, a, b, c,     G_w6c11, MD5S11);
 219     MD5_STEP0(MD5_Go, c, d, a, b,     G_wbc12, MD5S12);
 220     MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13);
 221     MD5_STEP0(MD5_Go, a, b, c, d,     G_w5c14, MD5S10);
 222     MD5_STEP0(MD5_Go, d, a, b, c,     G_wac15, MD5S11);
 223     MD5_STEP0(MD5_Go, c, d, a, b,     G_wfc16, MD5S12);
 224     MD5_STEP0(MD5_Go, b, c, d, a,     G_w4c17, MD5S13);
 225     MD5_STEP0(MD5_Go, a, b, c, d,     G_w9c18, MD5S10);
 226     MD5_STEP0(MD5_Go, d, a, b, c,     G_wec19, MD5S11);
 227     MD5_STEP0(MD5_Go, c, d, a, b,     G_w3c1a, MD5S12);
 228     MD5_STEP0(MD5_Go, b, c, d, a,     G_w8c1b, MD5S13);
 229     MD5_STEP0(MD5_Go, a, b, c, d,     G_wdc1c, MD5S10);
 230     MD5_STEP0(MD5_Go, d, a, b, c,     G_w2c1d, MD5S11);
 231     MD5_STEP0(MD5_Go, c, d, a, b,     G_w7c1e, MD5S12);
 232     MD5_STEP0(MD5_Go, b, c, d, a,     G_wcc1f, MD5S13);
 233
 234     MD5_STEP0(MD5_H1, a, b, c, d,     H_w5c20, MD5S20);
 235     MD5_STEP0(MD5_H2, d, a, b, c,     H_w8c21, MD5S21);
 236     MD5_STEP0(MD5_H1, c, d, a, b,     H_wbc22, MD5S22);
 237     MD5_STEP0(MD5_H2, b, c, d, a,     H_wec23, MD5S23);
 238     MD5_STEP0(MD5_H1, a, b, c, d,     H_w1c24, MD5S20);
 239     MD5_STEP0(MD5_H2, d, a, b, c,     H_w4c25, MD5S21);
 240     MD5_STEP0(MD5_H1, c, d, a, b,     H_w7c26, MD5S22);
 241     MD5_STEP0(MD5_H2, b, c, d, a,     H_wac27, MD5S23);
 242     MD5_STEP0(MD5_H1, a, b, c, d,     H_wdc28, MD5S20);
 243     MD5_STEP (MD5_H2, d, a, b, c, w0, H_w0c29, MD5S21);
 244     MD5_STEP0(MD5_H1, c, d, a, b,     H_w3c2a, MD5S22);
 245     MD5_STEP0(MD5_H2, b, c, d, a,     H_w6c2b, MD5S23);
 246     MD5_STEP0(MD5_H1, a, b, c, d,     H_w9c2c, MD5S20);
 247     MD5_STEP0(MD5_H2, d, a, b, c,     H_wcc2d, MD5S21);
 248     MD5_STEP0(MD5_H1, c, d, a, b,     H_wfc2e, MD5S22);
 249     MD5_STEP0(MD5_H2, b, c, d, a,     H_w2c2f, MD5S23);
 250
 251     MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30);
 252     MD5_STEP0(MD5_I , d, a, b, c,     I_w7c31, MD5S31);
 253     MD5_STEP0(MD5_I , c, d, a, b,     I_wec32, MD5S32);
 254     MD5_STEP0(MD5_I , b, c, d, a,     I_w5c33, MD5S33);
 255     MD5_STEP0(MD5_I , a, b, c, d,     I_wcc34, MD5S30);
 256     MD5_STEP0(MD5_I , d, a, b, c,     I_w3c35, MD5S31);
 257     MD5_STEP0(MD5_I , c, d, a, b,     I_wac36, MD5S32);
 258     MD5_STEP0(MD5_I , b, c, d, a,     I_w1c37, MD5S33);
 259     MD5_STEP0(MD5_I , a, b, c, d,     I_w8c38, MD5S30);
 260     MD5_STEP0(MD5_I , d, a, b, c,     I_wfc39, MD5S31);
 261     MD5_STEP0(MD5_I , c, d, a, b,     I_w6c3a, MD5S32);
 262     MD5_STEP0(MD5_I , b, c, d, a,     I_wdc3b, MD5S33);
 263     MD5_STEP0(MD5_I , a, b, c, d,     I_w4c3c, MD5S30);
 264     MD5_STEP0(MD5_I , d, a, b, c,     I_wbc3d, MD5S31);
 265     MD5_STEP0(MD5_I , c, d, a, b,     I_w2c3e, MD5S32);
 266     MD5_STEP0(MD5_I , b, c, d, a,     I_w9c3f, MD5S33);
 267
 268     a &= 0x00ffffff;
 269     d &= 0x00ffffff;
 270     c &= 0x00ffffff;
 271     b &= 0x00ffffff;
 272
 273     const u32x r0 = a;
 274     const u32x r1 = d;
 275     const u32x r2 = c;
 276     const u32x r3 = b;
 277
 278     #include VECT_COMPARE_M
 279   }
 280 }
 281
 282 __device__ static void m02410s (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 283 {
 284   /**
 285    * modifier
 286    */
 287
 288   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 289   const u32 lid = threadIdx.x;
 290
 291   /**
 292    * salt
 293    */
 294
 295   u32 salt_buf0[4];
 296
 297   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 298   salt_buf0[1] = 0;
 299   salt_buf0[2] = 0;
 300   salt_buf0[3] = 0;
 301
 302   u32 salt_buf1[4];
 303
 304   salt_buf1[0] = 0;
 305   salt_buf1[1] = 0;
 306   salt_buf1[2] = 0;
 307   salt_buf1[3] = 0;
 308
 309   u32 salt_buf2[4];
 310
 311   salt_buf2[0] = 0;
 312   salt_buf2[1] = 0;
 313   salt_buf2[2] = 0;
 314   salt_buf2[3] = 0;
 315
 316   u32 salt_buf3[4];
 317
 318   salt_buf3[0] = 0;
 319   salt_buf3[1] = 0;
 320   salt_buf3[2] = 0;
 321   salt_buf3[3] = 0;
 322
 323   switch_buffer_by_offset (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
 324
 325   w[0] |= salt_buf0[0];
 326   w[1] |= salt_buf0[1];
 327   w[2] |= salt_buf0[2];
 328   w[3] |= salt_buf0[3];
 329
 330   const u32 salt_len = (salt_bufs[salt_pos].salt_len < 4) ? salt_bufs[salt_pos].salt_len : 4;
 331
 332   const u32 pw_salt_len = pw_len + salt_len;
 333
 334   truncate_block (w, pw_salt_len);
 335
 336   /**
 337    * algorithm specific
 338    */
 339
 340   w[ 4] = 0x80;
 341   w[14] = 16 * 8;
 342
 343   /**
 344    * base
 345    */
 346
 347   const u32 F_w0c00 =    0  + MD5C00;
 348   const u32 F_w1c01 = w[ 1] + MD5C01;
 349   const u32 F_w2c02 = w[ 2] + MD5C02;
 350   const u32 F_w3c03 = w[ 3] + MD5C03;
 351   const u32 F_w4c04 = w[ 4] + MD5C04;
 352   const u32 F_w5c05 = w[ 5] + MD5C05;
 353   const u32 F_w6c06 = w[ 6] + MD5C06;
 354   const u32 F_w7c07 = w[ 7] + MD5C07;
 355   const u32 F_w8c08 = w[ 8] + MD5C08;
 356   const u32 F_w9c09 = w[ 9] + MD5C09;
 357   const u32 F_wac0a = w[10] + MD5C0a;
 358   const u32 F_wbc0b = w[11] + MD5C0b;
 359   const u32 F_wcc0c = w[12] + MD5C0c;
 360   const u32 F_wdc0d = w[13] + MD5C0d;
 361   const u32 F_wec0e = w[14] + MD5C0e;
 362   const u32 F_wfc0f = w[15] + MD5C0f;
 363
 364   const u32 G_w1c10 = w[ 1] + MD5C10;
 365   const u32 G_w6c11 = w[ 6] + MD5C11;
 366   const u32 G_wbc12 = w[11] + MD5C12;
 367   const u32 G_w0c13 =    0  + MD5C13;
 368   const u32 G_w5c14 = w[ 5] + MD5C14;
 369   const u32 G_wac15 = w[10] + MD5C15;
 370   const u32 G_wfc16 = w[15] + MD5C16;
 371   const u32 G_w4c17 = w[ 4] + MD5C17;
 372   const u32 G_w9c18 = w[ 9] + MD5C18;
 373   const u32 G_wec19 = w[14] + MD5C19;
 374   const u32 G_w3c1a = w[ 3] + MD5C1a;
 375   const u32 G_w8c1b = w[ 8] + MD5C1b;
 376   const u32 G_wdc1c = w[13] + MD5C1c;
 377   const u32 G_w2c1d = w[ 2] + MD5C1d;
 378   const u32 G_w7c1e = w[ 7] + MD5C1e;
 379   const u32 G_wcc1f = w[12] + MD5C1f;
 380
 381   const u32 H_w5c20 = w[ 5] + MD5C20;
 382   const u32 H_w8c21 = w[ 8] + MD5C21;
 383   const u32 H_wbc22 = w[11] + MD5C22;
 384   const u32 H_wec23 = w[14] + MD5C23;
 385   const u32 H_w1c24 = w[ 1] + MD5C24;
 386   const u32 H_w4c25 = w[ 4] + MD5C25;
 387   const u32 H_w7c26 = w[ 7] + MD5C26;
 388   const u32 H_wac27 = w[10] + MD5C27;
 389   const u32 H_wdc28 = w[13] + MD5C28;
 390   const u32 H_w0c29 =    0  + MD5C29;
 391   const u32 H_w3c2a = w[ 3] + MD5C2a;
 392   const u32 H_w6c2b = w[ 6] + MD5C2b;
 393   const u32 H_w9c2c = w[ 9] + MD5C2c;
 394   const u32 H_wcc2d = w[12] + MD5C2d;
 395   const u32 H_wfc2e = w[15] + MD5C2e;
 396   const u32 H_w2c2f = w[ 2] + MD5C2f;
 397
 398   const u32 I_w0c30 =    0  + MD5C30;
 399   const u32 I_w7c31 = w[ 7] + MD5C31;
 400   const u32 I_wec32 = w[14] + MD5C32;
 401   const u32 I_w5c33 = w[ 5] + MD5C33;
 402   const u32 I_wcc34 = w[12] + MD5C34;
 403   const u32 I_w3c35 = w[ 3] + MD5C35;
 404   const u32 I_wac36 = w[10] + MD5C36;
 405   const u32 I_w1c37 = w[ 1] + MD5C37;
 406   const u32 I_w8c38 = w[ 8] + MD5C38;
 407   const u32 I_wfc39 = w[15] + MD5C39;
 408   const u32 I_w6c3a = w[ 6] + MD5C3a;
 409   const u32 I_wdc3b = w[13] + MD5C3b;
 410   const u32 I_w4c3c = w[ 4] + MD5C3c;
 411   const u32 I_wbc3d = w[11] + MD5C3d;
 412   const u32 I_w2c3e = w[ 2] + MD5C3e;
 413   const u32 I_w9c3f = w[ 9] + MD5C3f;
 414
 415   /**
 416    * digest
 417    */
 418
 419   const u32 search[4] =
 420   {
 421     digests_buf[digests_offset].digest_buf[DGST_R0],
 422     digests_buf[digests_offset].digest_buf[DGST_R1],
 423     digests_buf[digests_offset].digest_buf[DGST_R2],
 424     digests_buf[digests_offset].digest_buf[DGST_R3]
 425   };
 426
 427   /**
 428    * loop
 429    */
 430
 431   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 432
 433   u32x w0l = w[0];
 434
 435   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 436   {
 437     const u32x w0r = c_bfs[il_pos];
 438
 439     const u32x w0 = w0l | w0r;
 440
 441     u32x tmp2;
 442
 443     u32x a = MD5M_A;
 444     u32x b = MD5M_B;
 445     u32x c = MD5M_C;
 446     u32x d = MD5M_D;
 447
 448     MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00);
 449     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w1c01, MD5S01);
 450     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w2c02, MD5S02);
 451     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w3c03, MD5S03);
 452     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w4c04, MD5S00);
 453     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w5c05, MD5S01);
 454     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w6c06, MD5S02);
 455     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w7c07, MD5S03);
 456     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w8c08, MD5S00);
 457     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w9c09, MD5S01);
 458     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wac0a, MD5S02);
 459     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wbc0b, MD5S03);
 460     MD5_STEP0(MD5_Fo, a, b, c, d,     F_wcc0c, MD5S00);
 461     MD5_STEP0(MD5_Fo, d, a, b, c,     F_wdc0d, MD5S01);
 462     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wec0e, MD5S02);
 463     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wfc0f, MD5S03);
 464
 465     MD5_STEP0(MD5_Go, a, b, c, d,     G_w1c10, MD5S10);
 466     MD5_STEP0(MD5_Go, d, a, b, c,     G_w6c11, MD5S11);
 467     MD5_STEP0(MD5_Go, c, d, a, b,     G_wbc12, MD5S12);
 468     MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13);
 469     MD5_STEP0(MD5_Go, a, b, c, d,     G_w5c14, MD5S10);
 470     MD5_STEP0(MD5_Go, d, a, b, c,     G_wac15, MD5S11);
 471     MD5_STEP0(MD5_Go, c, d, a, b,     G_wfc16, MD5S12);
 472     MD5_STEP0(MD5_Go, b, c, d, a,     G_w4c17, MD5S13);
 473     MD5_STEP0(MD5_Go, a, b, c, d,     G_w9c18, MD5S10);
 474     MD5_STEP0(MD5_Go, d, a, b, c,     G_wec19, MD5S11);
 475     MD5_STEP0(MD5_Go, c, d, a, b,     G_w3c1a, MD5S12);
 476     MD5_STEP0(MD5_Go, b, c, d, a,     G_w8c1b, MD5S13);
 477     MD5_STEP0(MD5_Go, a, b, c, d,     G_wdc1c, MD5S10);
 478     MD5_STEP0(MD5_Go, d, a, b, c,     G_w2c1d, MD5S11);
 479     MD5_STEP0(MD5_Go, c, d, a, b,     G_w7c1e, MD5S12);
 480     MD5_STEP0(MD5_Go, b, c, d, a,     G_wcc1f, MD5S13);
 481
 482     MD5_STEP0(MD5_H1, a, b, c, d,     H_w5c20, MD5S20);
 483     MD5_STEP0(MD5_H2, d, a, b, c,     H_w8c21, MD5S21);
 484     MD5_STEP0(MD5_H1, c, d, a, b,     H_wbc22, MD5S22);
 485     MD5_STEP0(MD5_H2, b, c, d, a,     H_wec23, MD5S23);
 486     MD5_STEP0(MD5_H1, a, b, c, d,     H_w1c24, MD5S20);
 487     MD5_STEP0(MD5_H2, d, a, b, c,     H_w4c25, MD5S21);
 488     MD5_STEP0(MD5_H1, c, d, a, b,     H_w7c26, MD5S22);
 489     MD5_STEP0(MD5_H2, b, c, d, a,     H_wac27, MD5S23);
 490     MD5_STEP0(MD5_H1, a, b, c, d,     H_wdc28, MD5S20);
 491     MD5_STEP (MD5_H2, d, a, b, c, w0, H_w0c29, MD5S21);
 492     MD5_STEP0(MD5_H1, c, d, a, b,     H_w3c2a, MD5S22);
 493     MD5_STEP0(MD5_H2, b, c, d, a,     H_w6c2b, MD5S23);
 494     MD5_STEP0(MD5_H1, a, b, c, d,     H_w9c2c, MD5S20);
 495     MD5_STEP0(MD5_H2, d, a, b, c,     H_wcc2d, MD5S21);
 496     MD5_STEP0(MD5_H1, c, d, a, b,     H_wfc2e, MD5S22);
 497     MD5_STEP0(MD5_H2, b, c, d, a,     H_w2c2f, MD5S23);
 498
 499     MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30);
 500     MD5_STEP0(MD5_I , d, a, b, c,     I_w7c31, MD5S31);
 501     MD5_STEP0(MD5_I , c, d, a, b,     I_wec32, MD5S32);
 502     MD5_STEP0(MD5_I , b, c, d, a,     I_w5c33, MD5S33);
 503     MD5_STEP0(MD5_I , a, b, c, d,     I_wcc34, MD5S30);
 504     MD5_STEP0(MD5_I , d, a, b, c,     I_w3c35, MD5S31);
 505     MD5_STEP0(MD5_I , c, d, a, b,     I_wac36, MD5S32);
 506     MD5_STEP0(MD5_I , b, c, d, a,     I_w1c37, MD5S33);
 507     MD5_STEP0(MD5_I , a, b, c, d,     I_w8c38, MD5S30);
 508     MD5_STEP0(MD5_I , d, a, b, c,     I_wfc39, MD5S31);
 509     MD5_STEP0(MD5_I , c, d, a, b,     I_w6c3a, MD5S32);
 510     MD5_STEP0(MD5_I , b, c, d, a,     I_wdc3b, MD5S33);
 511     MD5_STEP0(MD5_I , a, b, c, d,     I_w4c3c, MD5S30);
 512
 513     bool q_cond = ((a & 0x00ffffff) != search[0]);
 514
 515     if (q_cond) continue;
 516
 517     MD5_STEP0(MD5_I , d, a, b, c,     I_wbc3d, MD5S31);
 518     MD5_STEP0(MD5_I , c, d, a, b,     I_w2c3e, MD5S32);
 519     MD5_STEP0(MD5_I , b, c, d, a,     I_w9c3f, MD5S33);
 520
 521     a &= 0x00ffffff;
 522     d &= 0x00ffffff;
 523     c &= 0x00ffffff;
 524     b &= 0x00ffffff;
 525
 526     const u32x r0 = a;
 527     const u32x r1 = d;
 528     const u32x r2 = c;
 529     const u32x r3 = b;
 530
 531     #include VECT_COMPARE_S
 532   }
 533 }
 534
 535 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 536 {
 537   /**
 538    * base
 539    */
 540
 541   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 542
 543   if (gid >= gid_max) return;
 544
 545   u32 w[16];
 546
 547   w[ 0] = pws[gid].i[ 0];
 548   w[ 1] = pws[gid].i[ 1];
 549   w[ 2] = pws[gid].i[ 2];
 550   w[ 3] = pws[gid].i[ 3];
 551   w[ 4] = 0;
 552   w[ 5] = 0;
 553   w[ 6] = 0;
 554   w[ 7] = 0;
 555   w[ 8] = 0;
 556   w[ 9] = 0;
 557   w[10] = 0;
 558   w[11] = 0;
 559   w[12] = 0;
 560   w[13] = 0;
 561   w[14] = 0;
 562   w[15] = 0;
 563
 564   const u32 pw_len = pws[gid].pw_len;
 565
 566   /**
 567    * main
 568    */
 569
 570   m02410m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 571 }
 572
 573 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 574 {
 575 }
 576
 577 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 578 {
 579 }
 580
 581 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 582 {
 583   /**
 584    * base
 585    */
 586
 587   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 588
 589   if (gid >= gid_max) return;
 590
 591   u32 w[16];
 592
 593   w[ 0] = pws[gid].i[ 0];
 594   w[ 1] = pws[gid].i[ 1];
 595   w[ 2] = pws[gid].i[ 2];
 596   w[ 3] = pws[gid].i[ 3];
 597   w[ 4] = 0;
 598   w[ 5] = 0;
 599   w[ 6] = 0;
 600   w[ 7] = 0;
 601   w[ 8] = 0;
 602   w[ 9] = 0;
 603   w[10] = 0;
 604   w[11] = 0;
 605   w[12] = 0;
 606   w[13] = 0;
 607   w[14] = 0;
 608   w[15] = 0;
 609
 610   const u32 pw_len = pws[gid].pw_len;
 611
 612   /**
 613    * main
 614    */
 615
 616   m02410s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_scryptV_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, bfs_cnt, digests_cnt, digests_offset);
 617 }
 618
 619 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 620 {
 621 }
 622
 623 extern "C" __global__ void __launch_bounds__ (256, 1) m02410_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 624 {
 625 }