nv/m09900_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 m09900m (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)
  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    * base
  57    */
  58
  59   const u32 F_w0c00 =    0  + MD5C00;
  60   const u32 F_w1c01 = w[ 1] + MD5C01;
  61   const u32 F_w2c02 = w[ 2] + MD5C02;
  62   const u32 F_w3c03 = w[ 3] + MD5C03;
  63   const u32 F_w4c04 = w[ 4] + MD5C04;
  64   const u32 F_w5c05 = w[ 5] + MD5C05;
  65   const u32 F_w6c06 = w[ 6] + MD5C06;
  66   const u32 F_w7c07 = w[ 7] + MD5C07;
  67   const u32 F_w8c08 = w[ 8] + MD5C08;
  68   const u32 F_w9c09 = w[ 9] + MD5C09;
  69   const u32 F_wac0a = w[10] + MD5C0a;
  70   const u32 F_wbc0b = w[11] + MD5C0b;
  71   const u32 F_wcc0c = w[12] + MD5C0c;
  72   const u32 F_wdc0d = w[13] + MD5C0d;
  73   const u32 F_wec0e = w[14] + MD5C0e;
  74   const u32 F_wfc0f = w[15] + MD5C0f;
  75
  76   const u32 G_w1c10 = w[ 1] + MD5C10;
  77   const u32 G_w6c11 = w[ 6] + MD5C11;
  78   const u32 G_wbc12 = w[11] + MD5C12;
  79   const u32 G_w0c13 =    0  + MD5C13;
  80   const u32 G_w5c14 = w[ 5] + MD5C14;
  81   const u32 G_wac15 = w[10] + MD5C15;
  82   const u32 G_wfc16 = w[15] + MD5C16;
  83   const u32 G_w4c17 = w[ 4] + MD5C17;
  84   const u32 G_w9c18 = w[ 9] + MD5C18;
  85   const u32 G_wec19 = w[14] + MD5C19;
  86   const u32 G_w3c1a = w[ 3] + MD5C1a;
  87   const u32 G_w8c1b = w[ 8] + MD5C1b;
  88   const u32 G_wdc1c = w[13] + MD5C1c;
  89   const u32 G_w2c1d = w[ 2] + MD5C1d;
  90   const u32 G_w7c1e = w[ 7] + MD5C1e;
  91   const u32 G_wcc1f = w[12] + MD5C1f;
  92
  93   const u32 H_w5c20 = w[ 5] + MD5C20;
  94   const u32 H_w8c21 = w[ 8] + MD5C21;
  95   const u32 H_wbc22 = w[11] + MD5C22;
  96   const u32 H_wec23 = w[14] + MD5C23;
  97   const u32 H_w1c24 = w[ 1] + MD5C24;
  98   const u32 H_w4c25 = w[ 4] + MD5C25;
  99   const u32 H_w7c26 = w[ 7] + MD5C26;
 100   const u32 H_wac27 = w[10] + MD5C27;
 101   const u32 H_wdc28 = w[13] + MD5C28;
 102   const u32 H_w0c29 =    0  + MD5C29;
 103   const u32 H_w3c2a = w[ 3] + MD5C2a;
 104   const u32 H_w6c2b = w[ 6] + MD5C2b;
 105   const u32 H_w9c2c = w[ 9] + MD5C2c;
 106   const u32 H_wcc2d = w[12] + MD5C2d;
 107   const u32 H_wfc2e = w[15] + MD5C2e;
 108   const u32 H_w2c2f = w[ 2] + MD5C2f;
 109
 110   const u32 I_w0c30 =    0  + MD5C30;
 111   const u32 I_w7c31 = w[ 7] + MD5C31;
 112   const u32 I_wec32 = w[14] + MD5C32;
 113   const u32 I_w5c33 = w[ 5] + MD5C33;
 114   const u32 I_wcc34 = w[12] + MD5C34;
 115   const u32 I_w3c35 = w[ 3] + MD5C35;
 116   const u32 I_wac36 = w[10] + MD5C36;
 117   const u32 I_w1c37 = w[ 1] + MD5C37;
 118   const u32 I_w8c38 = w[ 8] + MD5C38;
 119   const u32 I_wfc39 = w[15] + MD5C39;
 120   const u32 I_w6c3a = w[ 6] + MD5C3a;
 121   const u32 I_wdc3b = w[13] + MD5C3b;
 122   const u32 I_w4c3c = w[ 4] + MD5C3c;
 123   const u32 I_wbc3d = w[11] + MD5C3d;
 124   const u32 I_w2c3e = w[ 2] + MD5C3e;
 125   const u32 I_w9c3f = w[ 9] + MD5C3f;
 126
 127   /**
 128    * loop
 129    */
 130
 131   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 132
 133   u32x w0l = w[0];
 134
 135   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 136   {
 137     const u32x w0r = c_bfs[il_pos];
 138
 139     const u32x w0 = w0l | w0r;
 140
 141     u32x a = MD5M_A;
 142     u32x b = MD5M_B;
 143     u32x c = MD5M_C;
 144     u32x d = MD5M_D;
 145
 146     MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00);
 147     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w1c01, MD5S01);
 148     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w2c02, MD5S02);
 149     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w3c03, MD5S03);
 150     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w4c04, MD5S00);
 151     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w5c05, MD5S01);
 152     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w6c06, MD5S02);
 153     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w7c07, MD5S03);
 154     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w8c08, MD5S00);
 155     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w9c09, MD5S01);
 156     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wac0a, MD5S02);
 157     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wbc0b, MD5S03);
 158     MD5_STEP0(MD5_Fo, a, b, c, d,     F_wcc0c, MD5S00);
 159     MD5_STEP0(MD5_Fo, d, a, b, c,     F_wdc0d, MD5S01);
 160     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wec0e, MD5S02);
 161     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wfc0f, MD5S03);
 162
 163     MD5_STEP0(MD5_Go, a, b, c, d,     G_w1c10, MD5S10);
 164     MD5_STEP0(MD5_Go, d, a, b, c,     G_w6c11, MD5S11);
 165     MD5_STEP0(MD5_Go, c, d, a, b,     G_wbc12, MD5S12);
 166     MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13);
 167     MD5_STEP0(MD5_Go, a, b, c, d,     G_w5c14, MD5S10);
 168     MD5_STEP0(MD5_Go, d, a, b, c,     G_wac15, MD5S11);
 169     MD5_STEP0(MD5_Go, c, d, a, b,     G_wfc16, MD5S12);
 170     MD5_STEP0(MD5_Go, b, c, d, a,     G_w4c17, MD5S13);
 171     MD5_STEP0(MD5_Go, a, b, c, d,     G_w9c18, MD5S10);
 172     MD5_STEP0(MD5_Go, d, a, b, c,     G_wec19, MD5S11);
 173     MD5_STEP0(MD5_Go, c, d, a, b,     G_w3c1a, MD5S12);
 174     MD5_STEP0(MD5_Go, b, c, d, a,     G_w8c1b, MD5S13);
 175     MD5_STEP0(MD5_Go, a, b, c, d,     G_wdc1c, MD5S10);
 176     MD5_STEP0(MD5_Go, d, a, b, c,     G_w2c1d, MD5S11);
 177     MD5_STEP0(MD5_Go, c, d, a, b,     G_w7c1e, MD5S12);
 178     MD5_STEP0(MD5_Go, b, c, d, a,     G_wcc1f, MD5S13);
 179
 180     MD5_STEP0(MD5_H , a, b, c, d,     H_w5c20, MD5S20);
 181     MD5_STEP0(MD5_H , d, a, b, c,     H_w8c21, MD5S21);
 182     MD5_STEP0(MD5_H , c, d, a, b,     H_wbc22, MD5S22);
 183     MD5_STEP0(MD5_H , b, c, d, a,     H_wec23, MD5S23);
 184     MD5_STEP0(MD5_H , a, b, c, d,     H_w1c24, MD5S20);
 185     MD5_STEP0(MD5_H , d, a, b, c,     H_w4c25, MD5S21);
 186     MD5_STEP0(MD5_H , c, d, a, b,     H_w7c26, MD5S22);
 187     MD5_STEP0(MD5_H , b, c, d, a,     H_wac27, MD5S23);
 188     MD5_STEP0(MD5_H , a, b, c, d,     H_wdc28, MD5S20);
 189     MD5_STEP (MD5_H , d, a, b, c, w0, H_w0c29, MD5S21);
 190     MD5_STEP0(MD5_H , c, d, a, b,     H_w3c2a, MD5S22);
 191     MD5_STEP0(MD5_H , b, c, d, a,     H_w6c2b, MD5S23);
 192     MD5_STEP0(MD5_H , a, b, c, d,     H_w9c2c, MD5S20);
 193     MD5_STEP0(MD5_H , d, a, b, c,     H_wcc2d, MD5S21);
 194     MD5_STEP0(MD5_H , c, d, a, b,     H_wfc2e, MD5S22);
 195     MD5_STEP0(MD5_H , b, c, d, a,     H_w2c2f, MD5S23);
 196
 197     MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30);
 198     MD5_STEP0(MD5_I , d, a, b, c,     I_w7c31, MD5S31);
 199     MD5_STEP0(MD5_I , c, d, a, b,     I_wec32, MD5S32);
 200     MD5_STEP0(MD5_I , b, c, d, a,     I_w5c33, MD5S33);
 201     MD5_STEP0(MD5_I , a, b, c, d,     I_wcc34, MD5S30);
 202     MD5_STEP0(MD5_I , d, a, b, c,     I_w3c35, MD5S31);
 203     MD5_STEP0(MD5_I , c, d, a, b,     I_wac36, MD5S32);
 204     MD5_STEP0(MD5_I , b, c, d, a,     I_w1c37, MD5S33);
 205     MD5_STEP0(MD5_I , a, b, c, d,     I_w8c38, MD5S30);
 206     MD5_STEP0(MD5_I , d, a, b, c,     I_wfc39, MD5S31);
 207     MD5_STEP0(MD5_I , c, d, a, b,     I_w6c3a, MD5S32);
 208     MD5_STEP0(MD5_I , b, c, d, a,     I_wdc3b, MD5S33);
 209     MD5_STEP0(MD5_I , a, b, c, d,     I_w4c3c, MD5S30);
 210     MD5_STEP0(MD5_I , d, a, b, c,     I_wbc3d, MD5S31);
 211     MD5_STEP0(MD5_I , c, d, a, b,     I_w2c3e, MD5S32);
 212     MD5_STEP0(MD5_I , b, c, d, a,     I_w9c3f, MD5S33);
 213
 214     a += MD5M_A;
 215     b += MD5M_B;
 216     c += MD5M_C;
 217     d += MD5M_D;
 218
 219     u32x r_a = a;
 220     u32x r_b = b;
 221     u32x r_c = c;
 222     u32x r_d = d;
 223
 224     u32x t0[4];
 225     u32x t1[4];
 226     u32x t2[4];
 227     u32x t3[4];
 228
 229     t0[0] = 0;
 230     t0[1] = 0;
 231     t0[2] = 0;
 232     t0[3] = 0;
 233     t1[0] = 0;
 234     t1[1] = 0;
 235     t1[2] = 0;
 236     t1[3] = 0;
 237     t2[0] = 0;
 238     t2[1] = 0x80;
 239     t2[2] = 0;
 240     t2[3] = 0;
 241     t3[0] = 0;
 242     t3[1] = 0;
 243     t3[2] = 100 * 8;
 244     t3[3] = 0;
 245
 246     MD5_STEP (MD5_Fo, a, b, c, d, t0[0], MD5C00, MD5S00);
 247     MD5_STEP (MD5_Fo, d, a, b, c, t0[1], MD5C01, MD5S01);
 248     MD5_STEP (MD5_Fo, c, d, a, b, t0[2], MD5C02, MD5S02);
 249     MD5_STEP (MD5_Fo, b, c, d, a, t0[3], MD5C03, MD5S03);
 250     MD5_STEP (MD5_Fo, a, b, c, d, t1[0], MD5C04, MD5S00);
 251     MD5_STEP (MD5_Fo, d, a, b, c, t1[1], MD5C05, MD5S01);
 252     MD5_STEP (MD5_Fo, c, d, a, b, t1[2], MD5C06, MD5S02);
 253     MD5_STEP (MD5_Fo, b, c, d, a, t1[3], MD5C07, MD5S03);
 254     MD5_STEP (MD5_Fo, a, b, c, d, t2[0], MD5C08, MD5S00);
 255     MD5_STEP (MD5_Fo, d, a, b, c, t2[1], MD5C09, MD5S01);
 256     MD5_STEP (MD5_Fo, c, d, a, b, t2[2], MD5C0a, MD5S02);
 257     MD5_STEP (MD5_Fo, b, c, d, a, t2[3], MD5C0b, MD5S03);
 258     MD5_STEP (MD5_Fo, a, b, c, d, t3[0], MD5C0c, MD5S00);
 259     MD5_STEP (MD5_Fo, d, a, b, c, t3[1], MD5C0d, MD5S01);
 260     MD5_STEP (MD5_Fo, c, d, a, b, t3[2], MD5C0e, MD5S02);
 261     MD5_STEP (MD5_Fo, b, c, d, a, t3[3], MD5C0f, MD5S03);
 262
 263     MD5_STEP (MD5_Go, a, b, c, d, t0[1], MD5C10, MD5S10);
 264     MD5_STEP (MD5_Go, d, a, b, c, t1[2], MD5C11, MD5S11);
 265     MD5_STEP (MD5_Go, c, d, a, b, t2[3], MD5C12, MD5S12);
 266     MD5_STEP (MD5_Go, b, c, d, a, t0[0], MD5C13, MD5S13);
 267     MD5_STEP (MD5_Go, a, b, c, d, t1[1], MD5C14, MD5S10);
 268     MD5_STEP (MD5_Go, d, a, b, c, t2[2], MD5C15, MD5S11);
 269     MD5_STEP (MD5_Go, c, d, a, b, t3[3], MD5C16, MD5S12);
 270     MD5_STEP (MD5_Go, b, c, d, a, t1[0], MD5C17, MD5S13);
 271     MD5_STEP (MD5_Go, a, b, c, d, t2[1], MD5C18, MD5S10);
 272     MD5_STEP (MD5_Go, d, a, b, c, t3[2], MD5C19, MD5S11);
 273     MD5_STEP (MD5_Go, c, d, a, b, t0[3], MD5C1a, MD5S12);
 274     MD5_STEP (MD5_Go, b, c, d, a, t2[0], MD5C1b, MD5S13);
 275     MD5_STEP (MD5_Go, a, b, c, d, t3[1], MD5C1c, MD5S10);
 276     MD5_STEP (MD5_Go, d, a, b, c, t0[2], MD5C1d, MD5S11);
 277     MD5_STEP (MD5_Go, c, d, a, b, t1[3], MD5C1e, MD5S12);
 278     MD5_STEP (MD5_Go, b, c, d, a, t3[0], MD5C1f, MD5S13);
 279
 280     MD5_STEP (MD5_H , a, b, c, d, t1[1], MD5C20, MD5S20);
 281     MD5_STEP (MD5_H , d, a, b, c, t2[0], MD5C21, MD5S21);
 282     MD5_STEP (MD5_H , c, d, a, b, t2[3], MD5C22, MD5S22);
 283     MD5_STEP (MD5_H , b, c, d, a, t3[2], MD5C23, MD5S23);
 284     MD5_STEP (MD5_H , a, b, c, d, t0[1], MD5C24, MD5S20);
 285     MD5_STEP (MD5_H , d, a, b, c, t1[0], MD5C25, MD5S21);
 286     MD5_STEP (MD5_H , c, d, a, b, t1[3], MD5C26, MD5S22);
 287     MD5_STEP (MD5_H , b, c, d, a, t2[2], MD5C27, MD5S23);
 288     MD5_STEP (MD5_H , a, b, c, d, t3[1], MD5C28, MD5S20);
 289     MD5_STEP (MD5_H , d, a, b, c, t0[0], MD5C29, MD5S21);
 290     MD5_STEP (MD5_H , c, d, a, b, t0[3], MD5C2a, MD5S22);
 291     MD5_STEP (MD5_H , b, c, d, a, t1[2], MD5C2b, MD5S23);
 292     MD5_STEP (MD5_H , a, b, c, d, t2[1], MD5C2c, MD5S20);
 293     MD5_STEP (MD5_H , d, a, b, c, t3[0], MD5C2d, MD5S21);
 294     MD5_STEP (MD5_H , c, d, a, b, t3[3], MD5C2e, MD5S22);
 295     MD5_STEP (MD5_H , b, c, d, a, t0[2], MD5C2f, MD5S23);
 296
 297     MD5_STEP (MD5_I , a, b, c, d, t0[0], MD5C30, MD5S30);
 298     MD5_STEP (MD5_I , d, a, b, c, t1[3], MD5C31, MD5S31);
 299     MD5_STEP (MD5_I , c, d, a, b, t3[2], MD5C32, MD5S32);
 300     MD5_STEP (MD5_I , b, c, d, a, t1[1], MD5C33, MD5S33);
 301     MD5_STEP (MD5_I , a, b, c, d, t3[0], MD5C34, MD5S30);
 302     MD5_STEP (MD5_I , d, a, b, c, t0[3], MD5C35, MD5S31);
 303     MD5_STEP (MD5_I , c, d, a, b, t2[2], MD5C36, MD5S32);
 304     MD5_STEP (MD5_I , b, c, d, a, t0[1], MD5C37, MD5S33);
 305     MD5_STEP (MD5_I , a, b, c, d, t2[0], MD5C38, MD5S30);
 306     MD5_STEP (MD5_I , d, a, b, c, t3[3], MD5C39, MD5S31);
 307     MD5_STEP (MD5_I , c, d, a, b, t1[2], MD5C3a, MD5S32);
 308     MD5_STEP (MD5_I , b, c, d, a, t3[1], MD5C3b, MD5S33);
 309     MD5_STEP (MD5_I , a, b, c, d, t1[0], MD5C3c, MD5S30);
 310     MD5_STEP (MD5_I , d, a, b, c, t2[3], MD5C3d, MD5S31);
 311     MD5_STEP (MD5_I , c, d, a, b, t0[2], MD5C3e, MD5S32);
 312     MD5_STEP (MD5_I , b, c, d, a, t2[1], MD5C3f, MD5S33);
 313
 314     a += r_a;
 315     b += r_b;
 316     c += r_c;
 317     d += r_d;
 318
 319     const u32x r0 = a;
 320     const u32x r1 = d;
 321     const u32x r2 = c;
 322     const u32x r3 = b;
 323
 324     #include VECT_COMPARE_M
 325   }
 326 }
 327
 328 __device__ static void m09900s (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)
 329 {
 330   /**
 331    * modifier
 332    */
 333
 334   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 335   const u32 lid = threadIdx.x;
 336
 337   /**
 338    * base
 339    */
 340
 341   const u32 F_w0c00 =    0  + MD5C00;
 342   const u32 F_w1c01 = w[ 1] + MD5C01;
 343   const u32 F_w2c02 = w[ 2] + MD5C02;
 344   const u32 F_w3c03 = w[ 3] + MD5C03;
 345   const u32 F_w4c04 = w[ 4] + MD5C04;
 346   const u32 F_w5c05 = w[ 5] + MD5C05;
 347   const u32 F_w6c06 = w[ 6] + MD5C06;
 348   const u32 F_w7c07 = w[ 7] + MD5C07;
 349   const u32 F_w8c08 = w[ 8] + MD5C08;
 350   const u32 F_w9c09 = w[ 9] + MD5C09;
 351   const u32 F_wac0a = w[10] + MD5C0a;
 352   const u32 F_wbc0b = w[11] + MD5C0b;
 353   const u32 F_wcc0c = w[12] + MD5C0c;
 354   const u32 F_wdc0d = w[13] + MD5C0d;
 355   const u32 F_wec0e = w[14] + MD5C0e;
 356   const u32 F_wfc0f = w[15] + MD5C0f;
 357
 358   const u32 G_w1c10 = w[ 1] + MD5C10;
 359   const u32 G_w6c11 = w[ 6] + MD5C11;
 360   const u32 G_wbc12 = w[11] + MD5C12;
 361   const u32 G_w0c13 =    0  + MD5C13;
 362   const u32 G_w5c14 = w[ 5] + MD5C14;
 363   const u32 G_wac15 = w[10] + MD5C15;
 364   const u32 G_wfc16 = w[15] + MD5C16;
 365   const u32 G_w4c17 = w[ 4] + MD5C17;
 366   const u32 G_w9c18 = w[ 9] + MD5C18;
 367   const u32 G_wec19 = w[14] + MD5C19;
 368   const u32 G_w3c1a = w[ 3] + MD5C1a;
 369   const u32 G_w8c1b = w[ 8] + MD5C1b;
 370   const u32 G_wdc1c = w[13] + MD5C1c;
 371   const u32 G_w2c1d = w[ 2] + MD5C1d;
 372   const u32 G_w7c1e = w[ 7] + MD5C1e;
 373   const u32 G_wcc1f = w[12] + MD5C1f;
 374
 375   const u32 H_w5c20 = w[ 5] + MD5C20;
 376   const u32 H_w8c21 = w[ 8] + MD5C21;
 377   const u32 H_wbc22 = w[11] + MD5C22;
 378   const u32 H_wec23 = w[14] + MD5C23;
 379   const u32 H_w1c24 = w[ 1] + MD5C24;
 380   const u32 H_w4c25 = w[ 4] + MD5C25;
 381   const u32 H_w7c26 = w[ 7] + MD5C26;
 382   const u32 H_wac27 = w[10] + MD5C27;
 383   const u32 H_wdc28 = w[13] + MD5C28;
 384   const u32 H_w0c29 =    0  + MD5C29;
 385   const u32 H_w3c2a = w[ 3] + MD5C2a;
 386   const u32 H_w6c2b = w[ 6] + MD5C2b;
 387   const u32 H_w9c2c = w[ 9] + MD5C2c;
 388   const u32 H_wcc2d = w[12] + MD5C2d;
 389   const u32 H_wfc2e = w[15] + MD5C2e;
 390   const u32 H_w2c2f = w[ 2] + MD5C2f;
 391
 392   const u32 I_w0c30 =    0  + MD5C30;
 393   const u32 I_w7c31 = w[ 7] + MD5C31;
 394   const u32 I_wec32 = w[14] + MD5C32;
 395   const u32 I_w5c33 = w[ 5] + MD5C33;
 396   const u32 I_wcc34 = w[12] + MD5C34;
 397   const u32 I_w3c35 = w[ 3] + MD5C35;
 398   const u32 I_wac36 = w[10] + MD5C36;
 399   const u32 I_w1c37 = w[ 1] + MD5C37;
 400   const u32 I_w8c38 = w[ 8] + MD5C38;
 401   const u32 I_wfc39 = w[15] + MD5C39;
 402   const u32 I_w6c3a = w[ 6] + MD5C3a;
 403   const u32 I_wdc3b = w[13] + MD5C3b;
 404   const u32 I_w4c3c = w[ 4] + MD5C3c;
 405   const u32 I_wbc3d = w[11] + MD5C3d;
 406   const u32 I_w2c3e = w[ 2] + MD5C3e;
 407   const u32 I_w9c3f = w[ 9] + MD5C3f;
 408
 409   /**
 410    * digest
 411    */
 412
 413   const u32 search[4] =
 414   {
 415     digests_buf[digests_offset].digest_buf[DGST_R0],
 416     digests_buf[digests_offset].digest_buf[DGST_R1],
 417     digests_buf[digests_offset].digest_buf[DGST_R2],
 418     digests_buf[digests_offset].digest_buf[DGST_R3]
 419   };
 420
 421   /**
 422    * loop
 423    */
 424
 425   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 426
 427   u32x w0l = w[0];
 428
 429   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 430   {
 431     const u32x w0r = c_bfs[il_pos];
 432
 433     const u32x w0 = w0l | w0r;
 434
 435     u32x a = MD5M_A;
 436     u32x b = MD5M_B;
 437     u32x c = MD5M_C;
 438     u32x d = MD5M_D;
 439
 440     MD5_STEP (MD5_Fo, a, b, c, d, w0, F_w0c00, MD5S00);
 441     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w1c01, MD5S01);
 442     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w2c02, MD5S02);
 443     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w3c03, MD5S03);
 444     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w4c04, MD5S00);
 445     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w5c05, MD5S01);
 446     MD5_STEP0(MD5_Fo, c, d, a, b,     F_w6c06, MD5S02);
 447     MD5_STEP0(MD5_Fo, b, c, d, a,     F_w7c07, MD5S03);
 448     MD5_STEP0(MD5_Fo, a, b, c, d,     F_w8c08, MD5S00);
 449     MD5_STEP0(MD5_Fo, d, a, b, c,     F_w9c09, MD5S01);
 450     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wac0a, MD5S02);
 451     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wbc0b, MD5S03);
 452     MD5_STEP0(MD5_Fo, a, b, c, d,     F_wcc0c, MD5S00);
 453     MD5_STEP0(MD5_Fo, d, a, b, c,     F_wdc0d, MD5S01);
 454     MD5_STEP0(MD5_Fo, c, d, a, b,     F_wec0e, MD5S02);
 455     MD5_STEP0(MD5_Fo, b, c, d, a,     F_wfc0f, MD5S03);
 456
 457     MD5_STEP0(MD5_Go, a, b, c, d,     G_w1c10, MD5S10);
 458     MD5_STEP0(MD5_Go, d, a, b, c,     G_w6c11, MD5S11);
 459     MD5_STEP0(MD5_Go, c, d, a, b,     G_wbc12, MD5S12);
 460     MD5_STEP (MD5_Go, b, c, d, a, w0, G_w0c13, MD5S13);
 461     MD5_STEP0(MD5_Go, a, b, c, d,     G_w5c14, MD5S10);
 462     MD5_STEP0(MD5_Go, d, a, b, c,     G_wac15, MD5S11);
 463     MD5_STEP0(MD5_Go, c, d, a, b,     G_wfc16, MD5S12);
 464     MD5_STEP0(MD5_Go, b, c, d, a,     G_w4c17, MD5S13);
 465     MD5_STEP0(MD5_Go, a, b, c, d,     G_w9c18, MD5S10);
 466     MD5_STEP0(MD5_Go, d, a, b, c,     G_wec19, MD5S11);
 467     MD5_STEP0(MD5_Go, c, d, a, b,     G_w3c1a, MD5S12);
 468     MD5_STEP0(MD5_Go, b, c, d, a,     G_w8c1b, MD5S13);
 469     MD5_STEP0(MD5_Go, a, b, c, d,     G_wdc1c, MD5S10);
 470     MD5_STEP0(MD5_Go, d, a, b, c,     G_w2c1d, MD5S11);
 471     MD5_STEP0(MD5_Go, c, d, a, b,     G_w7c1e, MD5S12);
 472     MD5_STEP0(MD5_Go, b, c, d, a,     G_wcc1f, MD5S13);
 473
 474     MD5_STEP0(MD5_H , a, b, c, d,     H_w5c20, MD5S20);
 475     MD5_STEP0(MD5_H , d, a, b, c,     H_w8c21, MD5S21);
 476     MD5_STEP0(MD5_H , c, d, a, b,     H_wbc22, MD5S22);
 477     MD5_STEP0(MD5_H , b, c, d, a,     H_wec23, MD5S23);
 478     MD5_STEP0(MD5_H , a, b, c, d,     H_w1c24, MD5S20);
 479     MD5_STEP0(MD5_H , d, a, b, c,     H_w4c25, MD5S21);
 480     MD5_STEP0(MD5_H , c, d, a, b,     H_w7c26, MD5S22);
 481     MD5_STEP0(MD5_H , b, c, d, a,     H_wac27, MD5S23);
 482     MD5_STEP0(MD5_H , a, b, c, d,     H_wdc28, MD5S20);
 483     MD5_STEP (MD5_H , d, a, b, c, w0, H_w0c29, MD5S21);
 484     MD5_STEP0(MD5_H , c, d, a, b,     H_w3c2a, MD5S22);
 485     MD5_STEP0(MD5_H , b, c, d, a,     H_w6c2b, MD5S23);
 486     MD5_STEP0(MD5_H , a, b, c, d,     H_w9c2c, MD5S20);
 487     MD5_STEP0(MD5_H , d, a, b, c,     H_wcc2d, MD5S21);
 488     MD5_STEP0(MD5_H , c, d, a, b,     H_wfc2e, MD5S22);
 489     MD5_STEP0(MD5_H , b, c, d, a,     H_w2c2f, MD5S23);
 490
 491     MD5_STEP (MD5_I , a, b, c, d, w0, I_w0c30, MD5S30);
 492     MD5_STEP0(MD5_I , d, a, b, c,     I_w7c31, MD5S31);
 493     MD5_STEP0(MD5_I , c, d, a, b,     I_wec32, MD5S32);
 494     MD5_STEP0(MD5_I , b, c, d, a,     I_w5c33, MD5S33);
 495     MD5_STEP0(MD5_I , a, b, c, d,     I_wcc34, MD5S30);
 496     MD5_STEP0(MD5_I , d, a, b, c,     I_w3c35, MD5S31);
 497     MD5_STEP0(MD5_I , c, d, a, b,     I_wac36, MD5S32);
 498     MD5_STEP0(MD5_I , b, c, d, a,     I_w1c37, MD5S33);
 499     MD5_STEP0(MD5_I , a, b, c, d,     I_w8c38, MD5S30);
 500     MD5_STEP0(MD5_I , d, a, b, c,     I_wfc39, MD5S31);
 501     MD5_STEP0(MD5_I , c, d, a, b,     I_w6c3a, MD5S32);
 502     MD5_STEP0(MD5_I , b, c, d, a,     I_wdc3b, MD5S33);
 503     MD5_STEP0(MD5_I , a, b, c, d,     I_w4c3c, MD5S30);
 504     MD5_STEP0(MD5_I , d, a, b, c,     I_wbc3d, MD5S31);
 505     MD5_STEP0(MD5_I , c, d, a, b,     I_w2c3e, MD5S32);
 506     MD5_STEP0(MD5_I , b, c, d, a,     I_w9c3f, MD5S33);
 507
 508     a += MD5M_A;
 509     b += MD5M_B;
 510     c += MD5M_C;
 511     d += MD5M_D;
 512
 513     u32x r_a = a;
 514     u32x r_b = b;
 515     u32x r_c = c;
 516     u32x r_d = d;
 517
 518     u32x t0[4];
 519     u32x t1[4];
 520     u32x t2[4];
 521     u32x t3[4];
 522
 523     t0[0] = 0;
 524     t0[1] = 0;
 525     t0[2] = 0;
 526     t0[3] = 0;
 527     t1[0] = 0;
 528     t1[1] = 0;
 529     t1[2] = 0;
 530     t1[3] = 0;
 531     t2[0] = 0;
 532     t2[1] = 0x80;
 533     t2[2] = 0;
 534     t2[3] = 0;
 535     t3[0] = 0;
 536     t3[1] = 0;
 537     t3[2] = 100 * 8;
 538     t3[3] = 0;
 539
 540     MD5_STEP (MD5_Fo, a, b, c, d, t0[0], MD5C00, MD5S00);
 541     MD5_STEP (MD5_Fo, d, a, b, c, t0[1], MD5C01, MD5S01);
 542     MD5_STEP (MD5_Fo, c, d, a, b, t0[2], MD5C02, MD5S02);
 543     MD5_STEP (MD5_Fo, b, c, d, a, t0[3], MD5C03, MD5S03);
 544     MD5_STEP (MD5_Fo, a, b, c, d, t1[0], MD5C04, MD5S00);
 545     MD5_STEP (MD5_Fo, d, a, b, c, t1[1], MD5C05, MD5S01);
 546     MD5_STEP (MD5_Fo, c, d, a, b, t1[2], MD5C06, MD5S02);
 547     MD5_STEP (MD5_Fo, b, c, d, a, t1[3], MD5C07, MD5S03);
 548     MD5_STEP (MD5_Fo, a, b, c, d, t2[0], MD5C08, MD5S00);
 549     MD5_STEP (MD5_Fo, d, a, b, c, t2[1], MD5C09, MD5S01);
 550     MD5_STEP (MD5_Fo, c, d, a, b, t2[2], MD5C0a, MD5S02);
 551     MD5_STEP (MD5_Fo, b, c, d, a, t2[3], MD5C0b, MD5S03);
 552     MD5_STEP (MD5_Fo, a, b, c, d, t3[0], MD5C0c, MD5S00);
 553     MD5_STEP (MD5_Fo, d, a, b, c, t3[1], MD5C0d, MD5S01);
 554     MD5_STEP (MD5_Fo, c, d, a, b, t3[2], MD5C0e, MD5S02);
 555     MD5_STEP (MD5_Fo, b, c, d, a, t3[3], MD5C0f, MD5S03);
 556
 557     MD5_STEP (MD5_Go, a, b, c, d, t0[1], MD5C10, MD5S10);
 558     MD5_STEP (MD5_Go, d, a, b, c, t1[2], MD5C11, MD5S11);
 559     MD5_STEP (MD5_Go, c, d, a, b, t2[3], MD5C12, MD5S12);
 560     MD5_STEP (MD5_Go, b, c, d, a, t0[0], MD5C13, MD5S13);
 561     MD5_STEP (MD5_Go, a, b, c, d, t1[1], MD5C14, MD5S10);
 562     MD5_STEP (MD5_Go, d, a, b, c, t2[2], MD5C15, MD5S11);
 563     MD5_STEP (MD5_Go, c, d, a, b, t3[3], MD5C16, MD5S12);
 564     MD5_STEP (MD5_Go, b, c, d, a, t1[0], MD5C17, MD5S13);
 565     MD5_STEP (MD5_Go, a, b, c, d, t2[1], MD5C18, MD5S10);
 566     MD5_STEP (MD5_Go, d, a, b, c, t3[2], MD5C19, MD5S11);
 567     MD5_STEP (MD5_Go, c, d, a, b, t0[3], MD5C1a, MD5S12);
 568     MD5_STEP (MD5_Go, b, c, d, a, t2[0], MD5C1b, MD5S13);
 569     MD5_STEP (MD5_Go, a, b, c, d, t3[1], MD5C1c, MD5S10);
 570     MD5_STEP (MD5_Go, d, a, b, c, t0[2], MD5C1d, MD5S11);
 571     MD5_STEP (MD5_Go, c, d, a, b, t1[3], MD5C1e, MD5S12);
 572     MD5_STEP (MD5_Go, b, c, d, a, t3[0], MD5C1f, MD5S13);
 573
 574     MD5_STEP (MD5_H , a, b, c, d, t1[1], MD5C20, MD5S20);
 575     MD5_STEP (MD5_H , d, a, b, c, t2[0], MD5C21, MD5S21);
 576     MD5_STEP (MD5_H , c, d, a, b, t2[3], MD5C22, MD5S22);
 577     MD5_STEP (MD5_H , b, c, d, a, t3[2], MD5C23, MD5S23);
 578     MD5_STEP (MD5_H , a, b, c, d, t0[1], MD5C24, MD5S20);
 579     MD5_STEP (MD5_H , d, a, b, c, t1[0], MD5C25, MD5S21);
 580     MD5_STEP (MD5_H , c, d, a, b, t1[3], MD5C26, MD5S22);
 581     MD5_STEP (MD5_H , b, c, d, a, t2[2], MD5C27, MD5S23);
 582     MD5_STEP (MD5_H , a, b, c, d, t3[1], MD5C28, MD5S20);
 583     MD5_STEP (MD5_H , d, a, b, c, t0[0], MD5C29, MD5S21);
 584     MD5_STEP (MD5_H , c, d, a, b, t0[3], MD5C2a, MD5S22);
 585     MD5_STEP (MD5_H , b, c, d, a, t1[2], MD5C2b, MD5S23);
 586     MD5_STEP (MD5_H , a, b, c, d, t2[1], MD5C2c, MD5S20);
 587     MD5_STEP (MD5_H , d, a, b, c, t3[0], MD5C2d, MD5S21);
 588     MD5_STEP (MD5_H , c, d, a, b, t3[3], MD5C2e, MD5S22);
 589     MD5_STEP (MD5_H , b, c, d, a, t0[2], MD5C2f, MD5S23);
 590
 591     MD5_STEP (MD5_I , a, b, c, d, t0[0], MD5C30, MD5S30);
 592     MD5_STEP (MD5_I , d, a, b, c, t1[3], MD5C31, MD5S31);
 593     MD5_STEP (MD5_I , c, d, a, b, t3[2], MD5C32, MD5S32);
 594     MD5_STEP (MD5_I , b, c, d, a, t1[1], MD5C33, MD5S33);
 595     MD5_STEP (MD5_I , a, b, c, d, t3[0], MD5C34, MD5S30);
 596     MD5_STEP (MD5_I , d, a, b, c, t0[3], MD5C35, MD5S31);
 597     MD5_STEP (MD5_I , c, d, a, b, t2[2], MD5C36, MD5S32);
 598     MD5_STEP (MD5_I , b, c, d, a, t0[1], MD5C37, MD5S33);
 599     MD5_STEP (MD5_I , a, b, c, d, t2[0], MD5C38, MD5S30);
 600     MD5_STEP (MD5_I , d, a, b, c, t3[3], MD5C39, MD5S31);
 601     MD5_STEP (MD5_I , c, d, a, b, t1[2], MD5C3a, MD5S32);
 602     MD5_STEP (MD5_I , b, c, d, a, t3[1], MD5C3b, MD5S33);
 603     MD5_STEP (MD5_I , a, b, c, d, t1[0], MD5C3c, MD5S30);
 604
 605     if ((a + r_a) != search[0]) continue;
 606
 607     MD5_STEP (MD5_I , d, a, b, c, t2[3], MD5C3d, MD5S31);
 608     MD5_STEP (MD5_I , c, d, a, b, t0[2], MD5C3e, MD5S32);
 609     MD5_STEP (MD5_I , b, c, d, a, t2[1], MD5C3f, MD5S33);
 610
 611     a += r_a;
 612     b += r_b;
 613     c += r_c;
 614     d += r_d;
 615
 616     const u32x r0 = a;
 617     const u32x r1 = d;
 618     const u32x r2 = c;
 619     const u32x r3 = b;
 620
 621     #include VECT_COMPARE_S
 622   }
 623 }
 624
 625 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 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)
 626 {
 627   /**
 628    * base
 629    */
 630
 631   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 632
 633   if (gid >= gid_max) return;
 634
 635   u32 w[16];
 636
 637   w[ 0] = pws[gid].i[ 0];
 638   w[ 1] = pws[gid].i[ 1];
 639   w[ 2] = pws[gid].i[ 2];
 640   w[ 3] = pws[gid].i[ 3];
 641   w[ 4] = 0;
 642   w[ 5] = 0;
 643   w[ 6] = 0;
 644   w[ 7] = 0;
 645   w[ 8] = 0;
 646   w[ 9] = 0;
 647   w[10] = 0;
 648   w[11] = 0;
 649   w[12] = 0;
 650   w[13] = 0;
 651   w[14] = pws[gid].i[14];
 652   w[15] = 0;
 653
 654   const u32 pw_len = pws[gid].pw_len;
 655
 656   /**
 657    * main
 658    */
 659
 660   m09900m (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);
 661 }
 662
 663 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 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)
 664 {
 665   /**
 666    * base
 667    */
 668
 669   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 670
 671   if (gid >= gid_max) return;
 672
 673   u32 w[16];
 674
 675   w[ 0] = pws[gid].i[ 0];
 676   w[ 1] = pws[gid].i[ 1];
 677   w[ 2] = pws[gid].i[ 2];
 678   w[ 3] = pws[gid].i[ 3];
 679   w[ 4] = pws[gid].i[ 4];
 680   w[ 5] = pws[gid].i[ 5];
 681   w[ 6] = pws[gid].i[ 6];
 682   w[ 7] = pws[gid].i[ 7];
 683   w[ 8] = 0;
 684   w[ 9] = 0;
 685   w[10] = 0;
 686   w[11] = 0;
 687   w[12] = 0;
 688   w[13] = 0;
 689   w[14] = pws[gid].i[14];
 690   w[15] = 0;
 691
 692   const u32 pw_len = pws[gid].pw_len;
 693
 694   /**
 695    * main
 696    */
 697
 698   m09900m (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);
 699 }
 700
 701 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 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)
 702 {
 703   /**
 704    * base
 705    */
 706
 707   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 708
 709   if (gid >= gid_max) return;
 710
 711   u32 w[16];
 712
 713   w[ 0] = pws[gid].i[ 0];
 714   w[ 1] = pws[gid].i[ 1];
 715   w[ 2] = pws[gid].i[ 2];
 716   w[ 3] = pws[gid].i[ 3];
 717   w[ 4] = pws[gid].i[ 4];
 718   w[ 5] = pws[gid].i[ 5];
 719   w[ 6] = pws[gid].i[ 6];
 720   w[ 7] = pws[gid].i[ 7];
 721   w[ 8] = pws[gid].i[ 8];
 722   w[ 9] = pws[gid].i[ 9];
 723   w[10] = pws[gid].i[10];
 724   w[11] = pws[gid].i[11];
 725   w[12] = pws[gid].i[12];
 726   w[13] = pws[gid].i[13];
 727   w[14] = pws[gid].i[14];
 728   w[15] = pws[gid].i[15];
 729
 730   const u32 pw_len = pws[gid].pw_len;
 731
 732   /**
 733    * main
 734    */
 735
 736   m09900m (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);
 737 }
 738
 739 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 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)
 740 {
 741   /**
 742    * base
 743    */
 744
 745   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 746
 747   if (gid >= gid_max) return;
 748
 749   u32 w[16];
 750
 751   w[ 0] = pws[gid].i[ 0];
 752   w[ 1] = pws[gid].i[ 1];
 753   w[ 2] = pws[gid].i[ 2];
 754   w[ 3] = pws[gid].i[ 3];
 755   w[ 4] = 0;
 756   w[ 5] = 0;
 757   w[ 6] = 0;
 758   w[ 7] = 0;
 759   w[ 8] = 0;
 760   w[ 9] = 0;
 761   w[10] = 0;
 762   w[11] = 0;
 763   w[12] = 0;
 764   w[13] = 0;
 765   w[14] = pws[gid].i[14];
 766   w[15] = 0;
 767
 768   const u32 pw_len = pws[gid].pw_len;
 769
 770   /**
 771    * main
 772    */
 773
 774   m09900s (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);
 775 }
 776
 777 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 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)
 778 {
 779   /**
 780    * base
 781    */
 782
 783   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 784
 785   if (gid >= gid_max) return;
 786
 787   u32 w[16];
 788
 789   w[ 0] = pws[gid].i[ 0];
 790   w[ 1] = pws[gid].i[ 1];
 791   w[ 2] = pws[gid].i[ 2];
 792   w[ 3] = pws[gid].i[ 3];
 793   w[ 4] = pws[gid].i[ 4];
 794   w[ 5] = pws[gid].i[ 5];
 795   w[ 6] = pws[gid].i[ 6];
 796   w[ 7] = pws[gid].i[ 7];
 797   w[ 8] = 0;
 798   w[ 9] = 0;
 799   w[10] = 0;
 800   w[11] = 0;
 801   w[12] = 0;
 802   w[13] = 0;
 803   w[14] = pws[gid].i[14];
 804   w[15] = 0;
 805
 806   const u32 pw_len = pws[gid].pw_len;
 807
 808   /**
 809    * main
 810    */
 811
 812   m09900s (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);
 813 }
 814
 815 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 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)
 816 {
 817   /**
 818    * base
 819    */
 820
 821   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 822
 823   if (gid >= gid_max) return;
 824
 825   u32 w[16];
 826
 827   w[ 0] = pws[gid].i[ 0];
 828   w[ 1] = pws[gid].i[ 1];
 829   w[ 2] = pws[gid].i[ 2];
 830   w[ 3] = pws[gid].i[ 3];
 831   w[ 4] = pws[gid].i[ 4];
 832   w[ 5] = pws[gid].i[ 5];
 833   w[ 6] = pws[gid].i[ 6];
 834   w[ 7] = pws[gid].i[ 7];
 835   w[ 8] = pws[gid].i[ 8];
 836   w[ 9] = pws[gid].i[ 9];
 837   w[10] = pws[gid].i[10];
 838   w[11] = pws[gid].i[11];
 839   w[12] = pws[gid].i[12];
 840   w[13] = pws[gid].i[13];
 841   w[14] = pws[gid].i[14];
 842   w[15] = pws[gid].i[15];
 843
 844   const u32 pw_len = pws[gid].pw_len;
 845
 846   /**
 847    * main
 848    */
 849
 850   m09900s (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);
 851 }