nv/m07600_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_SIZE1
  39 #define uint_to_hex_lower8(i) l_bin2asc[(i)]
  40 #endif
  41
  42 #ifdef VECT_SIZE2
  43 #define uint_to_hex_lower8(i) u32x (l_bin2asc[(i).x], l_bin2asc[(i).y])
  44 #endif
  45
  46 __device__ __constant__ char c_bin2asc[16] = { '0', '1', '2', '3', '4', '5', '6', '7', '8', '9', 'a', 'b', 'c', 'd', 'e', 'f' };
  47
  48 __device__ __shared__ short l_bin2asc[256];
  49
  50 __device__ __constant__ bf_t c_bfs[1024];
  51
  52 __device__ static void m07600m (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)
  53 {
  54   /**
  55    * modifier
  56    */
  57
  58   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
  59   const u32 lid = threadIdx.x;
  60
  61   /**
  62    * salt
  63    */
  64
  65   u32 salt_buf0[4];
  66
  67   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
  68   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
  69   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
  70   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
  71
  72   u32 salt_buf1[4];
  73
  74   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
  75   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
  76   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
  77   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
  78
  79   const u32 salt_len = salt_bufs[salt_pos].salt_len;
  80
  81   const u32 total_len = (salt_len + 40) * 8;
  82
  83   /**
  84    * loop
  85    */
  86
  87   u32x w0l = w0[0];
  88
  89   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
  90   {
  91     const u32 w0r = c_bfs[il_pos].i;
  92
  93     w0[0] = w0l | w0r;
  94
  95     /**
  96      * sha1
  97      */
  98
  99     u32x w0_t = w0[0];
 100     u32x w1_t = w0[1];
 101     u32x w2_t = w0[2];
 102     u32x w3_t = w0[3];
 103     u32x w4_t = w1[0];
 104     u32x w5_t = w1[1];
 105     u32x w6_t = w1[2];
 106     u32x w7_t = w1[3];
 107     u32x w8_t = w2[0];
 108     u32x w9_t = w2[1];
 109     u32x wa_t = w2[2];
 110     u32x wb_t = w2[3];
 111     u32x wc_t = w3[0];
 112     u32x wd_t = w3[1];
 113     u32x we_t = 0;
 114     u32x wf_t = pw_len * 8;
 115
 116     u32x a = SHA1M_A;
 117     u32x b = SHA1M_B;
 118     u32x c = SHA1M_C;
 119     u32x d = SHA1M_D;
 120     u32x e = SHA1M_E;
 121
 122     #undef K
 123     #define K SHA1C00
 124
 125     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 126     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 127     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 128     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 129     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 130     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 131     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 132     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 133     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 134     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 135     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 136     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 137     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 138     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 139     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 140     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 141     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 142     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 143     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 144     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 145
 146     #undef K
 147     #define K SHA1C01
 148
 149     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 150     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 151     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 152     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 153     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 154     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 155     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 156     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 157     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 158     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 159     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 160     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 161     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 162     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 163     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 164     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 165     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 166     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 167     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 168     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 169
 170     #undef K
 171     #define K SHA1C02
 172
 173     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 174     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 175     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 176     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 177     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 178     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 179     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 180     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 181     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 182     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 183     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 184     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 185     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 186     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 187     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 188     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 189     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 190     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 191     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 192     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 193
 194     #undef K
 195     #define K SHA1C03
 196
 197     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 198     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 199     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 200     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 201     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 202     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 203     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 204     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 205     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 206     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 207     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 208     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 209     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 210     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 211     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 212     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 213     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 214     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 215     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 216     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 217
 218     a += SHA1M_A;
 219     b += SHA1M_B;
 220     c += SHA1M_C;
 221     d += SHA1M_D;
 222     e += SHA1M_E;
 223
 224     /**
 225      * Prepend salt
 226      */
 227
 228     u32x w0t[4];
 229
 230     w0t[0] = uint_to_hex_lower8 ((a >> 24) & 255) <<  0
 231            | uint_to_hex_lower8 ((a >> 16) & 255) << 16;
 232     w0t[1] = uint_to_hex_lower8 ((a >>  8) & 255) <<  0
 233            | uint_to_hex_lower8 ((a >>  0) & 255) << 16;
 234     w0t[2] = uint_to_hex_lower8 ((b >> 24) & 255) <<  0
 235            | uint_to_hex_lower8 ((b >> 16) & 255) << 16;
 236     w0t[3] = uint_to_hex_lower8 ((b >>  8) & 255) <<  0
 237            | uint_to_hex_lower8 ((b >>  0) & 255) << 16;
 238
 239     u32x w1t[4];
 240
 241     w1t[0] = uint_to_hex_lower8 ((c >> 24) & 255) <<  0
 242            | uint_to_hex_lower8 ((c >> 16) & 255) << 16;
 243     w1t[1] = uint_to_hex_lower8 ((c >>  8) & 255) <<  0
 244            | uint_to_hex_lower8 ((c >>  0) & 255) << 16;
 245     w1t[2] = uint_to_hex_lower8 ((d >> 24) & 255) <<  0
 246            | uint_to_hex_lower8 ((d >> 16) & 255) << 16;
 247     w1t[3] = uint_to_hex_lower8 ((d >>  8) & 255) <<  0
 248            | uint_to_hex_lower8 ((d >>  0) & 255) << 16;
 249
 250     u32x w2t[2];
 251
 252     w2t[0] = uint_to_hex_lower8 ((e >> 24) & 255) <<  0
 253            | uint_to_hex_lower8 ((e >> 16) & 255) << 16;
 254     w2t[1] = uint_to_hex_lower8 ((e >>  8) & 255) <<  0
 255            | uint_to_hex_lower8 ((e >>  0) & 255) << 16;
 256
 257     w0_t = salt_buf0[0];
 258     w1_t = salt_buf0[1];
 259     w2_t = salt_buf0[2];
 260     w3_t = salt_buf0[3];
 261     w4_t = salt_buf1[0];
 262     w5_t = salt_buf1[1];
 263     w6_t = salt_buf1[2];
 264     w7_t = salt_buf1[3];
 265     w8_t = w0t[0];
 266     w9_t = w0t[1];
 267     wa_t = w0t[2];
 268     wb_t = w0t[3];
 269     wc_t = w1t[0];
 270     wd_t = w1t[1];
 271     we_t = w1t[2];
 272     wf_t = w1t[3];
 273
 274     /**
 275      * 2nd SHA1
 276      */
 277
 278     // 1st transform
 279
 280     w0_t = swap_workaround (w0_t);
 281     w1_t = swap_workaround (w1_t);
 282     w2_t = swap_workaround (w2_t);
 283     w3_t = swap_workaround (w3_t);
 284     w4_t = swap_workaround (w4_t);
 285     w5_t = swap_workaround (w5_t);
 286     w6_t = swap_workaround (w6_t);
 287     w7_t = swap_workaround (w7_t);
 288     w8_t = swap_workaround (w8_t);
 289     w9_t = swap_workaround (w9_t);
 290     wa_t = swap_workaround (wa_t);
 291     wb_t = swap_workaround (wb_t);
 292     wc_t = swap_workaround (wc_t);
 293     wd_t = swap_workaround (wd_t);
 294     we_t = swap_workaround (we_t);
 295     wf_t = swap_workaround (wf_t);
 296
 297     a = SHA1M_A;
 298     b = SHA1M_B;
 299     c = SHA1M_C;
 300     d = SHA1M_D;
 301     e = SHA1M_E;
 302
 303     #undef K
 304     #define K SHA1C00
 305
 306     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 307     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 308     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 309     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 310     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 311     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 312     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 313     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 314     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 315     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 316     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 317     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 318     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 319     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 320     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 321     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 322     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 323     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 324     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 325     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 326
 327     #undef K
 328     #define K SHA1C01
 329
 330     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 331     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 332     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 333     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 334     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 335     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 336     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 337     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 338     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 339     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 340     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 341     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 342     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 343     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 344     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 345     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 346     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 347     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 348     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 349     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 350
 351     #undef K
 352     #define K SHA1C02
 353
 354     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 355     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 356     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 357     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 358     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 359     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 360     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 361     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 362     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 363     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 364     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 365     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 366     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 367     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 368     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 369     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 370     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 371     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 372     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 373     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 374
 375     #undef K
 376     #define K SHA1C03
 377
 378     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 379     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 380     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 381     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 382     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 383     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 384     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 385     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 386     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 387     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 388     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 389     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 390     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 391     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 392     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 393     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 394     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 395     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 396     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 397     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 398
 399     a += SHA1M_A;
 400     b += SHA1M_B;
 401     c += SHA1M_C;
 402     d += SHA1M_D;
 403     e += SHA1M_E;
 404
 405     u32x r_a = a;
 406     u32x r_b = b;
 407     u32x r_c = c;
 408     u32x r_d = d;
 409     u32x r_e = e;
 410
 411     // 2nd transform
 412
 413     w0_t = swap_workaround (w2t[0]);
 414     w1_t = swap_workaround (w2t[1]);
 415     w2_t = 0x80000000;
 416     w3_t = 0;
 417     w4_t = 0;
 418     w5_t = 0;
 419     w6_t = 0;
 420     w7_t = 0;
 421     w8_t = 0;
 422     w9_t = 0;
 423     wa_t = 0;
 424     wb_t = 0;
 425     wc_t = 0;
 426     wd_t = 0;
 427     we_t = 0;
 428     wf_t = total_len;
 429
 430     #undef K
 431     #define K SHA1C00
 432
 433     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 434     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 435     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 436     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 437     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 438     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 439     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 440     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 441     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 442     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 443     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 444     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 445     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 446     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 447     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 448     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 449     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 450     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 451     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 452     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 453
 454     #undef K
 455     #define K SHA1C01
 456
 457     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 458     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 459     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 460     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 461     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 462     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 463     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 464     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 465     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 466     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 467     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 468     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 469     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 470     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 471     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 472     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 473     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 474     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 475     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 476     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 477
 478     #undef K
 479     #define K SHA1C02
 480
 481     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 482     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 483     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 484     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 485     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 486     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 487     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 488     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 489     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 490     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 491     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 492     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 493     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 494     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 495     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 496     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 497     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 498     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 499     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 500     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 501
 502     #undef K
 503     #define K SHA1C03
 504
 505     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 506     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 507     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 508     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 509     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 510     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 511     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 512     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 513     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 514     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 515     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 516     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 517     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 518     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 519     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 520     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 521     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 522     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 523     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 524     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 525
 526     a += r_a;
 527     b += r_b;
 528     c += r_c;
 529     d += r_d;
 530     e += r_e;
 531
 532     const u32x r0 = d;
 533     const u32x r1 = e;
 534     const u32x r2 = c;
 535     const u32x r3 = b;
 536
 537     #include VECT_COMPARE_M
 538   }
 539 }
 540
 541 __device__ static void m07600s (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)
 542 {
 543   /**
 544    * modifier
 545    */
 546
 547   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 548   const u32 lid = threadIdx.x;
 549
 550   /**
 551    * salt
 552    */
 553
 554   u32 salt_buf0[4];
 555
 556   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[ 0];
 557   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[ 1];
 558   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[ 2];
 559   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[ 3];
 560
 561   u32 salt_buf1[4];
 562
 563   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[ 4];
 564   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[ 5];
 565   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[ 6];
 566   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[ 7];
 567
 568   const u32 salt_len = salt_bufs[salt_pos].salt_len;
 569
 570   const u32 total_len = (salt_len + 40) * 8;
 571
 572   /**
 573    * digest
 574    */
 575
 576   const u32 search[4] =
 577   {
 578     digests_buf[digests_offset].digest_buf[DGST_R0],
 579     digests_buf[digests_offset].digest_buf[DGST_R1],
 580     digests_buf[digests_offset].digest_buf[DGST_R2],
 581     digests_buf[digests_offset].digest_buf[DGST_R3]
 582   };
 583
 584   /**
 585    * reverse
 586    */
 587
 588   const u32 e_rev = rotl32 (search[1], 2u);
 589
 590   /**
 591    * loop
 592    */
 593
 594   u32x w0l = w0[0];
 595
 596   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 597   {
 598     const u32 w0r = c_bfs[il_pos].i;
 599
 600     w0[0] = w0l | w0r;
 601
 602     /**
 603      * sha1
 604      */
 605
 606     u32x w0_t = w0[0];
 607     u32x w1_t = w0[1];
 608     u32x w2_t = w0[2];
 609     u32x w3_t = w0[3];
 610     u32x w4_t = w1[0];
 611     u32x w5_t = w1[1];
 612     u32x w6_t = w1[2];
 613     u32x w7_t = w1[3];
 614     u32x w8_t = w2[0];
 615     u32x w9_t = w2[1];
 616     u32x wa_t = w2[2];
 617     u32x wb_t = w2[3];
 618     u32x wc_t = w3[0];
 619     u32x wd_t = w3[1];
 620     u32x we_t = 0;
 621     u32x wf_t = pw_len * 8;
 622
 623     u32x a = SHA1M_A;
 624     u32x b = SHA1M_B;
 625     u32x c = SHA1M_C;
 626     u32x d = SHA1M_D;
 627     u32x e = SHA1M_E;
 628
 629     #undef K
 630     #define K SHA1C00
 631
 632     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 633     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 634     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 635     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 636     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 637     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 638     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 639     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 640     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 641     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 642     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 643     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 644     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 645     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 646     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 647     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 648     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 649     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 650     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 651     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 652
 653     #undef K
 654     #define K SHA1C01
 655
 656     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 657     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 658     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 659     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 660     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 661     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 662     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 663     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 664     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 665     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 666     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 667     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 668     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 669     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 670     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 671     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 672     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 673     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 674     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 675     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 676
 677     #undef K
 678     #define K SHA1C02
 679
 680     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 681     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 682     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 683     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 684     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 685     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 686     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 687     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 688     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 689     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 690     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 691     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 692     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 693     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 694     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 695     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 696     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 697     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 698     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 699     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 700
 701     #undef K
 702     #define K SHA1C03
 703
 704     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 705     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 706     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 707     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 708     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 709     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 710     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 711     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 712     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 713     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 714     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 715     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 716     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 717     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 718     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 719     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 720     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 721     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 722     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 723     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 724
 725     a += SHA1M_A;
 726     b += SHA1M_B;
 727     c += SHA1M_C;
 728     d += SHA1M_D;
 729     e += SHA1M_E;
 730
 731     /**
 732      * Prepend salt
 733      */
 734
 735     u32x w0t[4];
 736
 737     w0t[0] = uint_to_hex_lower8 ((a >> 24) & 255) <<  0
 738            | uint_to_hex_lower8 ((a >> 16) & 255) << 16;
 739     w0t[1] = uint_to_hex_lower8 ((a >>  8) & 255) <<  0
 740            | uint_to_hex_lower8 ((a >>  0) & 255) << 16;
 741     w0t[2] = uint_to_hex_lower8 ((b >> 24) & 255) <<  0
 742            | uint_to_hex_lower8 ((b >> 16) & 255) << 16;
 743     w0t[3] = uint_to_hex_lower8 ((b >>  8) & 255) <<  0
 744            | uint_to_hex_lower8 ((b >>  0) & 255) << 16;
 745
 746     u32x w1t[4];
 747
 748     w1t[0] = uint_to_hex_lower8 ((c >> 24) & 255) <<  0
 749            | uint_to_hex_lower8 ((c >> 16) & 255) << 16;
 750     w1t[1] = uint_to_hex_lower8 ((c >>  8) & 255) <<  0
 751            | uint_to_hex_lower8 ((c >>  0) & 255) << 16;
 752     w1t[2] = uint_to_hex_lower8 ((d >> 24) & 255) <<  0
 753            | uint_to_hex_lower8 ((d >> 16) & 255) << 16;
 754     w1t[3] = uint_to_hex_lower8 ((d >>  8) & 255) <<  0
 755            | uint_to_hex_lower8 ((d >>  0) & 255) << 16;
 756
 757     u32x w2t[2];
 758
 759     w2t[0] = uint_to_hex_lower8 ((e >> 24) & 255) <<  0
 760            | uint_to_hex_lower8 ((e >> 16) & 255) << 16;
 761     w2t[1] = uint_to_hex_lower8 ((e >>  8) & 255) <<  0
 762            | uint_to_hex_lower8 ((e >>  0) & 255) << 16;
 763
 764     w0_t = salt_buf0[0];
 765     w1_t = salt_buf0[1];
 766     w2_t = salt_buf0[2];
 767     w3_t = salt_buf0[3];
 768     w4_t = salt_buf1[0];
 769     w5_t = salt_buf1[1];
 770     w6_t = salt_buf1[2];
 771     w7_t = salt_buf1[3];
 772     w8_t = w0t[0];
 773     w9_t = w0t[1];
 774     wa_t = w0t[2];
 775     wb_t = w0t[3];
 776     wc_t = w1t[0];
 777     wd_t = w1t[1];
 778     we_t = w1t[2];
 779     wf_t = w1t[3];
 780
 781     /**
 782      * 2nd SHA1
 783      */
 784
 785     // 1st transform
 786
 787     w0_t = swap_workaround (w0_t);
 788     w1_t = swap_workaround (w1_t);
 789     w2_t = swap_workaround (w2_t);
 790     w3_t = swap_workaround (w3_t);
 791     w4_t = swap_workaround (w4_t);
 792     w5_t = swap_workaround (w5_t);
 793     w6_t = swap_workaround (w6_t);
 794     w7_t = swap_workaround (w7_t);
 795     w8_t = swap_workaround (w8_t);
 796     w9_t = swap_workaround (w9_t);
 797     wa_t = swap_workaround (wa_t);
 798     wb_t = swap_workaround (wb_t);
 799     wc_t = swap_workaround (wc_t);
 800     wd_t = swap_workaround (wd_t);
 801     we_t = swap_workaround (we_t);
 802     wf_t = swap_workaround (wf_t);
 803
 804     a = SHA1M_A;
 805     b = SHA1M_B;
 806     c = SHA1M_C;
 807     d = SHA1M_D;
 808     e = SHA1M_E;
 809
 810     #undef K
 811     #define K SHA1C00
 812
 813     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 814     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 815     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 816     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 817     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 818     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 819     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 820     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 821     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 822     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 823     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 824     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 825     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 826     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 827     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 828     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 829     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 830     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 831     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 832     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 833
 834     #undef K
 835     #define K SHA1C01
 836
 837     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 838     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 839     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 840     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 841     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 842     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 843     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 844     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 845     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 846     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 847     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 848     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 849     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 850     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 851     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 852     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 853     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 854     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 855     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 856     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 857
 858     #undef K
 859     #define K SHA1C02
 860
 861     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 862     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 863     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 864     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 865     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 866     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 867     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 868     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 869     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 870     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 871     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 872     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
 873     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
 874     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
 875     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
 876     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
 877     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
 878     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
 879     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
 880     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
 881
 882     #undef K
 883     #define K SHA1C03
 884
 885     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
 886     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
 887     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
 888     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
 889     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
 890     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
 891     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
 892     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
 893     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
 894     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
 895     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
 896     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
 897     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
 898     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
 899     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
 900     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
 901     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
 902     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
 903     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
 904     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
 905
 906     a += SHA1M_A;
 907     b += SHA1M_B;
 908     c += SHA1M_C;
 909     d += SHA1M_D;
 910     e += SHA1M_E;
 911
 912     u32x r_a = a;
 913     u32x r_b = b;
 914     u32x r_c = c;
 915     u32x r_d = d;
 916     u32x r_e = e;
 917
 918     // 2nd transform
 919
 920     w0_t = swap_workaround (w2t[0]);
 921     w1_t = swap_workaround (w2t[1]);
 922     w2_t = 0x80000000;
 923     w3_t = 0;
 924     w4_t = 0;
 925     w5_t = 0;
 926     w6_t = 0;
 927     w7_t = 0;
 928     w8_t = 0;
 929     w9_t = 0;
 930     wa_t = 0;
 931     wb_t = 0;
 932     wc_t = 0;
 933     wd_t = 0;
 934     we_t = 0;
 935     wf_t = total_len;
 936
 937     #undef K
 938     #define K SHA1C00
 939
 940     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w0_t);
 941     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w1_t);
 942     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w2_t);
 943     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w3_t);
 944     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w4_t);
 945     SHA1_STEP (SHA1_F0o, a, b, c, d, e, w5_t);
 946     SHA1_STEP (SHA1_F0o, e, a, b, c, d, w6_t);
 947     SHA1_STEP (SHA1_F0o, d, e, a, b, c, w7_t);
 948     SHA1_STEP (SHA1_F0o, c, d, e, a, b, w8_t);
 949     SHA1_STEP (SHA1_F0o, b, c, d, e, a, w9_t);
 950     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wa_t);
 951     SHA1_STEP (SHA1_F0o, e, a, b, c, d, wb_t);
 952     SHA1_STEP (SHA1_F0o, d, e, a, b, c, wc_t);
 953     SHA1_STEP (SHA1_F0o, c, d, e, a, b, wd_t);
 954     SHA1_STEP (SHA1_F0o, b, c, d, e, a, we_t);
 955     SHA1_STEP (SHA1_F0o, a, b, c, d, e, wf_t);
 956     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, e, a, b, c, d, w0_t);
 957     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, d, e, a, b, c, w1_t);
 958     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, c, d, e, a, b, w2_t);
 959     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, b, c, d, e, a, w3_t);
 960
 961     #undef K
 962     #define K SHA1C01
 963
 964     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w4_t);
 965     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w5_t);
 966     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w6_t);
 967     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w7_t);
 968     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w8_t);
 969     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w9_t);
 970     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wa_t);
 971     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wb_t);
 972     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wc_t);
 973     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wd_t);
 974     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, we_t);
 975     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wf_t);
 976     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w0_t);
 977     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w1_t);
 978     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w2_t);
 979     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w3_t);
 980     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w4_t);
 981     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w5_t);
 982     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w6_t);
 983     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w7_t);
 984
 985     #undef K
 986     #define K SHA1C02
 987
 988     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w8_t);
 989     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w9_t);
 990     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wa_t);
 991     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wb_t);
 992     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wc_t);
 993     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, wd_t);
 994     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, we_t);
 995     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, wf_t);
 996     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w0_t);
 997     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w1_t);
 998     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w2_t);
 999     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w3_t);
1000     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w4_t);
1001     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, w5_t);
1002     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, w6_t);
1003     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, a, b, c, d, e, w7_t);
1004     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, e, a, b, c, d, w8_t);
1005     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, d, e, a, b, c, w9_t);
1006     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, c, d, e, a, b, wa_t);
1007     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, b, c, d, e, a, wb_t);
1008
1009     #undef K
1010     #define K SHA1C03
1011
1012     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wc_t);
1013     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wd_t);
1014     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, we_t);
1015     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, wf_t);
1016     w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w0_t);
1017     w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w1_t);
1018     w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w2_t);
1019     w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w3_t);
1020     w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w4_t);
1021     w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, w5_t);
1022     w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, w6_t);
1023     w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, w7_t);
1024     w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, w8_t);
1025     w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, w9_t);
1026     wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wa_t);
1027     wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, a, b, c, d, e, wb_t);
1028     wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, e, a, b, c, d, wc_t);
1029     wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, d, e, a, b, c, wd_t);
1030     we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, c, d, e, a, b, we_t);
1031     wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, b, c, d, e, a, wf_t);
1032
1033     a += r_a;
1034     b += r_b;
1035     c += r_c;
1036     d += r_d;
1037     e += r_e;
1038
1039     const u32x r0 = d;
1040     const u32x r1 = e;
1041     const u32x r2 = c;
1042     const u32x r3 = b;
1043
1044     #include VECT_COMPARE_S
1045   }
1046 }
1047
1048 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1049 {
1050   /**
1051    * base
1052    */
1053
1054   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1055
1056   /**
1057    * modifier
1058    */
1059
1060   const u32 lid = threadIdx.x;
1061
1062
1063   u32x w0[4];
1064
1065   w0[0] = pws[gid].i[ 0];
1066   w0[1] = pws[gid].i[ 1];
1067   w0[2] = pws[gid].i[ 2];
1068   w0[3] = pws[gid].i[ 3];
1069
1070   u32x w1[4];
1071
1072   w1[0] = 0;
1073   w1[1] = 0;
1074   w1[2] = 0;
1075   w1[3] = 0;
1076
1077   u32x w2[4];
1078
1079   w2[0] = 0;
1080   w2[1] = 0;
1081   w2[2] = 0;
1082   w2[3] = 0;
1083
1084   u32x w3[4];
1085
1086   w3[0] = 0;
1087   w3[1] = 0;
1088   w3[2] = 0;
1089   w3[3] = pws[gid].i[15];
1090
1091   const u32 pw_len = pws[gid].pw_len;
1092
1093   /**
1094    * bin2asc table
1095    */
1096
1097   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1098                  | c_bin2asc[(lid >> 4) & 15] << 0;
1099
1100   __syncthreads ();
1101
1102   if (gid >= gid_max) return;
1103
1104   /**
1105    * main
1106    */
1107
1108   m07600m (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);
1109 }
1110
1111 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1112 {
1113   /**
1114    * base
1115    */
1116
1117   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1118
1119   /**
1120    * modifier
1121    */
1122
1123   const u32 lid = threadIdx.x;
1124
1125   u32x w0[4];
1126
1127   w0[0] = pws[gid].i[ 0];
1128   w0[1] = pws[gid].i[ 1];
1129   w0[2] = pws[gid].i[ 2];
1130   w0[3] = pws[gid].i[ 3];
1131
1132   u32x w1[4];
1133
1134   w1[0] = pws[gid].i[ 4];
1135   w1[1] = pws[gid].i[ 5];
1136   w1[2] = pws[gid].i[ 6];
1137   w1[3] = pws[gid].i[ 7];
1138
1139   u32x w2[4];
1140
1141   w2[0] = 0;
1142   w2[1] = 0;
1143   w2[2] = 0;
1144   w2[3] = 0;
1145
1146   u32x w3[4];
1147
1148   w3[0] = 0;
1149   w3[1] = 0;
1150   w3[2] = 0;
1151   w3[3] = pws[gid].i[15];
1152
1153   const u32 pw_len = pws[gid].pw_len;
1154
1155   /**
1156    * bin2asc table
1157    */
1158
1159   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1160                  | c_bin2asc[(lid >> 4) & 15] << 0;
1161
1162   __syncthreads ();
1163
1164   if (gid >= gid_max) return;
1165
1166   /**
1167    * main
1168    */
1169
1170   m07600m (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);
1171 }
1172
1173 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1174 {
1175   /**
1176    * base
1177    */
1178
1179   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1180
1181   /**
1182    * modifier
1183    */
1184
1185   const u32 lid = threadIdx.x;
1186
1187   u32x w0[4];
1188
1189   w0[0] = pws[gid].i[ 0];
1190   w0[1] = pws[gid].i[ 1];
1191   w0[2] = pws[gid].i[ 2];
1192   w0[3] = pws[gid].i[ 3];
1193
1194   u32x w1[4];
1195
1196   w1[0] = pws[gid].i[ 4];
1197   w1[1] = pws[gid].i[ 5];
1198   w1[2] = pws[gid].i[ 6];
1199   w1[3] = pws[gid].i[ 7];
1200
1201   u32x w2[4];
1202
1203   w2[0] = pws[gid].i[ 8];
1204   w2[1] = pws[gid].i[ 9];
1205   w2[2] = pws[gid].i[10];
1206   w2[3] = pws[gid].i[11];
1207
1208   u32x w3[4];
1209
1210   w3[0] = pws[gid].i[12];
1211   w3[1] = pws[gid].i[13];
1212   w3[2] = pws[gid].i[14];
1213   w3[3] = pws[gid].i[15];
1214
1215   const u32 pw_len = pws[gid].pw_len;
1216
1217   /**
1218    * bin2asc table
1219    */
1220
1221   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1222                  | c_bin2asc[(lid >> 4) & 15] << 0;
1223
1224   __syncthreads ();
1225
1226   if (gid >= gid_max) return;
1227
1228   /**
1229    * main
1230    */
1231
1232   m07600m (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);
1233 }
1234
1235 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1236 {
1237   /**
1238    * base
1239    */
1240
1241   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1242
1243   /**
1244    * modifier
1245    */
1246
1247   const u32 lid = threadIdx.x;
1248
1249   u32x w0[4];
1250
1251   w0[0] = pws[gid].i[ 0];
1252   w0[1] = pws[gid].i[ 1];
1253   w0[2] = pws[gid].i[ 2];
1254   w0[3] = pws[gid].i[ 3];
1255
1256   u32x w1[4];
1257
1258   w1[0] = 0;
1259   w1[1] = 0;
1260   w1[2] = 0;
1261   w1[3] = 0;
1262
1263   u32x w2[4];
1264
1265   w2[0] = 0;
1266   w2[1] = 0;
1267   w2[2] = 0;
1268   w2[3] = 0;
1269
1270   u32x w3[4];
1271
1272   w3[0] = 0;
1273   w3[1] = 0;
1274   w3[2] = 0;
1275   w3[3] = pws[gid].i[15];
1276
1277   const u32 pw_len = pws[gid].pw_len;
1278
1279   /**
1280    * bin2asc table
1281    */
1282
1283   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1284                  | c_bin2asc[(lid >> 4) & 15] << 0;
1285
1286   __syncthreads ();
1287
1288   if (gid >= gid_max) return;
1289
1290   /**
1291    * main
1292    */
1293
1294   m07600s (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);
1295 }
1296
1297 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1298 {
1299   /**
1300    * base
1301    */
1302
1303   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1304
1305   /**
1306    * modifier
1307    */
1308
1309   const u32 lid = threadIdx.x;
1310
1311   u32x w0[4];
1312
1313   w0[0] = pws[gid].i[ 0];
1314   w0[1] = pws[gid].i[ 1];
1315   w0[2] = pws[gid].i[ 2];
1316   w0[3] = pws[gid].i[ 3];
1317
1318   u32x w1[4];
1319
1320   w1[0] = pws[gid].i[ 4];
1321   w1[1] = pws[gid].i[ 5];
1322   w1[2] = pws[gid].i[ 6];
1323   w1[3] = pws[gid].i[ 7];
1324
1325   u32x w2[4];
1326
1327   w2[0] = 0;
1328   w2[1] = 0;
1329   w2[2] = 0;
1330   w2[3] = 0;
1331
1332   u32x w3[4];
1333
1334   w3[0] = 0;
1335   w3[1] = 0;
1336   w3[2] = 0;
1337   w3[3] = pws[gid].i[15];
1338
1339   const u32 pw_len = pws[gid].pw_len;
1340
1341   /**
1342    * bin2asc table
1343    */
1344
1345   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1346                  | c_bin2asc[(lid >> 4) & 15] << 0;
1347
1348   __syncthreads ();
1349
1350   if (gid >= gid_max) return;
1351
1352   /**
1353    * main
1354    */
1355
1356   m07600s (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);
1357 }
1358
1359 extern "C" __global__ void __launch_bounds__ (256, 1) m07600_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)
1360 {
1361   /**
1362    * base
1363    */
1364
1365   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1366
1367   /**
1368    * modifier
1369    */
1370
1371   const u32 lid = threadIdx.x;
1372
1373   u32x w0[4];
1374
1375   w0[0] = pws[gid].i[ 0];
1376   w0[1] = pws[gid].i[ 1];
1377   w0[2] = pws[gid].i[ 2];
1378   w0[3] = pws[gid].i[ 3];
1379
1380   u32x w1[4];
1381
1382   w1[0] = pws[gid].i[ 4];
1383   w1[1] = pws[gid].i[ 5];
1384   w1[2] = pws[gid].i[ 6];
1385   w1[3] = pws[gid].i[ 7];
1386
1387   u32x w2[4];
1388
1389   w2[0] = pws[gid].i[ 8];
1390   w2[1] = pws[gid].i[ 9];
1391   w2[2] = pws[gid].i[10];
1392   w2[3] = pws[gid].i[11];
1393
1394   u32x w3[4];
1395
1396   w3[0] = pws[gid].i[12];
1397   w3[1] = pws[gid].i[13];
1398   w3[2] = pws[gid].i[14];
1399   w3[3] = pws[gid].i[15];
1400
1401   const u32 pw_len = pws[gid].pw_len;
1402
1403   /**
1404    * bin2asc table
1405    */
1406
1407   l_bin2asc[lid] = c_bin2asc[(lid >> 0) & 15] << 8
1408                  | c_bin2asc[(lid >> 4) & 15] << 0;
1409
1410   __syncthreads ();
1411
1412   if (gid >= gid_max) return;
1413
1414   /**
1415    * main
1416    */
1417
1418   m07600s (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);
1419 }