nv/m08000_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA256_
   7 #define _SCALAR_
   8
   9 #include "include/constants.h"
  10 #include "include/kernel_vendor.h"
  11
  12 #ifdef  VLIW1
  13 #define VECT_SIZE1
  14 #endif
  15
  16 #ifdef  VLIW2
  17 #define VECT_SIZE1
  18 #endif
  19
  20 #define DGST_R0 3
  21 #define DGST_R1 7
  22 #define DGST_R2 2
  23 #define DGST_R3 6
  24
  25 #include "include/kernel_functions.c"
  26 #include "types_nv.c"
  27 #include "common_nv.c"
  28
  29 #ifdef  VECT_SIZE1
  30 #define VECT_COMPARE_S "check_single_vect1_comp4_warp.c"
  31 #define VECT_COMPARE_M "check_multi_vect1_comp4_warp.c"
  32 #endif
  33
  34 #ifdef  VECT_SIZE2
  35 #define VECT_COMPARE_S "check_single_vect2_comp4_warp.c"
  36 #define VECT_COMPARE_M "check_multi_vect2_comp4_warp.c"
  37 #endif
  38
  39 __device__ static void sha256_transform (u32x digest[8], const u32x w[16])
  40 {
  41   u32x w0_t = w[ 0];
  42   u32x w1_t = w[ 1];
  43   u32x w2_t = w[ 2];
  44   u32x w3_t = w[ 3];
  45   u32x w4_t = w[ 4];
  46   u32x w5_t = w[ 5];
  47   u32x w6_t = w[ 6];
  48   u32x w7_t = w[ 7];
  49   u32x w8_t = w[ 8];
  50   u32x w9_t = w[ 9];
  51   u32x wa_t = w[10];
  52   u32x wb_t = w[11];
  53   u32x wc_t = w[12];
  54   u32x wd_t = w[13];
  55   u32x we_t = w[14];
  56   u32x wf_t = w[15];
  57
  58   u32x a = digest[0];
  59   u32x b = digest[1];
  60   u32x c = digest[2];
  61   u32x d = digest[3];
  62   u32x e = digest[4];
  63   u32x f = digest[5];
  64   u32x g = digest[6];
  65   u32x h = digest[7];
  66
  67   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
  68   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
  69   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
  70   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
  71   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
  72   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
  73   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
  74   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
  75   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
  76   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
  77   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
  78   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
  79   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
  80   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
  81   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
  82   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
  83   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C10);
  84   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C11);
  85   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C12);
  86   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C13);
  87   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C14);
  88   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C15);
  89   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C16);
  90   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C17);
  91   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C18);
  92   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C19);
  93   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C1a);
  94   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C1b);
  95   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C1c);
  96   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C1d);
  97   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C1e);
  98   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C1f);
  99   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C20);
 100   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C21);
 101   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C22);
 102   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C23);
 103   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C24);
 104   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C25);
 105   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C26);
 106   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C27);
 107   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C28);
 108   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C29);
 109   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C2a);
 110   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C2b);
 111   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C2c);
 112   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C2d);
 113   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C2e);
 114   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C2f);
 115   w0_t = SHA256_S1(we_t) + w9_t + SHA256_S0(w1_t) + w0_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C30);
 116   w1_t = SHA256_S1(wf_t) + wa_t + SHA256_S0(w2_t) + w1_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C31);
 117   w2_t = SHA256_S1(w0_t) + wb_t + SHA256_S0(w3_t) + w2_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C32);
 118   w3_t = SHA256_S1(w1_t) + wc_t + SHA256_S0(w4_t) + w3_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C33);
 119   w4_t = SHA256_S1(w2_t) + wd_t + SHA256_S0(w5_t) + w4_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C34);
 120   w5_t = SHA256_S1(w3_t) + we_t + SHA256_S0(w6_t) + w5_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C35);
 121   w6_t = SHA256_S1(w4_t) + wf_t + SHA256_S0(w7_t) + w6_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C36);
 122   w7_t = SHA256_S1(w5_t) + w0_t + SHA256_S0(w8_t) + w7_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C37);
 123   w8_t = SHA256_S1(w6_t) + w1_t + SHA256_S0(w9_t) + w8_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C38);
 124   w9_t = SHA256_S1(w7_t) + w2_t + SHA256_S0(wa_t) + w9_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C39);
 125   wa_t = SHA256_S1(w8_t) + w3_t + SHA256_S0(wb_t) + wa_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C3a);
 126   wb_t = SHA256_S1(w9_t) + w4_t + SHA256_S0(wc_t) + wb_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C3b);
 127   wc_t = SHA256_S1(wa_t) + w5_t + SHA256_S0(wd_t) + wc_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C3c);
 128   wd_t = SHA256_S1(wb_t) + w6_t + SHA256_S0(we_t) + wd_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C3d);
 129   we_t = SHA256_S1(wc_t) + w7_t + SHA256_S0(wf_t) + we_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C3e);
 130   wf_t = SHA256_S1(wd_t) + w8_t + SHA256_S0(w0_t) + wf_t; SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C3f);
 131
 132   digest[0] += a;
 133   digest[1] += b;
 134   digest[2] += c;
 135   digest[3] += d;
 136   digest[4] += e;
 137   digest[5] += f;
 138   digest[6] += g;
 139   digest[7] += h;
 140 }
 141
 142 __device__ __constant__ u32x c_bfs[1024];
 143
 144 __device__ static void m08000m (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 145 {
 146   /**
 147    * modifier
 148    */
 149
 150   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 151   const u32 lid = threadIdx.x;
 152
 153   /**
 154    * salt
 155    */
 156
 157   const u32 salt_buf0 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 0]);
 158   const u32 salt_buf1 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 1]);
 159   const u32 salt_buf2 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 2]); // 0x80
 160
 161   /**
 162    * loop
 163    */
 164
 165   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 166
 167   u32x w0l = w[0];
 168
 169   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 170   {
 171     const u32x w0r = c_bfs[il_pos];
 172
 173     const u32x w0 = w0l | w0r;
 174
 175     u32x w_t[16];
 176
 177     w_t[ 0] = w0    >> 8;
 178     w_t[ 1] = w[ 1] >> 8;
 179     w_t[ 2] = w[ 2] >> 8;
 180     w_t[ 3] = w[ 3] >> 8;
 181     w_t[ 4] = w[ 4] >> 8;
 182     w_t[ 5] = w[ 5] >> 8;
 183     w_t[ 6] = w[ 6] >> 8;
 184     w_t[ 7] = w[ 7] >> 8;
 185     w_t[ 8] = w[ 8] >> 8;
 186     w_t[ 9] = w[ 9] >> 8;
 187     w_t[10] = w[10] >> 8;
 188     w_t[11] = w[11] >> 8;
 189     w_t[12] = w[12] >> 8;
 190     w_t[13] = w[13] >> 8;
 191     w_t[14] = w[14] >> 8;
 192     w_t[15] = w[15] >> 8;
 193
 194     u32x digest[8];
 195
 196     digest[0] = SHA256M_A;
 197     digest[1] = SHA256M_B;
 198     digest[2] = SHA256M_C;
 199     digest[3] = SHA256M_D;
 200     digest[4] = SHA256M_E;
 201     digest[5] = SHA256M_F;
 202     digest[6] = SHA256M_G;
 203     digest[7] = SHA256M_H;
 204
 205     sha256_transform (digest, w_t); //   0 - 64
 206
 207     w_t[ 0] = 0;
 208     w_t[ 1] = 0;
 209     w_t[ 2] = 0;
 210     w_t[ 3] = 0;
 211     w_t[ 4] = 0;
 212     w_t[ 5] = 0;
 213     w_t[ 6] = 0;
 214     w_t[ 7] = 0;
 215     w_t[ 8] = 0;
 216     w_t[ 9] = 0;
 217     w_t[10] = 0;
 218     w_t[11] = 0;
 219     w_t[12] = 0;
 220     w_t[13] = 0;
 221     w_t[14] = 0;
 222     w_t[15] = 0;
 223
 224     sha256_transform (digest, w_t); //  64 - 128
 225     sha256_transform (digest, w_t); // 128 - 192
 226     sha256_transform (digest, w_t); // 192 - 256
 227     sha256_transform (digest, w_t); // 256 - 320
 228     sha256_transform (digest, w_t); // 320 - 384
 229     sha256_transform (digest, w_t); // 384 - 448
 230
 231     w_t[15] =               0 | salt_buf0 >> 16;
 232
 233     sha256_transform (digest, w_t); // 448 - 512
 234
 235     w_t[ 0] = salt_buf0 << 16 | salt_buf1 >> 16;
 236     w_t[ 1] = salt_buf1 << 16 | salt_buf2 >> 16;
 237     w_t[ 2] = salt_buf2 << 16 | 0;
 238     w_t[15] = (510 + 8) * 8;
 239
 240     sha256_transform (digest, w_t); // 512 - 576
 241
 242     const u32x r0 = digest[3];
 243     const u32x r1 = digest[7];
 244     const u32x r2 = digest[2];
 245     const u32x r3 = digest[6];
 246
 247     #include VECT_COMPARE_M
 248   }
 249 }
 250
 251 __device__ static void m08000s (u32 w[16], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 252 {
 253   /**
 254    * modifier
 255    */
 256
 257   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 258   const u32 lid = threadIdx.x;
 259
 260   /**
 261    * digest
 262    */
 263
 264   const u32 search[4] =
 265   {
 266     digests_buf[digests_offset].digest_buf[DGST_R0],
 267     digests_buf[digests_offset].digest_buf[DGST_R1],
 268     digests_buf[digests_offset].digest_buf[DGST_R2],
 269     digests_buf[digests_offset].digest_buf[DGST_R3]
 270   };
 271
 272   /**
 273    * salt
 274    */
 275
 276   const u32 salt_buf0 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 0]);
 277   const u32 salt_buf1 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 1]);
 278   const u32 salt_buf2 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 2]); // 0x80
 279
 280   /**
 281    * loop
 282    */
 283
 284   const u32 bf_loops = ceil ((float) bfs_cnt / VECT_DIV);
 285
 286   u32x w0l = w[0];
 287
 288   for (u32 il_pos = 0; il_pos < bf_loops; il_pos++)
 289   {
 290     const u32x w0r = c_bfs[il_pos];
 291
 292     const u32x w0 = w0l | w0r;
 293
 294     u32x w_t[16];
 295
 296     w_t[ 0] = w0    >> 8;
 297     w_t[ 1] = w[ 1] >> 8;
 298     w_t[ 2] = w[ 2] >> 8;
 299     w_t[ 3] = w[ 3] >> 8;
 300     w_t[ 4] = w[ 4] >> 8;
 301     w_t[ 5] = w[ 5] >> 8;
 302     w_t[ 6] = w[ 6] >> 8;
 303     w_t[ 7] = w[ 7] >> 8;
 304     w_t[ 8] = w[ 8] >> 8;
 305     w_t[ 9] = w[ 9] >> 8;
 306     w_t[10] = w[10] >> 8;
 307     w_t[11] = w[11] >> 8;
 308     w_t[12] = w[12] >> 8;
 309     w_t[13] = w[13] >> 8;
 310     w_t[14] = w[14] >> 8;
 311     w_t[15] = w[15] >> 8;
 312
 313     u32x digest[8];
 314
 315     digest[0] = SHA256M_A;
 316     digest[1] = SHA256M_B;
 317     digest[2] = SHA256M_C;
 318     digest[3] = SHA256M_D;
 319     digest[4] = SHA256M_E;
 320     digest[5] = SHA256M_F;
 321     digest[6] = SHA256M_G;
 322     digest[7] = SHA256M_H;
 323
 324     sha256_transform (digest, w_t); //   0 - 64
 325
 326     w_t[ 0] = 0;
 327     w_t[ 1] = 0;
 328     w_t[ 2] = 0;
 329     w_t[ 3] = 0;
 330     w_t[ 4] = 0;
 331     w_t[ 5] = 0;
 332     w_t[ 6] = 0;
 333     w_t[ 7] = 0;
 334     w_t[ 8] = 0;
 335     w_t[ 9] = 0;
 336     w_t[10] = 0;
 337     w_t[11] = 0;
 338     w_t[12] = 0;
 339     w_t[13] = 0;
 340     w_t[14] = 0;
 341     w_t[15] = 0;
 342
 343     sha256_transform (digest, w_t); //  64 - 128
 344     sha256_transform (digest, w_t); // 128 - 192
 345     sha256_transform (digest, w_t); // 192 - 256
 346     sha256_transform (digest, w_t); // 256 - 320
 347     sha256_transform (digest, w_t); // 320 - 384
 348     sha256_transform (digest, w_t); // 384 - 448
 349
 350     w_t[15] =               0 | salt_buf0 >> 16;
 351
 352     sha256_transform (digest, w_t); // 448 - 512
 353
 354     w_t[ 0] = salt_buf0 << 16 | salt_buf1 >> 16;
 355     w_t[ 1] = salt_buf1 << 16 | salt_buf2 >> 16;
 356     w_t[ 2] = salt_buf2 << 16 | 0;
 357     w_t[15] = (510 + 8) * 8;
 358
 359     sha256_transform (digest, w_t); // 512 - 576
 360
 361     const u32x r0 = digest[3];
 362     const u32x r1 = digest[7];
 363     const u32x r2 = digest[2];
 364     const u32x r3 = digest[6];
 365
 366     #include VECT_COMPARE_S
 367   }
 368 }
 369
 370 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_m04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 371 {
 372   /**
 373    * base
 374    */
 375
 376   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 377
 378   if (gid >= gid_max) return;
 379
 380   u32 w[16];
 381
 382   w[ 0] = pws[gid].i[ 0];
 383   w[ 1] = pws[gid].i[ 1];
 384   w[ 2] = pws[gid].i[ 2];
 385   w[ 3] = pws[gid].i[ 3];
 386   w[ 4] = 0;
 387   w[ 5] = 0;
 388   w[ 6] = 0;
 389   w[ 7] = 0;
 390   w[ 8] = 0;
 391   w[ 9] = 0;
 392   w[10] = 0;
 393   w[11] = 0;
 394   w[12] = 0;
 395   w[13] = 0;
 396   w[14] = 0;
 397   w[15] = 0;
 398
 399   const u32 pw_len = pws[gid].pw_len;
 400
 401   /**
 402    * main
 403    */
 404
 405   m08000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 406 }
 407
 408 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_m08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 409 {
 410   /**
 411    * base
 412    */
 413
 414   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 415
 416   if (gid >= gid_max) return;
 417
 418   u32 w[16];
 419
 420   w[ 0] = pws[gid].i[ 0];
 421   w[ 1] = pws[gid].i[ 1];
 422   w[ 2] = pws[gid].i[ 2];
 423   w[ 3] = pws[gid].i[ 3];
 424   w[ 4] = pws[gid].i[ 4];
 425   w[ 5] = pws[gid].i[ 5];
 426   w[ 6] = pws[gid].i[ 6];
 427   w[ 7] = pws[gid].i[ 7];
 428   w[ 8] = 0;
 429   w[ 9] = 0;
 430   w[10] = 0;
 431   w[11] = 0;
 432   w[12] = 0;
 433   w[13] = 0;
 434   w[14] = 0;
 435   w[15] = 0;
 436
 437   const u32 pw_len = pws[gid].pw_len;
 438
 439   /**
 440    * main
 441    */
 442
 443   m08000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 444 }
 445
 446 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_m16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 447 {
 448   /**
 449    * base
 450    */
 451
 452   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 453
 454   if (gid >= gid_max) return;
 455
 456   u32 w[16];
 457
 458   w[ 0] = pws[gid].i[ 0];
 459   w[ 1] = pws[gid].i[ 1];
 460   w[ 2] = pws[gid].i[ 2];
 461   w[ 3] = pws[gid].i[ 3];
 462   w[ 4] = pws[gid].i[ 4];
 463   w[ 5] = pws[gid].i[ 5];
 464   w[ 6] = pws[gid].i[ 6];
 465   w[ 7] = pws[gid].i[ 7];
 466   w[ 8] = pws[gid].i[ 8];
 467   w[ 9] = pws[gid].i[ 9];
 468   w[10] = pws[gid].i[10];
 469   w[11] = pws[gid].i[11];
 470   w[12] = pws[gid].i[12];
 471   w[13] = pws[gid].i[13];
 472   w[14] = pws[gid].i[14];
 473   w[15] = pws[gid].i[15];
 474
 475   const u32 pw_len = pws[gid].pw_len;
 476
 477   /**
 478    * main
 479    */
 480
 481   m08000m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 482 }
 483
 484 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_s04 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 485 {
 486   /**
 487    * base
 488    */
 489
 490   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 491
 492   if (gid >= gid_max) return;
 493
 494   u32 w[16];
 495
 496   w[ 0] = pws[gid].i[ 0];
 497   w[ 1] = pws[gid].i[ 1];
 498   w[ 2] = pws[gid].i[ 2];
 499   w[ 3] = pws[gid].i[ 3];
 500   w[ 4] = 0;
 501   w[ 5] = 0;
 502   w[ 6] = 0;
 503   w[ 7] = 0;
 504   w[ 8] = 0;
 505   w[ 9] = 0;
 506   w[10] = 0;
 507   w[11] = 0;
 508   w[12] = 0;
 509   w[13] = 0;
 510   w[14] = 0;
 511   w[15] = 0;
 512
 513   const u32 pw_len = pws[gid].pw_len;
 514
 515   /**
 516    * main
 517    */
 518
 519   m08000s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 520 }
 521
 522 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_s08 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 523 {
 524   /**
 525    * base
 526    */
 527
 528   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 529
 530   if (gid >= gid_max) return;
 531
 532   u32 w[16];
 533
 534   w[ 0] = pws[gid].i[ 0];
 535   w[ 1] = pws[gid].i[ 1];
 536   w[ 2] = pws[gid].i[ 2];
 537   w[ 3] = pws[gid].i[ 3];
 538   w[ 4] = pws[gid].i[ 4];
 539   w[ 5] = pws[gid].i[ 5];
 540   w[ 6] = pws[gid].i[ 6];
 541   w[ 7] = pws[gid].i[ 7];
 542   w[ 8] = 0;
 543   w[ 9] = 0;
 544   w[10] = 0;
 545   w[11] = 0;
 546   w[12] = 0;
 547   w[13] = 0;
 548   w[14] = 0;
 549   w[15] = 0;
 550
 551   const u32 pw_len = pws[gid].pw_len;
 552
 553   /**
 554    * main
 555    */
 556
 557   m08000s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 558 }
 559
 560 extern "C" __global__ void __launch_bounds__ (256, 1) m08000_s16 (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *words_buf_r, 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)
 561 {
 562   /**
 563    * base
 564    */
 565
 566   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 567
 568   if (gid >= gid_max) return;
 569
 570   u32 w[16];
 571
 572   w[ 0] = pws[gid].i[ 0];
 573   w[ 1] = pws[gid].i[ 1];
 574   w[ 2] = pws[gid].i[ 2];
 575   w[ 3] = pws[gid].i[ 3];
 576   w[ 4] = pws[gid].i[ 4];
 577   w[ 5] = pws[gid].i[ 5];
 578   w[ 6] = pws[gid].i[ 6];
 579   w[ 7] = pws[gid].i[ 7];
 580   w[ 8] = pws[gid].i[ 8];
 581   w[ 9] = pws[gid].i[ 9];
 582   w[10] = pws[gid].i[10];
 583   w[11] = pws[gid].i[11];
 584   w[12] = pws[gid].i[12];
 585   w[13] = pws[gid].i[13];
 586   w[14] = pws[gid].i[14];
 587   w[15] = pws[gid].i[15];
 588
 589   const u32 pw_len = pws[gid].pw_len;
 590
 591   /**
 592    * main
 593    */
 594
 595   m08000s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, 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);
 596 }