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