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