nv/m08000_a0.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA256_
   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 7
  21 #define DGST_R2 2
  22 #define DGST_R3 6
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27 #include "include/rp_gpu.h"
  28 #include "rp_nv.c"
  29
  30 #ifdef  VECT_SIZE1
  31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  33 #endif
  34
  35 #ifdef  VECT_SIZE2
  36 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  37 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  38 #endif
  39
  40 __device__ static void sha256_transform (u32x digest[8], const u32x w[16])
  41 {
  42   u32x w0_t = w[ 0];
  43   u32x w1_t = w[ 1];
  44   u32x w2_t = w[ 2];
  45   u32x w3_t = w[ 3];
  46   u32x w4_t = w[ 4];
  47   u32x w5_t = w[ 5];
  48   u32x w6_t = w[ 6];
  49   u32x w7_t = w[ 7];
  50   u32x w8_t = w[ 8];
  51   u32x w9_t = w[ 9];
  52   u32x wa_t = w[10];
  53   u32x wb_t = w[11];
  54   u32x wc_t = w[12];
  55   u32x wd_t = w[13];
  56   u32x we_t = w[14];
  57   u32x wf_t = w[15];
  58
  59   u32x a = digest[0];
  60   u32x b = digest[1];
  61   u32x c = digest[2];
  62   u32x d = digest[3];
  63   u32x e = digest[4];
  64   u32x f = digest[5];
  65   u32x g = digest[6];
  66   u32x h = digest[7];
  67
  68   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, SHA256C00);
  69   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, SHA256C01);
  70   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, SHA256C02);
  71   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, SHA256C03);
  72   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, SHA256C04);
  73   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, SHA256C05);
  74   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, SHA256C06);
  75   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, SHA256C07);
  76   SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, SHA256C08);
  77   SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, SHA256C09);
  78   SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, SHA256C0a);
  79   SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, SHA256C0b);
  80   SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, SHA256C0c);
  81   SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, SHA256C0d);
  82   SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, SHA256C0e);
  83   SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, SHA256C0f);
  84   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);
  85   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);
  86   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);
  87   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);
  88   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);
  89   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);
  90   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);
  91   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);
  92   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);
  93   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);
  94   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);
  95   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);
  96   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);
  97   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);
  98   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);
  99   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);
 100   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);
 101   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);
 102   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);
 103   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);
 104   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);
 105   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);
 106   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);
 107   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);
 108   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);
 109   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);
 110   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);
 111   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);
 112   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);
 113   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);
 114   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);
 115   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);
 116   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);
 117   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);
 118   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);
 119   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);
 120   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);
 121   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);
 122   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);
 123   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);
 124   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);
 125   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);
 126   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);
 127   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);
 128   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);
 129   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);
 130   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);
 131   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);
 132
 133   digest[0] += a;
 134   digest[1] += b;
 135   digest[2] += c;
 136   digest[3] += d;
 137   digest[4] += e;
 138   digest[5] += f;
 139   digest[6] += g;
 140   digest[7] += h;
 141 }
 142
 143 __device__ __constant__ gpu_rule_t c_rules[1024];
 144
 145 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 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)
 146 {
 147   /**
 148    * modifier
 149    */
 150
 151   const u32 lid = threadIdx.x;
 152
 153   /**
 154    * base
 155    */
 156
 157   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 158
 159   if (gid >= gid_max) return;
 160
 161   u32x pw_buf0[4];
 162
 163   pw_buf0[0] = pws[gid].i[ 0];
 164   pw_buf0[1] = pws[gid].i[ 1];
 165   pw_buf0[2] = pws[gid].i[ 2];
 166   pw_buf0[3] = pws[gid].i[ 3];
 167
 168   u32x pw_buf1[4];
 169
 170   pw_buf1[0] = pws[gid].i[ 4];
 171   pw_buf1[1] = pws[gid].i[ 5];
 172   pw_buf1[2] = pws[gid].i[ 6];
 173   pw_buf1[3] = pws[gid].i[ 7];
 174
 175   const u32 pw_len = pws[gid].pw_len;
 176
 177   /**
 178    * salt
 179    */
 180
 181   const u32 salt_buf0 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 0]);
 182   const u32 salt_buf1 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 1]);
 183   const u32 salt_buf2 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 2]); // 0x80
 184
 185   /**
 186    * loop
 187    */
 188
 189   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 190   {
 191     u32x w0[4];
 192     u32x w1[4];
 193     u32x w2[4];
 194     u32x w3[4];
 195
 196     w0[0] = pw_buf0[0];
 197     w0[1] = pw_buf0[1];
 198     w0[2] = pw_buf0[2];
 199     w0[3] = pw_buf0[3];
 200     w1[0] = pw_buf1[0];
 201     w1[1] = pw_buf1[1];
 202     w1[2] = pw_buf1[2];
 203     w1[3] = pw_buf1[3];
 204     w2[0] = 0;
 205     w2[1] = 0;
 206     w2[2] = 0;
 207     w2[3] = 0;
 208     w3[0] = 0;
 209     w3[1] = 0;
 210     w3[2] = 0;
 211     w3[3] = 0;
 212
 213     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 214
 215     u32x w0_t[4];
 216     u32x w1_t[4];
 217     u32x w2_t[4];
 218     u32x w3_t[4];
 219
 220     make_unicode (w0, w0_t, w1_t);
 221
 222     make_unicode (w1, w2_t, w3_t);
 223
 224     u32x w_t[16];
 225
 226     w_t[ 0] = swap_workaround (w0_t[0]);
 227     w_t[ 1] = swap_workaround (w0_t[1]);
 228     w_t[ 2] = swap_workaround (w0_t[2]);
 229     w_t[ 3] = swap_workaround (w0_t[3]);
 230     w_t[ 4] = swap_workaround (w1_t[0]);
 231     w_t[ 5] = swap_workaround (w1_t[1]);
 232     w_t[ 6] = swap_workaround (w1_t[2]);
 233     w_t[ 7] = swap_workaround (w1_t[3]);
 234     w_t[ 8] = swap_workaround (w2_t[0]);
 235     w_t[ 9] = swap_workaround (w2_t[1]);
 236     w_t[10] = swap_workaround (w2_t[2]);
 237     w_t[11] = swap_workaround (w2_t[3]);
 238     w_t[12] = swap_workaround (w3_t[0]);
 239     w_t[13] = swap_workaround (w3_t[1]);
 240     w_t[14] = swap_workaround (w3_t[2]);
 241     w_t[15] = swap_workaround (w3_t[3]);
 242
 243     w_t[ 0] = w_t[ 0] >> 8;
 244     w_t[ 1] = w_t[ 1] >> 8;
 245     w_t[ 2] = w_t[ 2] >> 8;
 246     w_t[ 3] = w_t[ 3] >> 8;
 247     w_t[ 4] = w_t[ 4] >> 8;
 248     w_t[ 5] = w_t[ 5] >> 8;
 249     w_t[ 6] = w_t[ 6] >> 8;
 250     w_t[ 7] = w_t[ 7] >> 8;
 251     w_t[ 8] = w_t[ 8] >> 8;
 252     w_t[ 9] = w_t[ 9] >> 8;
 253     w_t[10] = w_t[10] >> 8;
 254     w_t[11] = w_t[11] >> 8;
 255     w_t[12] = w_t[12] >> 8;
 256     w_t[13] = w_t[13] >> 8;
 257     w_t[14] = w_t[14] >> 8;
 258     w_t[15] = w_t[15] >> 8;
 259
 260     u32x digest[8];
 261
 262     digest[0] = SHA256M_A;
 263     digest[1] = SHA256M_B;
 264     digest[2] = SHA256M_C;
 265     digest[3] = SHA256M_D;
 266     digest[4] = SHA256M_E;
 267     digest[5] = SHA256M_F;
 268     digest[6] = SHA256M_G;
 269     digest[7] = SHA256M_H;
 270
 271     sha256_transform (digest, w_t); //   0 - 64
 272
 273     w_t[ 0] = 0;
 274     w_t[ 1] = 0;
 275     w_t[ 2] = 0;
 276     w_t[ 3] = 0;
 277     w_t[ 4] = 0;
 278     w_t[ 5] = 0;
 279     w_t[ 6] = 0;
 280     w_t[ 7] = 0;
 281     w_t[ 8] = 0;
 282     w_t[ 9] = 0;
 283     w_t[10] = 0;
 284     w_t[11] = 0;
 285     w_t[12] = 0;
 286     w_t[13] = 0;
 287     w_t[14] = 0;
 288     w_t[15] = 0;
 289
 290     sha256_transform (digest, w_t); //  64 - 128
 291     sha256_transform (digest, w_t); // 128 - 192
 292     sha256_transform (digest, w_t); // 192 - 256
 293     sha256_transform (digest, w_t); // 256 - 320
 294     sha256_transform (digest, w_t); // 320 - 384
 295     sha256_transform (digest, w_t); // 384 - 448
 296
 297     w_t[15] =               0 | salt_buf0 >> 16;
 298
 299     sha256_transform (digest, w_t); // 448 - 512
 300
 301     w_t[ 0] = salt_buf0 << 16 | salt_buf1 >> 16;
 302     w_t[ 1] = salt_buf1 << 16 | salt_buf2 >> 16;
 303     w_t[ 2] = salt_buf2 << 16 | 0;
 304     w_t[15] = (510 + 8) * 8;
 305
 306     sha256_transform (digest, w_t); // 512 - 576
 307
 308     const u32x r0 = digest[3];
 309     const u32x r1 = digest[7];
 310     const u32x r2 = digest[2];
 311     const u32x r3 = digest[6];
 312
 313     #include VECT_COMPARE_M
 314   }
 315 }
 316
 317 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 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)
 318 {
 319 }
 320
 321 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 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)
 322 {
 323 }
 324
 325 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 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)
 326 {
 327   /**
 328    * modifier
 329    */
 330
 331   const u32 lid = threadIdx.x;
 332
 333   /**
 334    * base
 335    */
 336
 337   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 338
 339   if (gid >= gid_max) return;
 340
 341   u32x pw_buf0[4];
 342
 343   pw_buf0[0] = pws[gid].i[ 0];
 344   pw_buf0[1] = pws[gid].i[ 1];
 345   pw_buf0[2] = pws[gid].i[ 2];
 346   pw_buf0[3] = pws[gid].i[ 3];
 347
 348   u32x pw_buf1[4];
 349
 350   pw_buf1[0] = pws[gid].i[ 4];
 351   pw_buf1[1] = pws[gid].i[ 5];
 352   pw_buf1[2] = pws[gid].i[ 6];
 353   pw_buf1[3] = pws[gid].i[ 7];
 354
 355   const u32 pw_len = pws[gid].pw_len;
 356
 357   /**
 358    * salt
 359    */
 360
 361   const u32 salt_buf0 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 0]);
 362   const u32 salt_buf1 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 1]);
 363   const u32 salt_buf2 = swap_workaround (salt_bufs[salt_pos].salt_buf[ 2]); // 0x80
 364
 365   /**
 366    * digest
 367    */
 368
 369   const u32 search[4] =
 370   {
 371     digests_buf[digests_offset].digest_buf[DGST_R0],
 372     digests_buf[digests_offset].digest_buf[DGST_R1],
 373     digests_buf[digests_offset].digest_buf[DGST_R2],
 374     digests_buf[digests_offset].digest_buf[DGST_R3]
 375   };
 376
 377   /**
 378    * loop
 379    */
 380
 381   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 382   {
 383     u32x w0[4];
 384     u32x w1[4];
 385     u32x w2[4];
 386     u32x w3[4];
 387
 388     w0[0] = pw_buf0[0];
 389     w0[1] = pw_buf0[1];
 390     w0[2] = pw_buf0[2];
 391     w0[3] = pw_buf0[3];
 392     w1[0] = pw_buf1[0];
 393     w1[1] = pw_buf1[1];
 394     w1[2] = pw_buf1[2];
 395     w1[3] = pw_buf1[3];
 396     w2[0] = 0;
 397     w2[1] = 0;
 398     w2[2] = 0;
 399     w2[3] = 0;
 400     w3[0] = 0;
 401     w3[1] = 0;
 402     w3[2] = 0;
 403     w3[3] = 0;
 404
 405     const u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 406
 407     u32x w0_t[4];
 408     u32x w1_t[4];
 409     u32x w2_t[4];
 410     u32x w3_t[4];
 411
 412     make_unicode (w0, w0_t, w1_t);
 413
 414     make_unicode (w1, w2_t, w3_t);
 415
 416     u32x w_t[16];
 417
 418     w_t[ 0] = swap_workaround (w0_t[0]);
 419     w_t[ 1] = swap_workaround (w0_t[1]);
 420     w_t[ 2] = swap_workaround (w0_t[2]);
 421     w_t[ 3] = swap_workaround (w0_t[3]);
 422     w_t[ 4] = swap_workaround (w1_t[0]);
 423     w_t[ 5] = swap_workaround (w1_t[1]);
 424     w_t[ 6] = swap_workaround (w1_t[2]);
 425     w_t[ 7] = swap_workaround (w1_t[3]);
 426     w_t[ 8] = swap_workaround (w2_t[0]);
 427     w_t[ 9] = swap_workaround (w2_t[1]);
 428     w_t[10] = swap_workaround (w2_t[2]);
 429     w_t[11] = swap_workaround (w2_t[3]);
 430     w_t[12] = swap_workaround (w3_t[0]);
 431     w_t[13] = swap_workaround (w3_t[1]);
 432     w_t[14] = swap_workaround (w3_t[2]);
 433     w_t[15] = swap_workaround (w3_t[3]);
 434
 435     w_t[ 0] = w_t[ 0] >> 8;
 436     w_t[ 1] = w_t[ 1] >> 8;
 437     w_t[ 2] = w_t[ 2] >> 8;
 438     w_t[ 3] = w_t[ 3] >> 8;
 439     w_t[ 4] = w_t[ 4] >> 8;
 440     w_t[ 5] = w_t[ 5] >> 8;
 441     w_t[ 6] = w_t[ 6] >> 8;
 442     w_t[ 7] = w_t[ 7] >> 8;
 443     w_t[ 8] = w_t[ 8] >> 8;
 444     w_t[ 9] = w_t[ 9] >> 8;
 445     w_t[10] = w_t[10] >> 8;
 446     w_t[11] = w_t[11] >> 8;
 447     w_t[12] = w_t[12] >> 8;
 448     w_t[13] = w_t[13] >> 8;
 449     w_t[14] = w_t[14] >> 8;
 450     w_t[15] = w_t[15] >> 8;
 451
 452     u32x digest[8];
 453
 454     digest[0] = SHA256M_A;
 455     digest[1] = SHA256M_B;
 456     digest[2] = SHA256M_C;
 457     digest[3] = SHA256M_D;
 458     digest[4] = SHA256M_E;
 459     digest[5] = SHA256M_F;
 460     digest[6] = SHA256M_G;
 461     digest[7] = SHA256M_H;
 462
 463     sha256_transform (digest, w_t); //   0 - 64
 464
 465     w_t[ 0] = 0;
 466     w_t[ 1] = 0;
 467     w_t[ 2] = 0;
 468     w_t[ 3] = 0;
 469     w_t[ 4] = 0;
 470     w_t[ 5] = 0;
 471     w_t[ 6] = 0;
 472     w_t[ 7] = 0;
 473     w_t[ 8] = 0;
 474     w_t[ 9] = 0;
 475     w_t[10] = 0;
 476     w_t[11] = 0;
 477     w_t[12] = 0;
 478     w_t[13] = 0;
 479     w_t[14] = 0;
 480     w_t[15] = 0;
 481
 482     sha256_transform (digest, w_t); //  64 - 128
 483     sha256_transform (digest, w_t); // 128 - 192
 484     sha256_transform (digest, w_t); // 192 - 256
 485     sha256_transform (digest, w_t); // 256 - 320
 486     sha256_transform (digest, w_t); // 320 - 384
 487     sha256_transform (digest, w_t); // 384 - 448
 488
 489     w_t[15] =               0 | salt_buf0 >> 16;
 490
 491     sha256_transform (digest, w_t); // 448 - 512
 492
 493     w_t[ 0] = salt_buf0 << 16 | salt_buf1 >> 16;
 494     w_t[ 1] = salt_buf1 << 16 | salt_buf2 >> 16;
 495     w_t[ 2] = salt_buf2 << 16 | 0;
 496     w_t[15] = (510 + 8) * 8;
 497
 498     sha256_transform (digest, w_t); // 512 - 576
 499
 500     const u32x r0 = digest[3];
 501     const u32x r1 = digest[7];
 502     const u32x r2 = digest[2];
 503     const u32x r3 = digest[6];
 504
 505     #include VECT_COMPARE_S
 506   }
 507 }
 508
 509 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 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)
 510 {
 511 }
 512
 513 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 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)
 514 {
 515 }