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