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