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