nv/m09810_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _OLDOFFICE34_
   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 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  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_SIZE4
  39 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  40 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  41 #endif
  42
  43 typedef struct
  44 {
  45   u8 S[256];
  46
  47   u32 wtf_its_faster;
  48
  49 } RC4_KEY;
  50
  51 __device__ static void swap (RC4_KEY *rc4_key, const u32 i, const u32 j)
  52 {
  53   u8 tmp;
  54
  55   tmp           = rc4_key->S[i];
  56   rc4_key->S[i] = rc4_key->S[j];
  57   rc4_key->S[j] = tmp;
  58 }
  59
  60 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  61 {
  62   u32 v = 0x03020100;
  63   u32 a = 0x04040404;
  64
  65   u32 *ptr = (u32 *) rc4_key->S;
  66
  67   #pragma unroll 64
  68   for (u32 i = 0; i < 64; i++)
  69   {
  70     *ptr++ = v; v += a;
  71   }
  72
  73   u32 j = 0;
  74
  75   for (u32 i = 0; i < 16; i++)
  76   {
  77     u32 idx = i * 16;
  78
  79     u32 v;
  80
  81     v = data[0];
  82
  83     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  84     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  85     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
  86     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
  87
  88     v = data[1];
  89
  90     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  91     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  92     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
  93     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
  94
  95     v = data[2];
  96
  97     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  98     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  99     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
 100     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
 101
 102     v = data[3];
 103
 104     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
 105     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
 106     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
 107     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
 108   }
 109 }
 110
 111 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 112 {
 113   for (u32 k = 0; k < 4; k++)
 114   {
 115     u32 xor4 = 0;
 116
 117     u8 idx;
 118
 119     i += 1;
 120     j += rc4_key->S[i];
 121
 122     swap (rc4_key, i, j);
 123
 124     idx = rc4_key->S[i] + rc4_key->S[j];
 125
 126     xor4 |= rc4_key->S[idx] <<  0;
 127
 128     i += 1;
 129     j += rc4_key->S[i];
 130
 131     swap (rc4_key, i, j);
 132
 133     idx = rc4_key->S[i] + rc4_key->S[j];
 134
 135     xor4 |= rc4_key->S[idx] <<  8;
 136
 137     i += 1;
 138     j += rc4_key->S[i];
 139
 140     swap (rc4_key, i, j);
 141
 142     idx = rc4_key->S[i] + rc4_key->S[j];
 143
 144     xor4 |= rc4_key->S[idx] << 16;
 145
 146     i += 1;
 147     j += rc4_key->S[i];
 148
 149     swap (rc4_key, i, j);
 150
 151     idx = rc4_key->S[i] + rc4_key->S[j];
 152
 153     xor4 |= rc4_key->S[idx] << 24;
 154
 155     out[k] = in[k] ^ xor4;
 156   }
 157
 158   return j;
 159 }
 160
 161 __device__ static void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
 162 {
 163   u32x A = digest[0];
 164   u32x B = digest[1];
 165   u32x C = digest[2];
 166   u32x D = digest[3];
 167   u32x E = digest[4];
 168
 169   u32x w0_t = w0[0];
 170   u32x w1_t = w0[1];
 171   u32x w2_t = w0[2];
 172   u32x w3_t = w0[3];
 173   u32x w4_t = w1[0];
 174   u32x w5_t = w1[1];
 175   u32x w6_t = w1[2];
 176   u32x w7_t = w1[3];
 177   u32x w8_t = w2[0];
 178   u32x w9_t = w2[1];
 179   u32x wa_t = w2[2];
 180   u32x wb_t = w2[3];
 181   u32x wc_t = w3[0];
 182   u32x wd_t = w3[1];
 183   u32x we_t = w3[2];
 184   u32x wf_t = w3[3];
 185
 186   #undef K
 187   #define K SHA1C00
 188
 189   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
 190   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
 191   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
 192   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
 193   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
 194   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
 195   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
 196   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
 197   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
 198   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
 199   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
 200   SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
 201   SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
 202   SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
 203   SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
 204   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
 205   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
 206   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
 207   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
 208   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
 209
 210   #undef K
 211   #define K SHA1C01
 212
 213   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
 214   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
 215   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
 216   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
 217   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
 218   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
 219   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
 220   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
 221   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
 222   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
 223   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
 224   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
 225   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
 226   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
 227   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
 228   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
 229   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
 230   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
 231   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
 232   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
 233
 234   #undef K
 235   #define K SHA1C02
 236
 237   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
 238   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
 239   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
 240   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
 241   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
 242   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
 243   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
 244   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
 245   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
 246   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
 247   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
 248   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
 249   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
 250   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
 251   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
 252   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
 253   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
 254   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
 255   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
 256   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
 257
 258   #undef K
 259   #define K SHA1C03
 260
 261   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
 262   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
 263   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
 264   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
 265   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
 266   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
 267   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
 268   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
 269   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
 270   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
 271   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
 272   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
 273   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
 274   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
 275   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
 276   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
 277   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
 278   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
 279   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
 280   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
 281
 282   digest[0] += A;
 283   digest[1] += B;
 284   digest[2] += C;
 285   digest[3] += D;
 286   digest[4] += E;
 287 }
 288
 289 __device__ __constant__ comb_t c_combs[1024];
 290
 291 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 292 {
 293   /**
 294    * modifier
 295    */
 296
 297   const u32 lid = threadIdx.x;
 298
 299   __shared__ RC4_KEY rc4_keys[64];
 300
 301   RC4_KEY *rc4_key = &rc4_keys[lid];
 302
 303   /**
 304    * base
 305    */
 306
 307   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 308
 309   if (gid >= gid_max) return;
 310
 311   u32x wordl0[4];
 312
 313   wordl0[0] = pws[gid].i[ 0];
 314   wordl0[1] = pws[gid].i[ 1];
 315   wordl0[2] = pws[gid].i[ 2];
 316   wordl0[3] = pws[gid].i[ 3];
 317
 318   u32x wordl1[4];
 319
 320   wordl1[0] = pws[gid].i[ 4];
 321   wordl1[1] = pws[gid].i[ 5];
 322   wordl1[2] = pws[gid].i[ 6];
 323   wordl1[3] = pws[gid].i[ 7];
 324
 325   u32x wordl2[4];
 326
 327   wordl2[0] = 0;
 328   wordl2[1] = 0;
 329   wordl2[2] = 0;
 330   wordl2[3] = 0;
 331
 332   u32x wordl3[4];
 333
 334   wordl3[0] = 0;
 335   wordl3[1] = 0;
 336   wordl3[2] = 0;
 337   wordl3[3] = 0;
 338
 339   const u32 pw_l_len = pws[gid].pw_len;
 340
 341   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 342   {
 343     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 344   }
 345
 346   /**
 347    * esalt
 348    */
 349
 350   const u32 version = oldoffice34_bufs[salt_pos].version;
 351
 352   u32 encryptedVerifier[4];
 353
 354   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 355   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 356   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 357   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 358
 359   /**
 360    * loop
 361    */
 362
 363   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 364   {
 365     u32 wordr0[4];
 366
 367     wordr0[0] = c_combs[il_pos].i[0];
 368     wordr0[1] = c_combs[il_pos].i[1];
 369     wordr0[2] = c_combs[il_pos].i[2];
 370     wordr0[3] = c_combs[il_pos].i[3];
 371
 372     u32 wordr1[4];
 373
 374     wordr1[0] = c_combs[il_pos].i[4];
 375     wordr1[1] = c_combs[il_pos].i[5];
 376     wordr1[2] = c_combs[il_pos].i[6];
 377     wordr1[3] = c_combs[il_pos].i[7];
 378
 379     u32 wordr2[4];
 380
 381     wordr2[0] = 0;
 382     wordr2[1] = 0;
 383     wordr2[2] = 0;
 384     wordr2[3] = 0;
 385
 386     u32 wordr3[4];
 387
 388     wordr3[0] = 0;
 389     wordr3[1] = 0;
 390     wordr3[2] = 0;
 391     wordr3[3] = 0;
 392
 393     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 394     {
 395       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 396     }
 397
 398     u32x w0[4];
 399
 400     u32x key[4];
 401
 402     key[0] =  wordl0[0] | wordr0[0];
 403     key[1] = (wordl0[1] | wordr0[1]) & 0xff;
 404     key[2] = 0;
 405     key[3] = 0;
 406
 407     rc4_init_16 (rc4_key, key);
 408
 409     u32x out[4];
 410
 411     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 412
 413     u32x w0_t[4];
 414     u32x w1_t[4];
 415     u32x w2_t[4];
 416     u32x w3_t[4];
 417
 418     w0_t[0] = swap_workaround (out[0]);
 419     w0_t[1] = swap_workaround (out[1]);
 420     w0_t[2] = swap_workaround (out[2]);
 421     w0_t[3] = swap_workaround (out[3]);
 422     w1_t[0] = 0x80000000;
 423     w1_t[1] = 0;
 424     w1_t[2] = 0;
 425     w1_t[3] = 0;
 426     w2_t[0] = 0;
 427     w2_t[1] = 0;
 428     w2_t[2] = 0;
 429     w2_t[3] = 0;
 430     w3_t[0] = 0;
 431     w3_t[1] = 0;
 432     w3_t[2] = 0;
 433     w3_t[3] = 16 * 8;
 434
 435     u32x digest[5];
 436
 437     digest[0] = SHA1M_A;
 438     digest[1] = SHA1M_B;
 439     digest[2] = SHA1M_C;
 440     digest[3] = SHA1M_D;
 441     digest[4] = SHA1M_E;
 442
 443     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 444
 445     digest[0] = swap_workaround (digest[0]);
 446     digest[1] = swap_workaround (digest[1]);
 447     digest[2] = swap_workaround (digest[2]);
 448     digest[3] = swap_workaround (digest[3]);
 449
 450     rc4_next_16 (rc4_key, 16, j, digest, out);
 451
 452     const u32x r0 = out[0];
 453     const u32x r1 = out[1];
 454     const u32x r2 = out[2];
 455     const u32x r3 = out[3];
 456
 457     #include VECT_COMPARE_M
 458   }
 459 }
 460
 461 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_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)
 462 {
 463 }
 464
 465 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_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)
 466 {
 467 }
 468
 469 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 470 {
 471   /**
 472    * modifier
 473    */
 474
 475   const u32 lid = threadIdx.x;
 476
 477   __shared__ RC4_KEY rc4_keys[64];
 478
 479   RC4_KEY *rc4_key = &rc4_keys[lid];
 480
 481   /**
 482    * base
 483    */
 484
 485   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 486
 487   if (gid >= gid_max) return;
 488
 489   u32x wordl0[4];
 490
 491   wordl0[0] = pws[gid].i[ 0];
 492   wordl0[1] = pws[gid].i[ 1];
 493   wordl0[2] = pws[gid].i[ 2];
 494   wordl0[3] = pws[gid].i[ 3];
 495
 496   u32x wordl1[4];
 497
 498   wordl1[0] = pws[gid].i[ 4];
 499   wordl1[1] = pws[gid].i[ 5];
 500   wordl1[2] = pws[gid].i[ 6];
 501   wordl1[3] = pws[gid].i[ 7];
 502
 503   u32x wordl2[4];
 504
 505   wordl2[0] = 0;
 506   wordl2[1] = 0;
 507   wordl2[2] = 0;
 508   wordl2[3] = 0;
 509
 510   u32x wordl3[4];
 511
 512   wordl3[0] = 0;
 513   wordl3[1] = 0;
 514   wordl3[2] = 0;
 515   wordl3[3] = 0;
 516
 517   const u32 pw_l_len = pws[gid].pw_len;
 518
 519   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 520   {
 521     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 522   }
 523
 524   /**
 525    * digest
 526    */
 527
 528   const u32 search[4] =
 529   {
 530     digests_buf[digests_offset].digest_buf[DGST_R0],
 531     digests_buf[digests_offset].digest_buf[DGST_R1],
 532     digests_buf[digests_offset].digest_buf[DGST_R2],
 533     digests_buf[digests_offset].digest_buf[DGST_R3]
 534   };
 535
 536   /**
 537    * esalt
 538    */
 539
 540   const u32 version = oldoffice34_bufs[salt_pos].version;
 541
 542   u32 encryptedVerifier[4];
 543
 544   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 545   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 546   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 547   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 548
 549   /**
 550    * loop
 551    */
 552
 553   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 554   {
 555     u32 wordr0[4];
 556
 557     wordr0[0] = c_combs[il_pos].i[0];
 558     wordr0[1] = c_combs[il_pos].i[1];
 559     wordr0[2] = c_combs[il_pos].i[2];
 560     wordr0[3] = c_combs[il_pos].i[3];
 561
 562     u32 wordr1[4];
 563
 564     wordr1[0] = c_combs[il_pos].i[4];
 565     wordr1[1] = c_combs[il_pos].i[5];
 566     wordr1[2] = c_combs[il_pos].i[6];
 567     wordr1[3] = c_combs[il_pos].i[7];
 568
 569     u32 wordr2[4];
 570
 571     wordr2[0] = 0;
 572     wordr2[1] = 0;
 573     wordr2[2] = 0;
 574     wordr2[3] = 0;
 575
 576     u32 wordr3[4];
 577
 578     wordr3[0] = 0;
 579     wordr3[1] = 0;
 580     wordr3[2] = 0;
 581     wordr3[3] = 0;
 582
 583     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 584     {
 585       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 586     }
 587
 588     u32x w0[4];
 589
 590     u32x key[4];
 591
 592     key[0] =  wordl0[0] | wordr0[0];
 593     key[1] = (wordl0[1] | wordr0[1]) & 0xff;
 594     key[2] = 0;
 595     key[3] = 0;
 596
 597     rc4_init_16 (rc4_key, key);
 598
 599     u32x out[4];
 600
 601     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 602
 603     u32x w0_t[4];
 604     u32x w1_t[4];
 605     u32x w2_t[4];
 606     u32x w3_t[4];
 607
 608     w0_t[0] = swap_workaround (out[0]);
 609     w0_t[1] = swap_workaround (out[1]);
 610     w0_t[2] = swap_workaround (out[2]);
 611     w0_t[3] = swap_workaround (out[3]);
 612     w1_t[0] = 0x80000000;
 613     w1_t[1] = 0;
 614     w1_t[2] = 0;
 615     w1_t[3] = 0;
 616     w2_t[0] = 0;
 617     w2_t[1] = 0;
 618     w2_t[2] = 0;
 619     w2_t[3] = 0;
 620     w3_t[0] = 0;
 621     w3_t[1] = 0;
 622     w3_t[2] = 0;
 623     w3_t[3] = 16 * 8;
 624
 625     u32x digest[5];
 626
 627     digest[0] = SHA1M_A;
 628     digest[1] = SHA1M_B;
 629     digest[2] = SHA1M_C;
 630     digest[3] = SHA1M_D;
 631     digest[4] = SHA1M_E;
 632
 633     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 634
 635     digest[0] = swap_workaround (digest[0]);
 636     digest[1] = swap_workaround (digest[1]);
 637     digest[2] = swap_workaround (digest[2]);
 638     digest[3] = swap_workaround (digest[3]);
 639
 640     rc4_next_16 (rc4_key, 16, j, digest, out);
 641
 642     const u32x r0 = out[0];
 643     const u32x r1 = out[1];
 644     const u32x r2 = out[2];
 645     const u32x r3 = out[3];
 646
 647     #include VECT_COMPARE_S
 648   }
 649 }
 650
 651 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_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)
 652 {
 653 }
 654
 655 extern "C" __global__ void __launch_bounds__ (64, 1) m09810_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 oldoffice34_t *oldoffice34_bufs, u32 *d_return_buf, u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 656 {
 657 }