nv/m01000_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MD4_
   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 #define MD4_STEP_REV(f,a,b,c,d,x,t,s)   \
  45 {                                       \
  46   a  = rotr32 (a, s);                   \
  47   a -= f (b, c, d);                     \
  48   a -= x;                               \
  49   a -= t;                               \
  50 }
  51
  52 #define MD4_STEP_REV1(f,a,b,c,d,x,t,s)  \
  53 {                                       \
  54   a  = rotr32 (a, s);                   \
  55   a -= x;                               \
  56   a -= t;                               \
  57 }
  58
  59 __device__ __constant__ u32x c_bfs[1024];
  60
  61 __device__ static void m01000m (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, 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)
  62 {
  63   /**
  64    * modifier
  65    */
  66
  67   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  68   const u32 lid = threadIdx.x;
  69
  70   /**
  71    * base
  72    */
  73
  74   const u32 F_w0c00 =     0 + MD4C00;
  75   const u32 F_w1c00 = w[ 1] + MD4C00;
  76   const u32 F_w2c00 = w[ 2] + MD4C00;
  77   const u32 F_w3c00 = w[ 3] + MD4C00;
  78   const u32 F_w4c00 = w[ 4] + MD4C00;
  79   const u32 F_w5c00 = w[ 5] + MD4C00;
  80   const u32 F_w6c00 = w[ 6] + MD4C00;
  81   const u32 F_w7c00 = w[ 7] + MD4C00;
  82   const u32 F_w8c00 = w[ 8] + MD4C00;
  83   const u32 F_w9c00 = w[ 9] + MD4C00;
  84   const u32 F_wac00 = w[10] + MD4C00;
  85   const u32 F_wbc00 = w[11] + MD4C00;
  86   const u32 F_wcc00 = w[12] + MD4C00;
  87   const u32 F_wdc00 = w[13] + MD4C00;
  88   const u32 F_wec00 = w[14] + MD4C00;
  89   const u32 F_wfc00 = w[15] + MD4C00;
  90
  91   const u32 G_w0c01 =     0 + MD4C01;
  92   const u32 G_w4c01 = w[ 4] + MD4C01;
  93   const u32 G_w8c01 = w[ 8] + MD4C01;
  94   const u32 G_wcc01 = w[12] + MD4C01;
  95   const u32 G_w1c01 = w[ 1] + MD4C01;
  96   const u32 G_w5c01 = w[ 5] + MD4C01;
  97   const u32 G_w9c01 = w[ 9] + MD4C01;
  98   const u32 G_wdc01 = w[13] + MD4C01;
  99   const u32 G_w2c01 = w[ 2] + MD4C01;
 100   const u32 G_w6c01 = w[ 6] + MD4C01;
 101   const u32 G_wac01 = w[10] + MD4C01;
 102   const u32 G_wec01 = w[14] + MD4C01;
 103   const u32 G_w3c01 = w[ 3] + MD4C01;
 104   const u32 G_w7c01 = w[ 7] + MD4C01;
 105   const u32 G_wbc01 = w[11] + MD4C01;
 106   const u32 G_wfc01 = w[15] + MD4C01;
 107
 108   const u32 H_w0c02 =     0 + MD4C02;
 109   const u32 H_w8c02 = w[ 8] + MD4C02;
 110   const u32 H_w4c02 = w[ 4] + MD4C02;
 111   const u32 H_wcc02 = w[12] + MD4C02;
 112   const u32 H_w2c02 = w[ 2] + MD4C02;
 113   const u32 H_wac02 = w[10] + MD4C02;
 114   const u32 H_w6c02 = w[ 6] + MD4C02;
 115   const u32 H_wec02 = w[14] + MD4C02;
 116   const u32 H_w1c02 = w[ 1] + MD4C02;
 117   const u32 H_w9c02 = w[ 9] + MD4C02;
 118   const u32 H_w5c02 = w[ 5] + MD4C02;
 119   const u32 H_wdc02 = w[13] + MD4C02;
 120   const u32 H_w3c02 = w[ 3] + MD4C02;
 121   const u32 H_wbc02 = w[11] + MD4C02;
 122   const u32 H_w7c02 = w[ 7] + MD4C02;
 123   const u32 H_wfc02 = w[15] + MD4C02;
 124
 125   /**
 126    * loop
 127    */
 128
 129   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 130
 131   u32x w0l = w[0];
 132
 133   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 134   {
 135     const u32x w0r = c_bfs[il_pos];
 136
 137     const u32x w0 = w0l | w0r;
 138
 139     u32x tmp2;
 140
 141     u32x a = MD4M_A;
 142     u32x b = MD4M_B;
 143     u32x c = MD4M_C;
 144     u32x d = MD4M_D;
 145
 146     MD4_STEP (MD4_Fo, a, b, c, d, w0, F_w0c00, MD4S00);
 147     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w1c00, MD4S01);
 148     MD4_STEP0(MD4_Fo, c, d, a, b,     F_w2c00, MD4S02);
 149     MD4_STEP0(MD4_Fo, b, c, d, a,     F_w3c00, MD4S03);
 150     MD4_STEP0(MD4_Fo, a, b, c, d,     F_w4c00, MD4S00);
 151     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w5c00, MD4S01);
 152     MD4_STEP0(MD4_Fo, c, d, a, b,     F_w6c00, MD4S02);
 153     MD4_STEP0(MD4_Fo, b, c, d, a,     F_w7c00, MD4S03);
 154     MD4_STEP0(MD4_Fo, a, b, c, d,     F_w8c00, MD4S00);
 155     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w9c00, MD4S01);
 156     MD4_STEP0(MD4_Fo, c, d, a, b,     F_wac00, MD4S02);
 157     MD4_STEP0(MD4_Fo, b, c, d, a,     F_wbc00, MD4S03);
 158     MD4_STEP0(MD4_Fo, a, b, c, d,     F_wcc00, MD4S00);
 159     MD4_STEP0(MD4_Fo, d, a, b, c,     F_wdc00, MD4S01);
 160     MD4_STEP0(MD4_Fo, c, d, a, b,     F_wec00, MD4S02);
 161     MD4_STEP0(MD4_Fo, b, c, d, a,     F_wfc00, MD4S03);
 162
 163     MD4_STEP (MD4_Go, a, b, c, d, w0, G_w0c01, MD4S10);
 164     MD4_STEP0(MD4_Go, d, a, b, c,     G_w4c01, MD4S11);
 165     MD4_STEP0(MD4_Go, c, d, a, b,     G_w8c01, MD4S12);
 166     MD4_STEP0(MD4_Go, b, c, d, a,     G_wcc01, MD4S13);
 167     MD4_STEP0(MD4_Go, a, b, c, d,     G_w1c01, MD4S10);
 168     MD4_STEP0(MD4_Go, d, a, b, c,     G_w5c01, MD4S11);
 169     MD4_STEP0(MD4_Go, c, d, a, b,     G_w9c01, MD4S12);
 170     MD4_STEP0(MD4_Go, b, c, d, a,     G_wdc01, MD4S13);
 171     MD4_STEP0(MD4_Go, a, b, c, d,     G_w2c01, MD4S10);
 172     MD4_STEP0(MD4_Go, d, a, b, c,     G_w6c01, MD4S11);
 173     MD4_STEP0(MD4_Go, c, d, a, b,     G_wac01, MD4S12);
 174     MD4_STEP0(MD4_Go, b, c, d, a,     G_wec01, MD4S13);
 175     MD4_STEP0(MD4_Go, a, b, c, d,     G_w3c01, MD4S10);
 176     MD4_STEP0(MD4_Go, d, a, b, c,     G_w7c01, MD4S11);
 177     MD4_STEP0(MD4_Go, c, d, a, b,     G_wbc01, MD4S12);
 178     MD4_STEP0(MD4_Go, b, c, d, a,     G_wfc01, MD4S13);
 179
 180     MD4_STEP (MD4_H1, a, b, c, d, w0, H_w0c02, MD4S20);
 181     MD4_STEP0(MD4_H2, d, a, b, c,     H_w8c02, MD4S21);
 182     MD4_STEP0(MD4_H1, c, d, a, b,     H_w4c02, MD4S22);
 183     MD4_STEP0(MD4_H2, b, c, d, a,     H_wcc02, MD4S23);
 184     MD4_STEP0(MD4_H1, a, b, c, d,     H_w2c02, MD4S20);
 185     MD4_STEP0(MD4_H2, d, a, b, c,     H_wac02, MD4S21);
 186     MD4_STEP0(MD4_H1, c, d, a, b,     H_w6c02, MD4S22);
 187     MD4_STEP0(MD4_H2, b, c, d, a,     H_wec02, MD4S23);
 188     MD4_STEP0(MD4_H1, a, b, c, d,     H_w1c02, MD4S20);
 189     MD4_STEP0(MD4_H2, d, a, b, c,     H_w9c02, MD4S21);
 190     MD4_STEP0(MD4_H1, c, d, a, b,     H_w5c02, MD4S22);
 191     MD4_STEP0(MD4_H2, b, c, d, a,     H_wdc02, MD4S23);
 192     MD4_STEP0(MD4_H1, a, b, c, d,     H_w3c02, MD4S20);
 193     MD4_STEP0(MD4_H2, d, a, b, c,     H_wbc02, MD4S21);
 194     MD4_STEP0(MD4_H1, c, d, a, b,     H_w7c02, MD4S22);
 195     MD4_STEP0(MD4_H2, b, c, d, a,     H_wfc02, MD4S23);
 196
 197     const u32x r0 = a;
 198     const u32x r1 = d;
 199     const u32x r2 = c;
 200     const u32x r3 = b;
 201
 202     #include VECT_COMPARE_M
 203   }
 204 }
 205
 206 __device__ static void m01000s (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, 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)
 207 {
 208   /**
 209    * modifier
 210    */
 211
 212   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 213   const u32 lid = threadIdx.x;
 214
 215   /**
 216    * base
 217    */
 218
 219   const u32 F_w0c00 =     0 + MD4C00;
 220   const u32 F_w1c00 = w[ 1] + MD4C00;
 221   const u32 F_w2c00 = w[ 2] + MD4C00;
 222   const u32 F_w3c00 = w[ 3] + MD4C00;
 223   const u32 F_w4c00 = w[ 4] + MD4C00;
 224   const u32 F_w5c00 = w[ 5] + MD4C00;
 225   const u32 F_w6c00 = w[ 6] + MD4C00;
 226   const u32 F_w7c00 = w[ 7] + MD4C00;
 227   const u32 F_w8c00 = w[ 8] + MD4C00;
 228   const u32 F_w9c00 = w[ 9] + MD4C00;
 229   const u32 F_wac00 = w[10] + MD4C00;
 230   const u32 F_wbc00 = w[11] + MD4C00;
 231   const u32 F_wcc00 = w[12] + MD4C00;
 232   const u32 F_wdc00 = w[13] + MD4C00;
 233   const u32 F_wec00 = w[14] + MD4C00;
 234   const u32 F_wfc00 = w[15] + MD4C00;
 235
 236   const u32 G_w0c01 =     0 + MD4C01;
 237   const u32 G_w4c01 = w[ 4] + MD4C01;
 238   const u32 G_w8c01 = w[ 8] + MD4C01;
 239   const u32 G_wcc01 = w[12] + MD4C01;
 240   const u32 G_w1c01 = w[ 1] + MD4C01;
 241   const u32 G_w5c01 = w[ 5] + MD4C01;
 242   const u32 G_w9c01 = w[ 9] + MD4C01;
 243   const u32 G_wdc01 = w[13] + MD4C01;
 244   const u32 G_w2c01 = w[ 2] + MD4C01;
 245   const u32 G_w6c01 = w[ 6] + MD4C01;
 246   const u32 G_wac01 = w[10] + MD4C01;
 247   const u32 G_wec01 = w[14] + MD4C01;
 248   const u32 G_w3c01 = w[ 3] + MD4C01;
 249   const u32 G_w7c01 = w[ 7] + MD4C01;
 250   const u32 G_wbc01 = w[11] + MD4C01;
 251   const u32 G_wfc01 = w[15] + MD4C01;
 252
 253   const u32 H_w0c02 =     0 + MD4C02;
 254   const u32 H_w8c02 = w[ 8] + MD4C02;
 255   const u32 H_w4c02 = w[ 4] + MD4C02;
 256   const u32 H_wcc02 = w[12] + MD4C02;
 257   const u32 H_w2c02 = w[ 2] + MD4C02;
 258   const u32 H_wac02 = w[10] + MD4C02;
 259   const u32 H_w6c02 = w[ 6] + MD4C02;
 260   const u32 H_wec02 = w[14] + MD4C02;
 261   const u32 H_w1c02 = w[ 1] + MD4C02;
 262   const u32 H_w9c02 = w[ 9] + MD4C02;
 263   const u32 H_w5c02 = w[ 5] + MD4C02;
 264   const u32 H_wdc02 = w[13] + MD4C02;
 265   const u32 H_w3c02 = w[ 3] + MD4C02;
 266   const u32 H_wbc02 = w[11] + MD4C02;
 267   const u32 H_w7c02 = w[ 7] + MD4C02;
 268   const u32 H_wfc02 = w[15] + MD4C02;
 269
 270   /**
 271    * digest
 272    */
 273
 274   const u32 search[4] =
 275   {
 276     digests_buf[digests_offset].digest_buf[DGST_R0],
 277     digests_buf[digests_offset].digest_buf[DGST_R1],
 278     digests_buf[digests_offset].digest_buf[DGST_R2],
 279     digests_buf[digests_offset].digest_buf[DGST_R3]
 280   };
 281
 282   /**
 283    * reverse
 284    */
 285
 286   u32 a_rev = digests_buf[digests_offset].digest_buf[0];
 287   u32 b_rev = digests_buf[digests_offset].digest_buf[1];
 288   u32 c_rev = digests_buf[digests_offset].digest_buf[2];
 289   u32 d_rev = digests_buf[digests_offset].digest_buf[3];
 290
 291   MD4_STEP_REV (MD4_H, b_rev, c_rev, d_rev, a_rev, w[15], MD4C02, MD4S23);
 292   MD4_STEP_REV (MD4_H, c_rev, d_rev, a_rev, b_rev, w[ 7], MD4C02, MD4S22);
 293   MD4_STEP_REV (MD4_H, d_rev, a_rev, b_rev, c_rev, w[11], MD4C02, MD4S21);
 294   MD4_STEP_REV (MD4_H, a_rev, b_rev, c_rev, d_rev, w[ 3], MD4C02, MD4S20);
 295   MD4_STEP_REV (MD4_H, b_rev, c_rev, d_rev, a_rev, w[13], MD4C02, MD4S23);
 296   MD4_STEP_REV (MD4_H, c_rev, d_rev, a_rev, b_rev, w[ 5], MD4C02, MD4S22);
 297   MD4_STEP_REV (MD4_H, d_rev, a_rev, b_rev, c_rev, w[ 9], MD4C02, MD4S21);
 298   MD4_STEP_REV (MD4_H, a_rev, b_rev, c_rev, d_rev, w[ 1], MD4C02, MD4S20);
 299   MD4_STEP_REV (MD4_H, b_rev, c_rev, d_rev, a_rev, w[14], MD4C02, MD4S23);
 300   MD4_STEP_REV (MD4_H, c_rev, d_rev, a_rev, b_rev, w[ 6], MD4C02, MD4S22);
 301   MD4_STEP_REV (MD4_H, d_rev, a_rev, b_rev, c_rev, w[10], MD4C02, MD4S21);
 302   MD4_STEP_REV (MD4_H, a_rev, b_rev, c_rev, d_rev, w[ 2], MD4C02, MD4S20);
 303   MD4_STEP_REV (MD4_H, b_rev, c_rev, d_rev, a_rev, w[12], MD4C02, MD4S23);
 304   MD4_STEP_REV (MD4_H, c_rev, d_rev, a_rev, b_rev, w[ 4], MD4C02, MD4S22);
 305   MD4_STEP_REV (MD4_H, d_rev, a_rev, b_rev, c_rev, w[ 8], MD4C02, MD4S21);
 306   MD4_STEP_REV (MD4_H, a_rev, b_rev, c_rev, d_rev,     0, MD4C02, MD4S20);
 307
 308   const u32x sav_c = c_rev;
 309   const u32x sav_d = d_rev;
 310
 311   MD4_STEP_REV1(MD4_G, b_rev, c_rev, d_rev, a_rev, w[15], MD4C01, MD4S13);
 312   MD4_STEP_REV1(MD4_G, c_rev, d_rev, a_rev, b_rev, w[11], MD4C01, MD4S12);
 313
 314   /**
 315    * loop
 316    */
 317
 318   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 319
 320   u32x w0l = w[0];
 321
 322   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 323   {
 324     const u32x w0r = c_bfs[il_pos];
 325
 326     const u32x w0 = w0l | w0r;
 327
 328     u32x pre_a = a_rev;
 329     u32x pre_b = b_rev;
 330     u32x pre_c = c_rev;
 331
 332     pre_a = pre_a - w0;
 333     pre_b = pre_b - MD4_G (sav_c, sav_d, pre_a);
 334     pre_c = pre_c - MD4_G (sav_d, pre_a, pre_b);
 335
 336     u32x tmp2;
 337
 338     u32x a = MD4M_A;
 339     u32x b = MD4M_B;
 340     u32x c = MD4M_C;
 341     u32x d = MD4M_D;
 342
 343     MD4_STEP (MD4_Fo, a, b, c, d, w0, F_w0c00, MD4S00);
 344     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w1c00, MD4S01);
 345     MD4_STEP0(MD4_Fo, c, d, a, b,     F_w2c00, MD4S02);
 346     MD4_STEP0(MD4_Fo, b, c, d, a,     F_w3c00, MD4S03);
 347     MD4_STEP0(MD4_Fo, a, b, c, d,     F_w4c00, MD4S00);
 348     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w5c00, MD4S01);
 349     MD4_STEP0(MD4_Fo, c, d, a, b,     F_w6c00, MD4S02);
 350     MD4_STEP0(MD4_Fo, b, c, d, a,     F_w7c00, MD4S03);
 351     MD4_STEP0(MD4_Fo, a, b, c, d,     F_w8c00, MD4S00);
 352     MD4_STEP0(MD4_Fo, d, a, b, c,     F_w9c00, MD4S01);
 353     MD4_STEP0(MD4_Fo, c, d, a, b,     F_wac00, MD4S02);
 354     MD4_STEP0(MD4_Fo, b, c, d, a,     F_wbc00, MD4S03);
 355     MD4_STEP0(MD4_Fo, a, b, c, d,     F_wcc00, MD4S00);
 356     MD4_STEP0(MD4_Fo, d, a, b, c,     F_wdc00, MD4S01);
 357     MD4_STEP0(MD4_Fo, c, d, a, b,     F_wec00, MD4S02);
 358     MD4_STEP0(MD4_Fo, b, c, d, a,     F_wfc00, MD4S03);
 359
 360     MD4_STEP (MD4_Go, a, b, c, d, w0, G_w0c01, MD4S10);
 361     MD4_STEP0(MD4_Go, d, a, b, c,     G_w4c01, MD4S11);
 362     MD4_STEP0(MD4_Go, c, d, a, b,     G_w8c01, MD4S12);
 363     MD4_STEP0(MD4_Go, b, c, d, a,     G_wcc01, MD4S13);
 364     MD4_STEP0(MD4_Go, a, b, c, d,     G_w1c01, MD4S10);
 365     MD4_STEP0(MD4_Go, d, a, b, c,     G_w5c01, MD4S11);
 366     MD4_STEP0(MD4_Go, c, d, a, b,     G_w9c01, MD4S12);
 367     MD4_STEP0(MD4_Go, b, c, d, a,     G_wdc01, MD4S13);
 368     MD4_STEP0(MD4_Go, a, b, c, d,     G_w2c01, MD4S10);
 369     MD4_STEP0(MD4_Go, d, a, b, c,     G_w6c01, MD4S11);
 370     MD4_STEP0(MD4_Go, c, d, a, b,     G_wac01, MD4S12);
 371
 372     bool q_cond = (pre_c != c);
 373
 374     if (q_cond) continue;
 375
 376     MD4_STEP0(MD4_Go, b, c, d, a,     G_wec01, MD4S13);
 377     MD4_STEP0(MD4_Go, a, b, c, d,     G_w3c01, MD4S10);
 378     MD4_STEP0(MD4_Go, d, a, b, c,     G_w7c01, MD4S11);
 379     MD4_STEP0(MD4_Go, c, d, a, b,     G_wbc01, MD4S12);
 380     MD4_STEP0(MD4_Go, b, c, d, a,     G_wfc01, MD4S13);
 381
 382     MD4_STEP (MD4_H1, a, b, c, d, w0, H_w0c02, MD4S20);
 383     MD4_STEP0(MD4_H2, d, a, b, c,     H_w8c02, MD4S21);
 384     MD4_STEP0(MD4_H1, c, d, a, b,     H_w4c02, MD4S22);
 385     MD4_STEP0(MD4_H2, b, c, d, a,     H_wcc02, MD4S23);
 386     MD4_STEP0(MD4_H1, a, b, c, d,     H_w2c02, MD4S20);
 387     MD4_STEP0(MD4_H2, d, a, b, c,     H_wac02, MD4S21);
 388     MD4_STEP0(MD4_H1, c, d, a, b,     H_w6c02, MD4S22);
 389     MD4_STEP0(MD4_H2, b, c, d, a,     H_wec02, MD4S23);
 390     MD4_STEP0(MD4_H1, a, b, c, d,     H_w1c02, MD4S20);
 391     MD4_STEP0(MD4_H2, d, a, b, c,     H_w9c02, MD4S21);
 392     MD4_STEP0(MD4_H1, c, d, a, b,     H_w5c02, MD4S22);
 393     MD4_STEP0(MD4_H2, b, c, d, a,     H_wdc02, MD4S23);
 394     MD4_STEP0(MD4_H1, a, b, c, d,     H_w3c02, MD4S20);
 395     MD4_STEP0(MD4_H2, d, a, b, c,     H_wbc02, MD4S21);
 396     MD4_STEP0(MD4_H1, c, d, a, b,     H_w7c02, MD4S22);
 397     MD4_STEP0(MD4_H2, b, c, d, a,     H_wfc02, MD4S23);
 398
 399     const u32x r0 = a;
 400     const u32x r1 = d;
 401     const u32x r2 = c;
 402     const u32x r3 = b;
 403
 404     #include VECT_COMPARE_S
 405   }
 406 }
 407
 408 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 409 {
 410   /**
 411    * base
 412    */
 413
 414   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 415
 416   if (gid >= gid_max) return;
 417
 418   u32 w[16];
 419
 420   w[ 0] = pws[gid].i[ 0];
 421   w[ 1] = pws[gid].i[ 1];
 422   w[ 2] = pws[gid].i[ 2];
 423   w[ 3] = pws[gid].i[ 3];
 424   w[ 4] = 0;
 425   w[ 5] = 0;
 426   w[ 6] = 0;
 427   w[ 7] = 0;
 428   w[ 8] = 0;
 429   w[ 9] = 0;
 430   w[10] = 0;
 431   w[11] = 0;
 432   w[12] = 0;
 433   w[13] = 0;
 434   w[14] = pws[gid].i[14];
 435   w[15] = 0;
 436
 437   const u32 pw_len = pws[gid].pw_len;
 438
 439   /**
 440    * main
 441    */
 442
 443   m01000m (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);
 444 }
 445
 446 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 447 {
 448   /**
 449    * base
 450    */
 451
 452   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 453
 454   if (gid >= gid_max) return;
 455
 456   u32 w[16];
 457
 458   w[ 0] = pws[gid].i[ 0];
 459   w[ 1] = pws[gid].i[ 1];
 460   w[ 2] = pws[gid].i[ 2];
 461   w[ 3] = pws[gid].i[ 3];
 462   w[ 4] = pws[gid].i[ 4];
 463   w[ 5] = pws[gid].i[ 5];
 464   w[ 6] = pws[gid].i[ 6];
 465   w[ 7] = pws[gid].i[ 7];
 466   w[ 8] = 0;
 467   w[ 9] = 0;
 468   w[10] = 0;
 469   w[11] = 0;
 470   w[12] = 0;
 471   w[13] = 0;
 472   w[14] = pws[gid].i[14];
 473   w[15] = 0;
 474
 475   const u32 pw_len = pws[gid].pw_len;
 476
 477   /**
 478    * main
 479    */
 480
 481   m01000m (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);
 482 }
 483
 484 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 485 {
 486   /**
 487    * base
 488    */
 489
 490   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 491
 492   if (gid >= gid_max) return;
 493
 494   u32 w[16];
 495
 496   w[ 0] = pws[gid].i[ 0];
 497   w[ 1] = pws[gid].i[ 1];
 498   w[ 2] = pws[gid].i[ 2];
 499   w[ 3] = pws[gid].i[ 3];
 500   w[ 4] = pws[gid].i[ 4];
 501   w[ 5] = pws[gid].i[ 5];
 502   w[ 6] = pws[gid].i[ 6];
 503   w[ 7] = pws[gid].i[ 7];
 504   w[ 8] = pws[gid].i[ 8];
 505   w[ 9] = pws[gid].i[ 9];
 506   w[10] = pws[gid].i[10];
 507   w[11] = pws[gid].i[11];
 508   w[12] = pws[gid].i[12];
 509   w[13] = pws[gid].i[13];
 510   w[14] = pws[gid].i[14];
 511   w[15] = pws[gid].i[15];
 512
 513   const u32 pw_len = pws[gid].pw_len;
 514
 515   /**
 516    * main
 517    */
 518
 519   m01000m (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);
 520 }
 521
 522 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 523 {
 524   /**
 525    * base
 526    */
 527
 528   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 529
 530   if (gid >= gid_max) return;
 531
 532   u32 w[16];
 533
 534   w[ 0] = pws[gid].i[ 0];
 535   w[ 1] = pws[gid].i[ 1];
 536   w[ 2] = pws[gid].i[ 2];
 537   w[ 3] = pws[gid].i[ 3];
 538   w[ 4] = 0;
 539   w[ 5] = 0;
 540   w[ 6] = 0;
 541   w[ 7] = 0;
 542   w[ 8] = 0;
 543   w[ 9] = 0;
 544   w[10] = 0;
 545   w[11] = 0;
 546   w[12] = 0;
 547   w[13] = 0;
 548   w[14] = pws[gid].i[14];
 549   w[15] = 0;
 550
 551   const u32 pw_len = pws[gid].pw_len;
 552
 553   /**
 554    * main
 555    */
 556
 557   m01000s (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);
 558 }
 559
 560 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 561 {
 562   /**
 563    * base
 564    */
 565
 566   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 567
 568   if (gid >= gid_max) return;
 569
 570   u32 w[16];
 571
 572   w[ 0] = pws[gid].i[ 0];
 573   w[ 1] = pws[gid].i[ 1];
 574   w[ 2] = pws[gid].i[ 2];
 575   w[ 3] = pws[gid].i[ 3];
 576   w[ 4] = pws[gid].i[ 4];
 577   w[ 5] = pws[gid].i[ 5];
 578   w[ 6] = pws[gid].i[ 6];
 579   w[ 7] = pws[gid].i[ 7];
 580   w[ 8] = 0;
 581   w[ 9] = 0;
 582   w[10] = 0;
 583   w[11] = 0;
 584   w[12] = 0;
 585   w[13] = 0;
 586   w[14] = pws[gid].i[14];
 587   w[15] = 0;
 588
 589   const u32 pw_len = pws[gid].pw_len;
 590
 591   /**
 592    * main
 593    */
 594
 595   m01000s (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);
 596 }
 597
 598 extern "C" __global__ void __launch_bounds__ (256, 1) m01000_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r,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)
 599 {
 600   /**
 601    * base
 602    */
 603
 604   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 605
 606   if (gid >= gid_max) return;
 607
 608   u32 w[16];
 609
 610   w[ 0] = pws[gid].i[ 0];
 611   w[ 1] = pws[gid].i[ 1];
 612   w[ 2] = pws[gid].i[ 2];
 613   w[ 3] = pws[gid].i[ 3];
 614   w[ 4] = pws[gid].i[ 4];
 615   w[ 5] = pws[gid].i[ 5];
 616   w[ 6] = pws[gid].i[ 6];
 617   w[ 7] = pws[gid].i[ 7];
 618   w[ 8] = pws[gid].i[ 8];
 619   w[ 9] = pws[gid].i[ 9];
 620   w[10] = pws[gid].i[10];
 621   w[11] = pws[gid].i[11];
 622   w[12] = pws[gid].i[12];
 623   w[13] = pws[gid].i[13];
 624   w[14] = pws[gid].i[14];
 625   w[15] = pws[gid].i[15];
 626
 627   const u32 pw_len = pws[gid].pw_len;
 628
 629   /**
 630    * main
 631    */
 632
 633   m01000s (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);
 634 }