nv/m08300_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA1_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 3
  20 #define DGST_R1 4
  21 #define DGST_R2 2
  22 #define DGST_R3 1
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 #ifdef  VECT_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 __device__ __constant__ bf_t c_bfs[1024];
  39
  40 __device__ static void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
  41 {
  42   u32x A = digest[0];
  43   u32x B = digest[1];
  44   u32x C = digest[2];
  45   u32x D = digest[3];
  46   u32x E = digest[4];
  47
  48   u32x w0_t = w0[0];
  49   u32x w1_t = w0[1];
  50   u32x w2_t = w0[2];
  51   u32x w3_t = w0[3];
  52   u32x w4_t = w1[0];
  53   u32x w5_t = w1[1];
  54   u32x w6_t = w1[2];
  55   u32x w7_t = w1[3];
  56   u32x w8_t = w2[0];
  57   u32x w9_t = w2[1];
  58   u32x wa_t = w2[2];
  59   u32x wb_t = w2[3];
  60   u32x wc_t = w3[0];
  61   u32x wd_t = w3[1];
  62   u32x we_t = w3[2];
  63   u32x wf_t = w3[3];
  64
  65   #undef K
  66   #define K SHA1C00
  67
  68   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
  69   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
  70   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
  71   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
  72   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
  73   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
  74   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
  75   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
  76   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
  77   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
  78   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
  79   SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
  80   SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
  81   SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
  82   SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
  83   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
  84   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
  85   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
  86   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
  87   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
  88
  89   #undef K
  90   #define K SHA1C01
  91
  92   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
  93   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
  94   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
  95   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
  96   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
  97   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
  98   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
  99   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
 100   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
 101   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
 102   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
 103   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
 104   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
 105   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
 106   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
 107   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
 108   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
 109   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
 110   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
 111   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
 112
 113   #undef K
 114   #define K SHA1C02
 115
 116   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
 117   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
 118   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
 119   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
 120   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
 121   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
 122   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
 123   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
 124   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
 125   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
 126   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
 127   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
 128   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
 129   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
 130   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
 131   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
 132   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
 133   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
 134   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
 135   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
 136
 137   #undef K
 138   #define K SHA1C03
 139
 140   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
 141   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
 142   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
 143   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
 144   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
 145   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
 146   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
 147   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
 148   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
 149   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
 150   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
 151   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
 152   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
 153   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
 154   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
 155   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
 156   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
 157   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
 158   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
 159   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
 160
 161   digest[0] += A;
 162   digest[1] += B;
 163   digest[2] += C;
 164   digest[3] += D;
 165   digest[4] += E;
 166 }
 167
 168 __device__ static void m08300m (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 169 {
 170   /**
 171    * modifier
 172    */
 173
 174   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 175   const u32 lid = threadIdx.x;
 176
 177   /**
 178    * salt
 179    */
 180
 181   const u32 salt_iter = salt_bufs[salt_pos].salt_iter;
 182
 183   u32 salt_buf0[4];
 184
 185   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 186   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
 187   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
 188   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
 189
 190   u32 salt_buf1[4];
 191
 192   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
 193   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
 194   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
 195   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
 196
 197   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 198
 199   u32 domain_buf0[4];
 200
 201   domain_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[ 0];
 202   domain_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[ 1];
 203   domain_buf0[2] = salt_bufs[salt_pos].salt_buf_pc[ 2];
 204   domain_buf0[3] = salt_bufs[salt_pos].salt_buf_pc[ 3];
 205
 206   u32 domain_buf1[4];
 207
 208   domain_buf1[0] = salt_bufs[salt_pos].salt_buf_pc[ 4];
 209   domain_buf1[1] = salt_bufs[salt_pos].salt_buf_pc[ 5];
 210   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
 211   domain_buf1[3] = 0;
 212
 213   const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
 214
 215   /**
 216    * base
 217    */
 218
 219   u32 s0[4];
 220
 221   s0[0] = salt_buf0[0];
 222   s0[1] = salt_buf0[1];
 223   s0[2] = salt_buf0[2];
 224   s0[3] = salt_buf0[3];
 225
 226   u32 s1[4];
 227
 228   s1[0] = salt_buf1[0];
 229   s1[1] = salt_buf1[1];
 230   s1[2] = salt_buf1[2];
 231   s1[3] = salt_buf1[3];
 232
 233   u32 s2[4];
 234
 235   s2[0] = 0;
 236   s2[1] = 0;
 237   s2[2] = 0;
 238   s2[3] = 0;
 239
 240   u32 s3[4];
 241
 242   s3[0] = 0;
 243   s3[1] = 0;
 244   s3[2] = 0;
 245   s3[3] = 0;
 246
 247   switch_buffer_by_offset (s0, s1, s2, s3, 1 + pw_len + domain_len + 1);
 248
 249   u32 d0[4];
 250
 251   d0[0] = domain_buf0[0];
 252   d0[1] = domain_buf0[1];
 253   d0[2] = domain_buf0[2];
 254   d0[3] = domain_buf0[3];
 255
 256   u32 d1[4];
 257
 258   d1[0] = domain_buf1[0];
 259   d1[1] = domain_buf1[1];
 260   d1[2] = domain_buf1[2];
 261   d1[3] = 0;
 262
 263   u32 d2[4];
 264
 265   d2[0] = 0;
 266   d2[1] = 0;
 267   d2[2] = 0;
 268   d2[3] = 0;
 269
 270   u32 d3[4];
 271
 272   d3[0] = 0;
 273   d3[1] = 0;
 274   d3[2] = 0;
 275   d3[3] = 0;
 276
 277   switch_buffer_by_offset (d0, d1, d2, d3, 1 + pw_len);
 278
 279   /**
 280    * loop
 281    */
 282
 283   u32x w0l = w0[0];
 284
 285   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 286   {
 287     const u32 w0r = c_bfs[il_pos].i;
 288
 289     w0[0] = w0l | w0r;
 290
 291     u32x w0_t[4];
 292
 293     w0_t[0] = w0[0];
 294     w0_t[1] = w0[1];
 295     w0_t[2] = w0[2];
 296     w0_t[3] = w0[3];
 297
 298     u32x w1_t[4];
 299
 300     w1_t[0] = w1[0];
 301     w1_t[1] = w1[1];
 302     w1_t[2] = w1[2];
 303     w1_t[3] = w1[3];
 304
 305     u32x w2_t[4];
 306
 307     w2_t[0] = w2[0];
 308     w2_t[1] = w2[1];
 309     w2_t[2] = w2[2];
 310     w2_t[3] = w2[3];
 311
 312     u32x w3_t[4];
 313
 314     w3_t[0] = w3[0];
 315     w3_t[1] = w3[1];
 316     w3_t[2] = w3[2];
 317     w3_t[3] = w3[3];
 318
 319     switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, 1);
 320
 321     w0_t[0] |= pw_len & 0xff;
 322
 323     /**
 324      * sha1
 325      */
 326
 327     u32x w0_t2[4];
 328
 329     w0_t2[0] = swap_workaround (w0_t[0] | d0[0] | s0[0]);
 330     w0_t2[1] = swap_workaround (w0_t[1] | d0[1] | s0[1]);
 331     w0_t2[2] = swap_workaround (w0_t[2] | d0[2] | s0[2]);
 332     w0_t2[3] = swap_workaround (w0_t[3] | d0[3] | s0[3]);
 333
 334     u32x w1_t2[4];
 335
 336     w1_t2[0] = swap_workaround (w1_t[0] | d1[0] | s1[0]);
 337     w1_t2[1] = swap_workaround (w1_t[1] | d1[1] | s1[1]);
 338     w1_t2[2] = swap_workaround (w1_t[2] | d1[2] | s1[2]);
 339     w1_t2[3] = swap_workaround (w1_t[3] | d1[3] | s1[3]);
 340
 341     u32x w2_t2[4];
 342
 343     w2_t2[0] = swap_workaround (w2_t[0] | d2[0] | s2[0]);
 344     w2_t2[1] = swap_workaround (w2_t[1] | d2[1] | s2[1]);
 345     w2_t2[2] = swap_workaround (w2_t[2] | d2[2] | s2[2]);
 346     w2_t2[3] = swap_workaround (w2_t[3] | d2[3] | s2[3]);
 347
 348     u32x w3_t2[4];
 349
 350     w3_t2[0] = swap_workaround (w3_t[0] | d3[0] | s3[0]);
 351     w3_t2[1] = swap_workaround (w3_t[1] | d3[1] | s3[1]);
 352     w3_t2[2] = 0;
 353     w3_t2[3] = (1 + pw_len + domain_len + 1 + salt_len) * 8;
 354
 355     u32x digest[5];
 356
 357     digest[0] = SHA1M_A;
 358     digest[1] = SHA1M_B;
 359     digest[2] = SHA1M_C;
 360     digest[3] = SHA1M_D;
 361     digest[4] = SHA1M_E;
 362
 363     sha1_transform (w0_t2, w1_t2, w2_t2, w3_t2, digest);
 364
 365     // iterations
 366
 367     for (u32 i = 0; i < salt_iter; i++)
 368     {
 369       u32x w0_t3[4];
 370
 371       w0_t3[0] = digest[0];
 372       w0_t3[1] = digest[1];
 373       w0_t3[2] = digest[2];
 374       w0_t3[3] = digest[3];
 375
 376       u32x w1_t3[4];
 377
 378       w1_t3[0] = digest[4];
 379       w1_t3[1] = swap_workaround (salt_buf0[0]);
 380       w1_t3[2] = swap_workaround (salt_buf0[1]);
 381       w1_t3[3] = swap_workaround (salt_buf0[2]);
 382
 383       u32x w2_t3[4];
 384
 385       w2_t3[0] = swap_workaround (salt_buf0[3]);
 386       w2_t3[1] = swap_workaround (salt_buf1[0]);
 387       w2_t3[2] = swap_workaround (salt_buf1[1]);
 388       w2_t3[3] = swap_workaround (salt_buf1[2]);
 389
 390       u32x w3_t3[4];
 391
 392       w3_t3[0] = swap_workaround (salt_buf1[3]);
 393       w3_t3[1] = 0;
 394       w3_t3[2] = 0;
 395       w3_t3[3] = (20 + salt_len) * 8;
 396
 397       digest[0] = SHA1M_A;
 398       digest[1] = SHA1M_B;
 399       digest[2] = SHA1M_C;
 400       digest[3] = SHA1M_D;
 401       digest[4] = SHA1M_E;
 402
 403       sha1_transform (w0_t3, w1_t3, w2_t3, w3_t3, digest);
 404     }
 405
 406     const u32x r0 = digest[3];
 407     const u32x r1 = digest[4];
 408     const u32x r2 = digest[2];
 409     const u32x r3 = digest[1];
 410
 411     #include VECT_COMPARE_M
 412   }
 413 }
 414
 415 __device__ static void m08300s (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 416 {
 417   /**
 418    * modifier
 419    */
 420
 421   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 422   const u32 lid = threadIdx.x;
 423
 424   /**
 425    * digest
 426    */
 427
 428   const u32 search[4] =
 429   {
 430     digests_buf[digests_offset].digest_buf[DGST_R0],
 431     digests_buf[digests_offset].digest_buf[DGST_R1],
 432     digests_buf[digests_offset].digest_buf[DGST_R2],
 433     digests_buf[digests_offset].digest_buf[DGST_R3]
 434   };
 435
 436   /**
 437    * salt
 438    */
 439
 440   const u32 salt_iter = salt_bufs[salt_pos].salt_iter;
 441
 442   u32 salt_buf0[4];
 443
 444   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 445   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
 446   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
 447   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
 448
 449   u32 salt_buf1[4];
 450
 451   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
 452   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
 453   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
 454   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
 455
 456   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 457
 458   u32 domain_buf0[4];
 459
 460   domain_buf0[0] = salt_bufs[salt_pos].salt_buf_pc[ 0];
 461   domain_buf0[1] = salt_bufs[salt_pos].salt_buf_pc[ 1];
 462   domain_buf0[2] = salt_bufs[salt_pos].salt_buf_pc[ 2];
 463   domain_buf0[3] = salt_bufs[salt_pos].salt_buf_pc[ 3];
 464
 465   u32 domain_buf1[4];
 466
 467   domain_buf1[0] = salt_bufs[salt_pos].salt_buf_pc[ 4];
 468   domain_buf1[1] = salt_bufs[salt_pos].salt_buf_pc[ 5];
 469   domain_buf1[2] = salt_bufs[salt_pos].salt_buf_pc[ 6];
 470   domain_buf1[3] = 0;
 471
 472   const u32 domain_len = salt_bufs[salt_pos].salt_buf_pc[ 7];
 473
 474   /**
 475    * base
 476    */
 477
 478   u32 s0[4];
 479
 480   s0[0] = salt_buf0[0];
 481   s0[1] = salt_buf0[1];
 482   s0[2] = salt_buf0[2];
 483   s0[3] = salt_buf0[3];
 484
 485   u32 s1[4];
 486
 487   s1[0] = salt_buf1[0];
 488   s1[1] = salt_buf1[1];
 489   s1[2] = salt_buf1[2];
 490   s1[3] = salt_buf1[3];
 491
 492   u32 s2[4];
 493
 494   s2[0] = 0;
 495   s2[1] = 0;
 496   s2[2] = 0;
 497   s2[3] = 0;
 498
 499   u32 s3[4];
 500
 501   s3[0] = 0;
 502   s3[1] = 0;
 503   s3[2] = 0;
 504   s3[3] = 0;
 505
 506   switch_buffer_by_offset (s0, s1, s2, s3, 1 + pw_len + domain_len + 1);
 507
 508   u32 d0[4];
 509
 510   d0[0] = domain_buf0[0];
 511   d0[1] = domain_buf0[1];
 512   d0[2] = domain_buf0[2];
 513   d0[3] = domain_buf0[3];
 514
 515   u32 d1[4];
 516
 517   d1[0] = domain_buf1[0];
 518   d1[1] = domain_buf1[1];
 519   d1[2] = domain_buf1[2];
 520   d1[3] = 0;
 521
 522   u32 d2[4];
 523
 524   d2[0] = 0;
 525   d2[1] = 0;
 526   d2[2] = 0;
 527   d2[3] = 0;
 528
 529   u32 d3[4];
 530
 531   d3[0] = 0;
 532   d3[1] = 0;
 533   d3[2] = 0;
 534   d3[3] = 0;
 535
 536   switch_buffer_by_offset (d0, d1, d2, d3, 1 + pw_len);
 537
 538   /**
 539    * loop
 540    */
 541
 542   u32x w0l = w0[0];
 543
 544   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 545   {
 546     const u32 w0r = c_bfs[il_pos].i;
 547
 548     w0[0] = w0l | w0r;
 549
 550     u32x w0_t[4];
 551
 552     w0_t[0] = w0[0];
 553     w0_t[1] = w0[1];
 554     w0_t[2] = w0[2];
 555     w0_t[3] = w0[3];
 556
 557     u32x w1_t[4];
 558
 559     w1_t[0] = w1[0];
 560     w1_t[1] = w1[1];
 561     w1_t[2] = w1[2];
 562     w1_t[3] = w1[3];
 563
 564     u32x w2_t[4];
 565
 566     w2_t[0] = w2[0];
 567     w2_t[1] = w2[1];
 568     w2_t[2] = w2[2];
 569     w2_t[3] = w2[3];
 570
 571     u32x w3_t[4];
 572
 573     w3_t[0] = w3[0];
 574     w3_t[1] = w3[1];
 575     w3_t[2] = w3[2];
 576     w3_t[3] = w3[3];
 577
 578     switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, 1);
 579
 580     w0_t[0] |= pw_len & 0xff;
 581
 582     /**
 583      * sha1
 584      */
 585
 586     u32x w0_t2[4];
 587
 588     w0_t2[0] = swap_workaround (w0_t[0] | d0[0] | s0[0]);
 589     w0_t2[1] = swap_workaround (w0_t[1] | d0[1] | s0[1]);
 590     w0_t2[2] = swap_workaround (w0_t[2] | d0[2] | s0[2]);
 591     w0_t2[3] = swap_workaround (w0_t[3] | d0[3] | s0[3]);
 592
 593     u32x w1_t2[4];
 594
 595     w1_t2[0] = swap_workaround (w1_t[0] | d1[0] | s1[0]);
 596     w1_t2[1] = swap_workaround (w1_t[1] | d1[1] | s1[1]);
 597     w1_t2[2] = swap_workaround (w1_t[2] | d1[2] | s1[2]);
 598     w1_t2[3] = swap_workaround (w1_t[3] | d1[3] | s1[3]);
 599
 600     u32x w2_t2[4];
 601
 602     w2_t2[0] = swap_workaround (w2_t[0] | d2[0] | s2[0]);
 603     w2_t2[1] = swap_workaround (w2_t[1] | d2[1] | s2[1]);
 604     w2_t2[2] = swap_workaround (w2_t[2] | d2[2] | s2[2]);
 605     w2_t2[3] = swap_workaround (w2_t[3] | d2[3] | s2[3]);
 606
 607     u32x w3_t2[4];
 608
 609     w3_t2[0] = swap_workaround (w3_t[0] | d3[0] | s3[0]);
 610     w3_t2[1] = swap_workaround (w3_t[1] | d3[1] | s3[1]);
 611     w3_t2[2] = 0;
 612     w3_t2[3] = (1 + pw_len + domain_len + 1 + salt_len) * 8;
 613
 614     u32x digest[5];
 615
 616     digest[0] = SHA1M_A;
 617     digest[1] = SHA1M_B;
 618     digest[2] = SHA1M_C;
 619     digest[3] = SHA1M_D;
 620     digest[4] = SHA1M_E;
 621
 622     sha1_transform (w0_t2, w1_t2, w2_t2, w3_t2, digest);
 623
 624     // iterations
 625
 626     for (u32 i = 0; i < salt_iter; i++)
 627     {
 628       u32x w0_t3[4];
 629
 630       w0_t3[0] = digest[0];
 631       w0_t3[1] = digest[1];
 632       w0_t3[2] = digest[2];
 633       w0_t3[3] = digest[3];
 634
 635       u32x w1_t3[4];
 636
 637       w1_t3[0] = digest[4];
 638       w1_t3[1] = swap_workaround (salt_buf0[0]);
 639       w1_t3[2] = swap_workaround (salt_buf0[1]);
 640       w1_t3[3] = swap_workaround (salt_buf0[2]);
 641
 642       u32x w2_t3[4];
 643
 644       w2_t3[0] = swap_workaround (salt_buf0[3]);
 645       w2_t3[1] = swap_workaround (salt_buf1[0]);
 646       w2_t3[2] = swap_workaround (salt_buf1[1]);
 647       w2_t3[3] = swap_workaround (salt_buf1[2]);
 648
 649       u32x w3_t3[4];
 650
 651       w3_t3[0] = swap_workaround (salt_buf1[3]);
 652       w3_t3[1] = 0;
 653       w3_t3[2] = 0;
 654       w3_t3[3] = (20 + salt_len) * 8;
 655
 656       digest[0] = SHA1M_A;
 657       digest[1] = SHA1M_B;
 658       digest[2] = SHA1M_C;
 659       digest[3] = SHA1M_D;
 660       digest[4] = SHA1M_E;
 661
 662       sha1_transform (w0_t3, w1_t3, w2_t3, w3_t3, digest);
 663     }
 664
 665     const u32x r0 = digest[3];
 666     const u32x r1 = digest[4];
 667     const u32x r2 = digest[2];
 668     const u32x r3 = digest[1];
 669
 670     #include VECT_COMPARE_S
 671   }
 672 }
 673
 674 extern "C" __global__ void __launch_bounds__ (256, 1) m08300_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 675 {
 676   /**
 677    * base
 678    */
 679
 680   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 681
 682   if (gid >= gid_max) return;
 683
 684   u32x w0[4];
 685
 686   w0[0] = pws[gid].i[ 0];
 687   w0[1] = pws[gid].i[ 1];
 688   w0[2] = pws[gid].i[ 2];
 689   w0[3] = pws[gid].i[ 3];
 690
 691   u32x w1[4];
 692
 693   w1[0] = 0;
 694   w1[1] = 0;
 695   w1[2] = 0;
 696   w1[3] = 0;
 697
 698   u32x w2[4];
 699
 700   w2[0] = 0;
 701   w2[1] = 0;
 702   w2[2] = 0;
 703   w2[3] = 0;
 704
 705   u32x w3[4];
 706
 707   w3[0] = 0;
 708   w3[1] = 0;
 709   w3[2] = 0;
 710   w3[3] = 0;
 711
 712   const u32 pw_len = pws[gid].pw_len;
 713
 714   /**
 715    * main
 716    */
 717
 718   m08300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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);
 719 }
 720
 721 extern "C" __global__ void __launch_bounds__ (256, 1) m08300_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 722 {
 723   /**
 724    * base
 725    */
 726
 727   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 728
 729   if (gid >= gid_max) return;
 730
 731   u32x w0[4];
 732
 733   w0[0] = pws[gid].i[ 0];
 734   w0[1] = pws[gid].i[ 1];
 735   w0[2] = pws[gid].i[ 2];
 736   w0[3] = pws[gid].i[ 3];
 737
 738   u32x w1[4];
 739
 740   w1[0] = pws[gid].i[ 4];
 741   w1[1] = pws[gid].i[ 5];
 742   w1[2] = pws[gid].i[ 6];
 743   w1[3] = pws[gid].i[ 7];
 744
 745   u32x w2[4];
 746
 747   w2[0] = 0;
 748   w2[1] = 0;
 749   w2[2] = 0;
 750   w2[3] = 0;
 751
 752   u32x w3[4];
 753
 754   w3[0] = 0;
 755   w3[1] = 0;
 756   w3[2] = 0;
 757   w3[3] = 0;
 758
 759   const u32 pw_len = pws[gid].pw_len;
 760
 761   /**
 762    * main
 763    */
 764
 765   m08300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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);
 766 }
 767
 768 extern "C" __global__ void __launch_bounds__ (256, 1) m08300_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 769 {
 770   /**
 771    * base
 772    */
 773
 774   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 775
 776   if (gid >= gid_max) return;
 777
 778   u32x w0[4];
 779
 780   w0[0] = pws[gid].i[ 0];
 781   w0[1] = pws[gid].i[ 1];
 782   w0[2] = pws[gid].i[ 2];
 783   w0[3] = pws[gid].i[ 3];
 784
 785   u32x w1[4];
 786
 787   w1[0] = pws[gid].i[ 4];
 788   w1[1] = pws[gid].i[ 5];
 789   w1[2] = pws[gid].i[ 6];
 790   w1[3] = pws[gid].i[ 7];
 791
 792   u32x w2[4];
 793
 794   w2[0] = pws[gid].i[ 8];
 795   w2[1] = pws[gid].i[ 9];
 796   w2[2] = pws[gid].i[10];
 797   w2[3] = pws[gid].i[11];
 798
 799   u32x w3[4];
 800
 801   w3[0] = pws[gid].i[12];
 802   w3[1] = pws[gid].i[13];
 803   w3[2] = 0;
 804   w3[3] = 0;
 805
 806   const u32 pw_len = pws[gid].pw_len;
 807
 808   /**
 809    * main
 810    */
 811
 812   m08300m (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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) m08300_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 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   u32x w0[4];
 826
 827   w0[0] = pws[gid].i[ 0];
 828   w0[1] = pws[gid].i[ 1];
 829   w0[2] = pws[gid].i[ 2];
 830   w0[3] = pws[gid].i[ 3];
 831
 832   u32x w1[4];
 833
 834   w1[0] = 0;
 835   w1[1] = 0;
 836   w1[2] = 0;
 837   w1[3] = 0;
 838
 839   u32x w2[4];
 840
 841   w2[0] = 0;
 842   w2[1] = 0;
 843   w2[2] = 0;
 844   w2[3] = 0;
 845
 846   u32x w3[4];
 847
 848   w3[0] = 0;
 849   w3[1] = 0;
 850   w3[2] = 0;
 851   w3[3] = 0;
 852
 853   const u32 pw_len = pws[gid].pw_len;
 854
 855   /**
 856    * main
 857    */
 858
 859   m08300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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);
 860 }
 861
 862 extern "C" __global__ void __launch_bounds__ (256, 1) m08300_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 863 {
 864   /**
 865    * base
 866    */
 867
 868   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 869
 870   if (gid >= gid_max) return;
 871
 872   u32x w0[4];
 873
 874   w0[0] = pws[gid].i[ 0];
 875   w0[1] = pws[gid].i[ 1];
 876   w0[2] = pws[gid].i[ 2];
 877   w0[3] = pws[gid].i[ 3];
 878
 879   u32x w1[4];
 880
 881   w1[0] = pws[gid].i[ 4];
 882   w1[1] = pws[gid].i[ 5];
 883   w1[2] = pws[gid].i[ 6];
 884   w1[3] = pws[gid].i[ 7];
 885
 886   u32x w2[4];
 887
 888   w2[0] = 0;
 889   w2[1] = 0;
 890   w2[2] = 0;
 891   w2[3] = 0;
 892
 893   u32x w3[4];
 894
 895   w3[0] = 0;
 896   w3[1] = 0;
 897   w3[2] = 0;
 898   w3[3] = 0;
 899
 900   const u32 pw_len = pws[gid].pw_len;
 901
 902   /**
 903    * main
 904    */
 905
 906   m08300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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);
 907 }
 908
 909 extern "C" __global__ void __launch_bounds__ (256, 1) m08300_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, const void *tmps, void *hooks, const u32 *bitmaps_buf_s1_a, const u32 *bitmaps_buf_s1_b, const u32 *bitmaps_buf_s1_c, const u32 *bitmaps_buf_s1_d, const u32 *bitmaps_buf_s2_a, const u32 *bitmaps_buf_s2_b, const u32 *bitmaps_buf_s2_c, const u32 *bitmaps_buf_s2_d, plain_t *plains_buf, const digest_t *digests_buf, u32 *hashes_shown, const salt_t *salt_bufs, const void *esalt_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 910 {
 911   /**
 912    * base
 913    */
 914
 915   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 916
 917   if (gid >= gid_max) return;
 918
 919   u32x w0[4];
 920
 921   w0[0] = pws[gid].i[ 0];
 922   w0[1] = pws[gid].i[ 1];
 923   w0[2] = pws[gid].i[ 2];
 924   w0[3] = pws[gid].i[ 3];
 925
 926   u32x w1[4];
 927
 928   w1[0] = pws[gid].i[ 4];
 929   w1[1] = pws[gid].i[ 5];
 930   w1[2] = pws[gid].i[ 6];
 931   w1[3] = pws[gid].i[ 7];
 932
 933   u32x w2[4];
 934
 935   w2[0] = pws[gid].i[ 8];
 936   w2[1] = pws[gid].i[ 9];
 937   w2[2] = pws[gid].i[10];
 938   w2[3] = pws[gid].i[11];
 939
 940   u32x w3[4];
 941
 942   w3[0] = pws[gid].i[12];
 943   w3[1] = pws[gid].i[13];
 944   w3[2] = 0;
 945   w3[3] = 0;
 946
 947   const u32 pw_len = pws[gid].pw_len;
 948
 949   /**
 950    * main
 951    */
 952
 953   m08300s (w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, 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);
 954 }