nv/m09810_a3.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_SIZE4
  34 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  36 #endif
  37
  38 typedef struct
  39 {
  40   u8 S[256];
  41
  42   u32 wtf_its_faster;
  43
  44 } RC4_KEY;
  45
  46 __device__ static void swap (RC4_KEY *rc4_key, const u32 i, const u32 j)
  47 {
  48   u8 tmp;
  49
  50   tmp           = rc4_key->S[i];
  51   rc4_key->S[i] = rc4_key->S[j];
  52   rc4_key->S[j] = tmp;
  53 }
  54
  55 __device__ static void rc4_init_16 (RC4_KEY *rc4_key, const u32 data[4])
  56 {
  57   u32 v = 0x03020100;
  58   u32 a = 0x04040404;
  59
  60   u32 *ptr = (u32 *) rc4_key->S;
  61
  62   #pragma unroll 64
  63   for (u32 i = 0; i < 64; i++)
  64   {
  65     *ptr++ = v; v += a;
  66   }
  67
  68   u32 j = 0;
  69
  70   for (u32 i = 0; i < 16; i++)
  71   {
  72     u32 idx = i * 16;
  73
  74     u32 v;
  75
  76     v = data[0];
  77
  78     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  79     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  80     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
  81     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
  82
  83     v = data[1];
  84
  85     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  86     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  87     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
  88     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
  89
  90     v = data[2];
  91
  92     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
  93     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
  94     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
  95     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
  96
  97     v = data[3];
  98
  99     j += rc4_key->S[idx] + (v >>  0); swap (rc4_key, idx, j & 0xff); idx++;
 100     j += rc4_key->S[idx] + (v >>  8); swap (rc4_key, idx, j & 0xff); idx++;
 101     j += rc4_key->S[idx] + (v >> 16); swap (rc4_key, idx, j & 0xff); idx++;
 102     j += rc4_key->S[idx] + (v >> 24); swap (rc4_key, idx, j & 0xff); idx++;
 103   }
 104 }
 105
 106 __device__ static u8 rc4_next_16 (RC4_KEY *rc4_key, u8 i, u8 j, const u32 in[4], u32 out[4])
 107 {
 108   for (u32 k = 0; k < 4; k++)
 109   {
 110     u32 xor4 = 0;
 111
 112     u8 idx;
 113
 114     i += 1;
 115     j += rc4_key->S[i];
 116
 117     swap (rc4_key, i, j);
 118
 119     idx = rc4_key->S[i] + rc4_key->S[j];
 120
 121     xor4 |= rc4_key->S[idx] <<  0;
 122
 123     i += 1;
 124     j += rc4_key->S[i];
 125
 126     swap (rc4_key, i, j);
 127
 128     idx = rc4_key->S[i] + rc4_key->S[j];
 129
 130     xor4 |= rc4_key->S[idx] <<  8;
 131
 132     i += 1;
 133     j += rc4_key->S[i];
 134
 135     swap (rc4_key, i, j);
 136
 137     idx = rc4_key->S[i] + rc4_key->S[j];
 138
 139     xor4 |= rc4_key->S[idx] << 16;
 140
 141     i += 1;
 142     j += rc4_key->S[i];
 143
 144     swap (rc4_key, i, j);
 145
 146     idx = rc4_key->S[i] + rc4_key->S[j];
 147
 148     xor4 |= rc4_key->S[idx] << 24;
 149
 150     out[k] = in[k] ^ xor4;
 151   }
 152
 153   return j;
 154 }
 155
 156 __device__ static void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
 157 {
 158   u32x A = digest[0];
 159   u32x B = digest[1];
 160   u32x C = digest[2];
 161   u32x D = digest[3];
 162   u32x E = digest[4];
 163
 164   u32x w0_t = w0[0];
 165   u32x w1_t = w0[1];
 166   u32x w2_t = w0[2];
 167   u32x w3_t = w0[3];
 168   u32x w4_t = w1[0];
 169   u32x w5_t = w1[1];
 170   u32x w6_t = w1[2];
 171   u32x w7_t = w1[3];
 172   u32x w8_t = w2[0];
 173   u32x w9_t = w2[1];
 174   u32x wa_t = w2[2];
 175   u32x wb_t = w2[3];
 176   u32x wc_t = w3[0];
 177   u32x wd_t = w3[1];
 178   u32x we_t = w3[2];
 179   u32x wf_t = w3[3];
 180
 181   #undef K
 182   #define K SHA1C00
 183
 184   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
 185   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
 186   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
 187   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
 188   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
 189   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
 190   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
 191   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
 192   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
 193   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
 194   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
 195   SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
 196   SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
 197   SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
 198   SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
 199   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
 200   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
 201   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
 202   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
 203   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
 204
 205   #undef K
 206   #define K SHA1C01
 207
 208   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
 209   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
 210   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
 211   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
 212   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
 213   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
 214   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
 215   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
 216   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
 217   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
 218   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
 219   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
 220   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
 221   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
 222   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
 223   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
 224   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
 225   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
 226   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
 227   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
 228
 229   #undef K
 230   #define K SHA1C02
 231
 232   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
 233   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
 234   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
 235   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
 236   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
 237   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
 238   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
 239   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
 240   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
 241   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
 242   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
 243   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
 244   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
 245   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
 246   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
 247   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
 248   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
 249   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
 250   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
 251   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
 252
 253   #undef K
 254   #define K SHA1C03
 255
 256   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
 257   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
 258   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
 259   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
 260   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
 261   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
 262   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
 263   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
 264   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
 265   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
 266   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
 267   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
 268   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
 269   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
 270   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
 271   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
 272   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
 273   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
 274   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
 275   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
 276
 277   digest[0] += A;
 278   digest[1] += B;
 279   digest[2] += C;
 280   digest[3] += D;
 281   digest[4] += E;
 282 }
 283
 284 __device__ __constant__ bf_t c_bfs[1024];
 285
 286 __device__ static void m09810m (RC4_KEY rc4_keys[64], u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 287 {
 288   /**
 289    * modifier
 290    */
 291
 292   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 293   const u32 lid = threadIdx.x;
 294
 295   RC4_KEY *rc4_key = &rc4_keys[lid];
 296
 297   /**
 298    * esalt
 299    */
 300
 301   const u32 version = oldoffice34_bufs[salt_pos].version;
 302
 303   u32 encryptedVerifier[4];
 304
 305   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 306   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 307   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 308   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 309
 310   /**
 311    * loop
 312    */
 313
 314   u32x w0l = w0[0];
 315
 316   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 317   {
 318     const u32 w0r = c_bfs[il_pos].i;
 319
 320     w0[0] = w0l | w0r;
 321
 322     u32x key[4];
 323
 324     key[0] = w0[0];
 325     key[1] = w0[1] & 0xff;
 326     key[2] = 0;
 327     key[3] = 0;
 328
 329     rc4_init_16 (rc4_key, key);
 330
 331     u32x out[4];
 332
 333     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 334
 335     u32x w0_t[4];
 336     u32x w1_t[4];
 337     u32x w2_t[4];
 338     u32x w3_t[4];
 339
 340     w0_t[0] = swap_workaround (out[0]);
 341     w0_t[1] = swap_workaround (out[1]);
 342     w0_t[2] = swap_workaround (out[2]);
 343     w0_t[3] = swap_workaround (out[3]);
 344     w1_t[0] = 0x80000000;
 345     w1_t[1] = 0;
 346     w1_t[2] = 0;
 347     w1_t[3] = 0;
 348     w2_t[0] = 0;
 349     w2_t[1] = 0;
 350     w2_t[2] = 0;
 351     w2_t[3] = 0;
 352     w3_t[0] = 0;
 353     w3_t[1] = 0;
 354     w3_t[2] = 0;
 355     w3_t[3] = 16 * 8;
 356
 357     u32x digest[5];
 358
 359     digest[0] = SHA1M_A;
 360     digest[1] = SHA1M_B;
 361     digest[2] = SHA1M_C;
 362     digest[3] = SHA1M_D;
 363     digest[4] = SHA1M_E;
 364
 365     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 366
 367     digest[0] = swap_workaround (digest[0]);
 368     digest[1] = swap_workaround (digest[1]);
 369     digest[2] = swap_workaround (digest[2]);
 370     digest[3] = swap_workaround (digest[3]);
 371
 372     rc4_next_16 (rc4_key, 16, j, digest, out);
 373
 374     const u32x r0 = out[0];
 375     const u32x r1 = out[1];
 376     const u32x r2 = out[2];
 377     const u32x r3 = out[3];
 378
 379     #include VECT_COMPARE_M
 380   }
 381 }
 382
 383 __device__ static void m09810s (RC4_KEY rc4_keys[64], u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const u32 pw_len, const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset)
 384 {
 385   /**
 386    * modifier
 387    */
 388
 389   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 390   const u32 lid = threadIdx.x;
 391
 392   RC4_KEY *rc4_key = &rc4_keys[lid];
 393
 394   /**
 395    * digest
 396    */
 397
 398   const u32 search[4] =
 399   {
 400     digests_buf[digests_offset].digest_buf[DGST_R0],
 401     digests_buf[digests_offset].digest_buf[DGST_R1],
 402     digests_buf[digests_offset].digest_buf[DGST_R2],
 403     digests_buf[digests_offset].digest_buf[DGST_R3]
 404   };
 405
 406   /**
 407    * esalt
 408    */
 409
 410   const u32 version = oldoffice34_bufs[salt_pos].version;
 411
 412   u32 encryptedVerifier[4];
 413
 414   encryptedVerifier[0] = oldoffice34_bufs[salt_pos].encryptedVerifier[0];
 415   encryptedVerifier[1] = oldoffice34_bufs[salt_pos].encryptedVerifier[1];
 416   encryptedVerifier[2] = oldoffice34_bufs[salt_pos].encryptedVerifier[2];
 417   encryptedVerifier[3] = oldoffice34_bufs[salt_pos].encryptedVerifier[3];
 418
 419   /**
 420    * loop
 421    */
 422
 423   u32x w0l = w0[0];
 424
 425   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 426   {
 427     const u32 w0r = c_bfs[il_pos].i;
 428
 429     w0[0] = w0l | w0r;
 430
 431     u32x key[4];
 432
 433     key[0] = w0[0];
 434     key[1] = w0[1] & 0xff;
 435     key[2] = 0;
 436     key[3] = 0;
 437
 438     rc4_init_16 (rc4_key, key);
 439
 440     u32x out[4];
 441
 442     u8 j = rc4_next_16 (rc4_key, 0, 0, encryptedVerifier, out);
 443
 444     u32x w0_t[4];
 445     u32x w1_t[4];
 446     u32x w2_t[4];
 447     u32x w3_t[4];
 448
 449     w0_t[0] = swap_workaround (out[0]);
 450     w0_t[1] = swap_workaround (out[1]);
 451     w0_t[2] = swap_workaround (out[2]);
 452     w0_t[3] = swap_workaround (out[3]);
 453     w1_t[0] = 0x80000000;
 454     w1_t[1] = 0;
 455     w1_t[2] = 0;
 456     w1_t[3] = 0;
 457     w2_t[0] = 0;
 458     w2_t[1] = 0;
 459     w2_t[2] = 0;
 460     w2_t[3] = 0;
 461     w3_t[0] = 0;
 462     w3_t[1] = 0;
 463     w3_t[2] = 0;
 464     w3_t[3] = 16 * 8;
 465
 466     u32x digest[5];
 467
 468     digest[0] = SHA1M_A;
 469     digest[1] = SHA1M_B;
 470     digest[2] = SHA1M_C;
 471     digest[3] = SHA1M_D;
 472     digest[4] = SHA1M_E;
 473
 474     sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
 475
 476     digest[0] = swap_workaround (digest[0]);
 477     digest[1] = swap_workaround (digest[1]);
 478     digest[2] = swap_workaround (digest[2]);
 479     digest[3] = swap_workaround (digest[3]);
 480
 481     rc4_next_16 (rc4_key, 16, j, digest, out);
 482
 483     const u32x r0 = out[0];
 484     const u32x r1 = out[1];
 485     const u32x r2 = out[2];
 486     const u32x r3 = out[3];
 487
 488     #include VECT_COMPARE_S
 489   }
 490 }
 491
 492 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 493 {
 494   /**
 495    * base
 496    */
 497
 498   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 499
 500   if (gid >= gid_max) return;
 501
 502   u32x w0[4];
 503
 504   w0[0] = pws[gid].i[ 0];
 505   w0[1] = pws[gid].i[ 1];
 506   w0[2] = pws[gid].i[ 2];
 507   w0[3] = pws[gid].i[ 3];
 508
 509   u32x w1[4];
 510
 511   w1[0] = 0;
 512   w1[1] = 0;
 513   w1[2] = 0;
 514   w1[3] = 0;
 515
 516   u32x w2[4];
 517
 518   w2[0] = 0;
 519   w2[1] = 0;
 520   w2[2] = 0;
 521   w2[3] = 0;
 522
 523   u32x w3[4];
 524
 525   w3[0] = 0;
 526   w3[1] = 0;
 527   w3[2] = 0;
 528   w3[3] = 0;
 529
 530   const u32 pw_len = pws[gid].pw_len;
 531
 532   /**
 533    * main
 534    */
 535
 536   __shared__ RC4_KEY rc4_keys[64];
 537
 538   m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 539 }
 540
 541 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 542 {
 543   /**
 544    * base
 545    */
 546
 547   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 548
 549   if (gid >= gid_max) return;
 550
 551   u32x w0[4];
 552
 553   w0[0] = pws[gid].i[ 0];
 554   w0[1] = pws[gid].i[ 1];
 555   w0[2] = pws[gid].i[ 2];
 556   w0[3] = pws[gid].i[ 3];
 557
 558   u32x w1[4];
 559
 560   w1[0] = pws[gid].i[ 4];
 561   w1[1] = pws[gid].i[ 5];
 562   w1[2] = pws[gid].i[ 6];
 563   w1[3] = pws[gid].i[ 7];
 564
 565   u32x w2[4];
 566
 567   w2[0] = 0;
 568   w2[1] = 0;
 569   w2[2] = 0;
 570   w2[3] = 0;
 571
 572   u32x w3[4];
 573
 574   w3[0] = 0;
 575   w3[1] = 0;
 576   w3[2] = 0;
 577   w3[3] = 0;
 578
 579   const u32 pw_len = pws[gid].pw_len;
 580
 581   /**
 582    * main
 583    */
 584
 585   __shared__ RC4_KEY rc4_keys[64];
 586
 587   m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 588 }
 589
 590 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 591 {
 592   /**
 593    * base
 594    */
 595
 596   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 597
 598   if (gid >= gid_max) return;
 599
 600   u32x w0[4];
 601
 602   w0[0] = pws[gid].i[ 0];
 603   w0[1] = pws[gid].i[ 1];
 604   w0[2] = pws[gid].i[ 2];
 605   w0[3] = pws[gid].i[ 3];
 606
 607   u32x w1[4];
 608
 609   w1[0] = pws[gid].i[ 4];
 610   w1[1] = pws[gid].i[ 5];
 611   w1[2] = pws[gid].i[ 6];
 612   w1[3] = pws[gid].i[ 7];
 613
 614   u32x w2[4];
 615
 616   w2[0] = pws[gid].i[ 8];
 617   w2[1] = pws[gid].i[ 9];
 618   w2[2] = pws[gid].i[10];
 619   w2[3] = pws[gid].i[11];
 620
 621   u32x w3[4];
 622
 623   w3[0] = pws[gid].i[12];
 624   w3[1] = pws[gid].i[13];
 625   w3[2] = 0;
 626   w3[3] = 0;
 627
 628   const u32 pw_len = pws[gid].pw_len;
 629
 630   /**
 631    * main
 632    */
 633
 634   __shared__ RC4_KEY rc4_keys[64];
 635
 636   m09810m (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 637 }
 638
 639 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 640 {
 641   /**
 642    * base
 643    */
 644
 645   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 646
 647   if (gid >= gid_max) return;
 648
 649   u32x w0[4];
 650
 651   w0[0] = pws[gid].i[ 0];
 652   w0[1] = pws[gid].i[ 1];
 653   w0[2] = pws[gid].i[ 2];
 654   w0[3] = pws[gid].i[ 3];
 655
 656   u32x w1[4];
 657
 658   w1[0] = 0;
 659   w1[1] = 0;
 660   w1[2] = 0;
 661   w1[3] = 0;
 662
 663   u32x w2[4];
 664
 665   w2[0] = 0;
 666   w2[1] = 0;
 667   w2[2] = 0;
 668   w2[3] = 0;
 669
 670   u32x w3[4];
 671
 672   w3[0] = 0;
 673   w3[1] = 0;
 674   w3[2] = 0;
 675   w3[3] = 0;
 676
 677   const u32 pw_len = pws[gid].pw_len;
 678
 679   /**
 680    * main
 681    */
 682
 683   __shared__ RC4_KEY rc4_keys[64];
 684
 685   m09810s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 686 }
 687
 688 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 689 {
 690   /**
 691    * base
 692    */
 693
 694   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 695
 696   if (gid >= gid_max) return;
 697
 698   u32x w0[4];
 699
 700   w0[0] = pws[gid].i[ 0];
 701   w0[1] = pws[gid].i[ 1];
 702   w0[2] = pws[gid].i[ 2];
 703   w0[3] = pws[gid].i[ 3];
 704
 705   u32x w1[4];
 706
 707   w1[0] = pws[gid].i[ 4];
 708   w1[1] = pws[gid].i[ 5];
 709   w1[2] = pws[gid].i[ 6];
 710   w1[3] = pws[gid].i[ 7];
 711
 712   u32x w2[4];
 713
 714   w2[0] = 0;
 715   w2[1] = 0;
 716   w2[2] = 0;
 717   w2[3] = 0;
 718
 719   u32x w3[4];
 720
 721   w3[0] = 0;
 722   w3[1] = 0;
 723   w3[2] = 0;
 724   w3[3] = 0;
 725
 726   const u32 pw_len = pws[gid].pw_len;
 727
 728   /**
 729    * main
 730    */
 731
 732   __shared__ RC4_KEY rc4_keys[64];
 733
 734   m09810s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 735 }
 736
 737 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 u32x *bfs_buf, 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 bfs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 738 {
 739   /**
 740    * base
 741    */
 742
 743   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 744
 745   if (gid >= gid_max) return;
 746
 747   u32x w0[4];
 748
 749   w0[0] = pws[gid].i[ 0];
 750   w0[1] = pws[gid].i[ 1];
 751   w0[2] = pws[gid].i[ 2];
 752   w0[3] = pws[gid].i[ 3];
 753
 754   u32x w1[4];
 755
 756   w1[0] = pws[gid].i[ 4];
 757   w1[1] = pws[gid].i[ 5];
 758   w1[2] = pws[gid].i[ 6];
 759   w1[3] = pws[gid].i[ 7];
 760
 761   u32x w2[4];
 762
 763   w2[0] = pws[gid].i[ 8];
 764   w2[1] = pws[gid].i[ 9];
 765   w2[2] = pws[gid].i[10];
 766   w2[3] = pws[gid].i[11];
 767
 768   u32x w3[4];
 769
 770   w3[0] = pws[gid].i[12];
 771   w3[1] = pws[gid].i[13];
 772   w3[2] = 0;
 773   w3[3] = 0;
 774
 775   const u32 pw_len = pws[gid].pw_len;
 776
 777   /**
 778    * main
 779    */
 780
 781   __shared__ RC4_KEY rc4_keys[64];
 782
 783   m09810s (rc4_keys, w0, w1, w2, w3, pw_len, pws, rules_buf, combs_buf, bfs_buf, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, oldoffice34_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);
 784 }