nv/m09800_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) m09800_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    * salt
 348    */
 349
 350   u32 salt_buf[4];
 351
 352   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 353   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 354   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 355   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 356
 357   const u32 salt_len = 16;
 358
 359   /**
 360    * esalt
 361    */
 362
 363   const u32 version = oldoffice34_bufs[salt_pos].version;
 364
 365   u32 encryptedVerifier[4];
 366
 367   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 368   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 369   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 370   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 371
 372   /**
 373    * loop
 374    */
 375
 376   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 377   {
 378     const u32 pw_r_len = c_combs[il_pos].pw_len;
 379
 380     const u32 pw_len = pw_l_len + pw_r_len;
 381
 382     const u32 pw_salt_len = (pw_len * 2) + salt_len;
 383
 384     u32 wordr0[4];
 385
 386     wordr0[0] = c_combs[il_pos].i[0];
 387     wordr0[1] = c_combs[il_pos].i[1];
 388     wordr0[2] = c_combs[il_pos].i[2];
 389     wordr0[3] = c_combs[il_pos].i[3];
 390
 391     u32 wordr1[4];
 392
 393     wordr1[0] = c_combs[il_pos].i[4];
 394     wordr1[1] = c_combs[il_pos].i[5];
 395     wordr1[2] = c_combs[il_pos].i[6];
 396     wordr1[3] = c_combs[il_pos].i[7];
 397
 398     u32 wordr2[4];
 399
 400     wordr2[0] = 0;
 401     wordr2[1] = 0;
 402     wordr2[2] = 0;
 403     wordr2[3] = 0;
 404
 405     u32 wordr3[4];
 406
 407     wordr3[0] = 0;
 408     wordr3[1] = 0;
 409     wordr3[2] = 0;
 410     wordr3[3] = 0;
 411
 412     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 413     {
 414       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 415     }
 416
 417     u32x w0[4];
 418
 419     w0[0] = wordl0[0] | wordr0[0];
 420     w0[1] = wordl0[1] | wordr0[1];
 421     w0[2] = wordl0[2] | wordr0[2];
 422     w0[3] = wordl0[3] | wordr0[3];
 423
 424     u32x w1[4];
 425
 426     w1[0] = wordl1[0] | wordr1[0];
 427     w1[1] = wordl1[1] | wordr1[1];
 428     w1[2] = wordl1[2] | wordr1[2];
 429     w1[3] = wordl1[3] | wordr1[3];
 430
 431     u32x w2[4];
 432
 433     w2[0] = wordl2[0] | wordr2[0];
 434     w2[1] = wordl2[1] | wordr2[1];
 435     w2[2] = wordl2[2] | wordr2[2];
 436     w2[3] = wordl2[3] | wordr2[3];
 437
 438     u32x w3[4];
 439
 440     w3[0] = wordl3[0] | wordr3[0];
 441     w3[1] = wordl3[1] | wordr3[1];
 442     w3[2] = 0;
 443     w3[3] = 0;
 444
 445     append_0x80_2 (w0, w1, pw_len);
 446
 447     u32x w0_t[4];
 448     u32x w1_t[4];
 449     u32x w2_t[4];
 450     u32x w3_t[4];
 451
 452     make_unicode (w0, w0_t, w1_t);
 453     make_unicode (w1, w2_t, w3_t);
 454
 455     switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
 456
 457     w0_t[0] = salt_buf[0];
 458     w0_t[1] = salt_buf[1];
 459     w0_t[2] = salt_buf[2];
 460     w0_t[3] = salt_buf[3];
 461     w1_t[0] = swap_workaround (w1_t[0]);
 462     w1_t[1] = swap_workaround (w1_t[1]);
 463     w1_t[2] = swap_workaround (w1_t[2]);
 464     w1_t[3] = swap_workaround (w1_t[3]);
 465     w2_t[0] = swap_workaround (w2_t[0]);
 466     w2_t[1] = swap_workaround (w2_t[1]);
 467     w2_t[2] = swap_workaround (w2_t[2]);
 468     w2_t[3] = swap_workaround (w2_t[3]);
 469     w3_t[0] = swap_workaround (w3_t[0]);
 470     w3_t[1] = swap_workaround (w3_t[1]);
 471     w3_t[2] = 0;
 472     w3_t[3] = pw_salt_len * 8;
 473
 474     u32x digest[5];
 475
 476     digest[0] = SHA1M_A;
 477     digest[1] = SHA1M_B;
 478     digest[2] = SHA1M_C;
 479     digest[3] = SHA1M_D;
 480     digest[4] = SHA1M_E;
 481
 482     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 483
 484     w0_t[0] = digest[0];
 485     w0_t[1] = digest[1];
 486     w0_t[2] = digest[2];
 487     w0_t[3] = digest[3];
 488     w1_t[0] = digest[4];
 489     w1_t[1] = 0;
 490     w1_t[2] = 0x80000000;
 491     w1_t[3] = 0;
 492     w2_t[0] = 0;
 493     w2_t[1] = 0;
 494     w2_t[2] = 0;
 495     w2_t[3] = 0;
 496     w3_t[0] = 0;
 497     w3_t[1] = 0;
 498     w3_t[2] = 0;
 499     w3_t[3] = (20 + 4) * 8;
 500
 501     digest[0] = SHA1M_A;
 502     digest[1] = SHA1M_B;
 503     digest[2] = SHA1M_C;
 504     digest[3] = SHA1M_D;
 505     digest[4] = SHA1M_E;
 506
 507     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 508
 509     u32x key[4];
 510
 511     key[0] = swap_workaround (digest[0]);
 512     key[1] = swap_workaround (digest[1]);
 513     key[2] = swap_workaround (digest[2]);
 514     key[3] = swap_workaround (digest[3]);
 515
 516     if (version == 3)
 517     {
 518       key[1] &= 0xff;
 519       key[2]  = 0;
 520       key[3]  = 0;
 521     }
 522
 523     rc4_init_16 (rc4_key, key);
 524
 525     u32x out[4];
 526
 527     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 528
 529     w0_t[0] = swap_workaround (out[0]);
 530     w0_t[1] = swap_workaround (out[1]);
 531     w0_t[2] = swap_workaround (out[2]);
 532     w0_t[3] = swap_workaround (out[3]);
 533     w1_t[0] = 0x80000000;
 534     w1_t[1] = 0;
 535     w1_t[2] = 0;
 536     w1_t[3] = 0;
 537     w2_t[0] = 0;
 538     w2_t[1] = 0;
 539     w2_t[2] = 0;
 540     w2_t[3] = 0;
 541     w3_t[0] = 0;
 542     w3_t[1] = 0;
 543     w3_t[2] = 0;
 544     w3_t[3] = 16 * 8;
 545
 546     digest[0] = SHA1M_A;
 547     digest[1] = SHA1M_B;
 548     digest[2] = SHA1M_C;
 549     digest[3] = SHA1M_D;
 550     digest[4] = SHA1M_E;
 551
 552     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 553
 554     digest[0] = swap_workaround (digest[0]);
 555     digest[1] = swap_workaround (digest[1]);
 556     digest[2] = swap_workaround (digest[2]);
 557     digest[3] = swap_workaround (digest[3]);
 558
 559     rc4_next_16 (rc4_key, 16, j, digest, out);
 560
 561     const u32x r0 = out[0];
 562     const u32x r1 = out[1];
 563     const u32x r2 = out[2];
 564     const u32x r3 = out[3];
 565
 566     #include VECT_COMPARE_M
 567   }
 568 }
 569
 570 extern "C" __global__ void __launch_bounds__ (64, 1) m09800_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)
 571 {
 572 }
 573
 574 extern "C" __global__ void __launch_bounds__ (64, 1) m09800_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)
 575 {
 576 }
 577
 578 extern "C" __global__ void __launch_bounds__ (64, 1) m09800_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)
 579 {
 580   /**
 581    * modifier
 582    */
 583
 584   const u32 lid = threadIdx.x;
 585
 586   __shared__ RC4_KEY rc4_keys[64];
 587
 588   RC4_KEY *rc4_key = &rc4_keys[lid];
 589
 590   /**
 591    * base
 592    */
 593
 594   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 595
 596   if (gid >= gid_max) return;
 597
 598   u32x wordl0[4];
 599
 600   wordl0[0] = pws[gid].i[ 0];
 601   wordl0[1] = pws[gid].i[ 1];
 602   wordl0[2] = pws[gid].i[ 2];
 603   wordl0[3] = pws[gid].i[ 3];
 604
 605   u32x wordl1[4];
 606
 607   wordl1[0] = pws[gid].i[ 4];
 608   wordl1[1] = pws[gid].i[ 5];
 609   wordl1[2] = pws[gid].i[ 6];
 610   wordl1[3] = pws[gid].i[ 7];
 611
 612   u32x wordl2[4];
 613
 614   wordl2[0] = 0;
 615   wordl2[1] = 0;
 616   wordl2[2] = 0;
 617   wordl2[3] = 0;
 618
 619   u32x wordl3[4];
 620
 621   wordl3[0] = 0;
 622   wordl3[1] = 0;
 623   wordl3[2] = 0;
 624   wordl3[3] = 0;
 625
 626   const u32 pw_l_len = pws[gid].pw_len;
 627
 628   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 629   {
 630     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 631   }
 632
 633   /**
 634    * digest
 635    */
 636
 637   const u32 search[4] =
 638   {
 639     digests_buf[digests_offset].digest_buf[DGST_R0],
 640     digests_buf[digests_offset].digest_buf[DGST_R1],
 641     digests_buf[digests_offset].digest_buf[DGST_R2],
 642     digests_buf[digests_offset].digest_buf[DGST_R3]
 643   };
 644
 645   /**
 646    * salt
 647    */
 648
 649   u32 salt_buf[4];
 650
 651   salt_buf[0] = salt_bufs[salt_pos].salt_buf[0];
 652   salt_buf[1] = salt_bufs[salt_pos].salt_buf[1];
 653   salt_buf[2] = salt_bufs[salt_pos].salt_buf[2];
 654   salt_buf[3] = salt_bufs[salt_pos].salt_buf[3];
 655
 656   const u32 salt_len = 16;
 657
 658   /**
 659    * esalt
 660    */
 661
 662   const u32 version = oldoffice34_bufs[salt_pos].version;
 663
 664   u32 encryptedVerifier[4];
 665
 666   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 667   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 668   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 669   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 670
 671   /**
 672    * loop
 673    */
 674
 675   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 676   {
 677     const u32 pw_r_len = c_combs[il_pos].pw_len;
 678
 679     const u32 pw_len = pw_l_len + pw_r_len;
 680
 681     const u32 pw_salt_len = (pw_len * 2) + salt_len;
 682
 683     u32 wordr0[4];
 684
 685     wordr0[0] = c_combs[il_pos].i[0];
 686     wordr0[1] = c_combs[il_pos].i[1];
 687     wordr0[2] = c_combs[il_pos].i[2];
 688     wordr0[3] = c_combs[il_pos].i[3];
 689
 690     u32 wordr1[4];
 691
 692     wordr1[0] = c_combs[il_pos].i[4];
 693     wordr1[1] = c_combs[il_pos].i[5];
 694     wordr1[2] = c_combs[il_pos].i[6];
 695     wordr1[3] = c_combs[il_pos].i[7];
 696
 697     u32 wordr2[4];
 698
 699     wordr2[0] = 0;
 700     wordr2[1] = 0;
 701     wordr2[2] = 0;
 702     wordr2[3] = 0;
 703
 704     u32 wordr3[4];
 705
 706     wordr3[0] = 0;
 707     wordr3[1] = 0;
 708     wordr3[2] = 0;
 709     wordr3[3] = 0;
 710
 711     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 712     {
 713       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 714     }
 715
 716     u32x w0[4];
 717
 718     w0[0] = wordl0[0] | wordr0[0];
 719     w0[1] = wordl0[1] | wordr0[1];
 720     w0[2] = wordl0[2] | wordr0[2];
 721     w0[3] = wordl0[3] | wordr0[3];
 722
 723     u32x w1[4];
 724
 725     w1[0] = wordl1[0] | wordr1[0];
 726     w1[1] = wordl1[1] | wordr1[1];
 727     w1[2] = wordl1[2] | wordr1[2];
 728     w1[3] = wordl1[3] | wordr1[3];
 729
 730     u32x w2[4];
 731
 732     w2[0] = wordl2[0] | wordr2[0];
 733     w2[1] = wordl2[1] | wordr2[1];
 734     w2[2] = wordl2[2] | wordr2[2];
 735     w2[3] = wordl2[3] | wordr2[3];
 736
 737     u32x w3[4];
 738
 739     w3[0] = wordl3[0] | wordr3[0];
 740     w3[1] = wordl3[1] | wordr3[1];
 741     w3[2] = 0;
 742     w3[3] = 0;
 743
 744     append_0x80_2 (w0, w1, pw_len);
 745
 746     u32x w0_t[4];
 747     u32x w1_t[4];
 748     u32x w2_t[4];
 749     u32x w3_t[4];
 750
 751     make_unicode (w0, w0_t, w1_t);
 752     make_unicode (w1, w2_t, w3_t);
 753
 754     switch_buffer_by_offset (w0_t, w1_t, w2_t, w3_t, salt_len);
 755
 756     w0_t[0] = salt_buf[0];
 757     w0_t[1] = salt_buf[1];
 758     w0_t[2] = salt_buf[2];
 759     w0_t[3] = salt_buf[3];
 760     w1_t[0] = swap_workaround (w1_t[0]);
 761     w1_t[1] = swap_workaround (w1_t[1]);
 762     w1_t[2] = swap_workaround (w1_t[2]);
 763     w1_t[3] = swap_workaround (w1_t[3]);
 764     w2_t[0] = swap_workaround (w2_t[0]);
 765     w2_t[1] = swap_workaround (w2_t[1]);
 766     w2_t[2] = swap_workaround (w2_t[2]);
 767     w2_t[3] = swap_workaround (w2_t[3]);
 768     w3_t[0] = swap_workaround (w3_t[0]);
 769     w3_t[1] = swap_workaround (w3_t[1]);
 770     w3_t[2] = 0;
 771     w3_t[3] = pw_salt_len * 8;
 772
 773     u32x digest[5];
 774
 775     digest[0] = SHA1M_A;
 776     digest[1] = SHA1M_B;
 777     digest[2] = SHA1M_C;
 778     digest[3] = SHA1M_D;
 779     digest[4] = SHA1M_E;
 780
 781     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 782
 783     w0_t[0] = digest[0];
 784     w0_t[1] = digest[1];
 785     w0_t[2] = digest[2];
 786     w0_t[3] = digest[3];
 787     w1_t[0] = digest[4];
 788     w1_t[1] = 0;
 789     w1_t[2] = 0x80000000;
 790     w1_t[3] = 0;
 791     w2_t[0] = 0;
 792     w2_t[1] = 0;
 793     w2_t[2] = 0;
 794     w2_t[3] = 0;
 795     w3_t[0] = 0;
 796     w3_t[1] = 0;
 797     w3_t[2] = 0;
 798     w3_t[3] = (20 + 4) * 8;
 799
 800     digest[0] = SHA1M_A;
 801     digest[1] = SHA1M_B;
 802     digest[2] = SHA1M_C;
 803     digest[3] = SHA1M_D;
 804     digest[4] = SHA1M_E;
 805
 806     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 807
 808     u32x key[4];
 809
 810     key[0] = swap_workaround (digest[0]);
 811     key[1] = swap_workaround (digest[1]);
 812     key[2] = swap_workaround (digest[2]);
 813     key[3] = swap_workaround (digest[3]);
 814
 815     if (version == 3)
 816     {
 817       key[1] &= 0xff;
 818       key[2]  = 0;
 819       key[3]  = 0;
 820     }
 821
 822     rc4_init_16 (rc4_key, key);
 823
 824     u32x out[4];
 825
 826     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 827
 828     w0_t[0] = swap_workaround (out[0]);
 829     w0_t[1] = swap_workaround (out[1]);
 830     w0_t[2] = swap_workaround (out[2]);
 831     w0_t[3] = swap_workaround (out[3]);
 832     w1_t[0] = 0x80000000;
 833     w1_t[1] = 0;
 834     w1_t[2] = 0;
 835     w1_t[3] = 0;
 836     w2_t[0] = 0;
 837     w2_t[1] = 0;
 838     w2_t[2] = 0;
 839     w2_t[3] = 0;
 840     w3_t[0] = 0;
 841     w3_t[1] = 0;
 842     w3_t[2] = 0;
 843     w3_t[3] = 16 * 8;
 844
 845     digest[0] = SHA1M_A;
 846     digest[1] = SHA1M_B;
 847     digest[2] = SHA1M_C;
 848     digest[3] = SHA1M_D;
 849     digest[4] = SHA1M_E;
 850
 851     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 852
 853     digest[0] = swap_workaround (digest[0]);
 854     digest[1] = swap_workaround (digest[1]);
 855     digest[2] = swap_workaround (digest[2]);
 856     digest[3] = swap_workaround (digest[3]);
 857
 858     rc4_next_16 (rc4_key, 16, j, digest, out);
 859
 860     const u32x r0 = out[0];
 861     const u32x r1 = out[1];
 862     const u32x r2 = out[2];
 863     const u32x r3 = out[3];
 864
 865     #include VECT_COMPARE_S
 866   }
 867 }
 868
 869 extern "C" __global__ void __launch_bounds__ (64, 1) m09800_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)
 870 {
 871 }
 872
 873 extern "C" __global__ void __launch_bounds__ (64, 1) m09800_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)
 874 {
 875 }