nv/m08400_a3.cu

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