nv/m05600_a3.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _NETNTLMV2_
   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 3
  21 #define DGST_R2 2
  22 #define DGST_R3 1
  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 __device__ static void md4_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
  39 {
  40   u32x a = digest[0];
  41   u32x b = digest[1];
  42   u32x c = digest[2];
  43   u32x d = digest[3];
  44
  45   u32x w0_t = w0[0];
  46   u32x w1_t = w0[1];
  47   u32x w2_t = w0[2];
  48   u32x w3_t = w0[3];
  49   u32x w4_t = w1[0];
  50   u32x w5_t = w1[1];
  51   u32x w6_t = w1[2];
  52   u32x w7_t = w1[3];
  53   u32x w8_t = w2[0];
  54   u32x w9_t = w2[1];
  55   u32x wa_t = w2[2];
  56   u32x wb_t = w2[3];
  57   u32x wc_t = w3[0];
  58   u32x wd_t = w3[1];
  59   u32x we_t = w3[2];
  60   u32x wf_t = w3[3];
  61
  62   MD4_STEP (MD4_Fo, a, b, c, d, w0_t, MD4C00, MD4S00);
  63   MD4_STEP (MD4_Fo, d, a, b, c, w1_t, MD4C00, MD4S01);
  64   MD4_STEP (MD4_Fo, c, d, a, b, w2_t, MD4C00, MD4S02);
  65   MD4_STEP (MD4_Fo, b, c, d, a, w3_t, MD4C00, MD4S03);
  66   MD4_STEP (MD4_Fo, a, b, c, d, w4_t, MD4C00, MD4S00);
  67   MD4_STEP (MD4_Fo, d, a, b, c, w5_t, MD4C00, MD4S01);
  68   MD4_STEP (MD4_Fo, c, d, a, b, w6_t, MD4C00, MD4S02);
  69   MD4_STEP (MD4_Fo, b, c, d, a, w7_t, MD4C00, MD4S03);
  70   MD4_STEP (MD4_Fo, a, b, c, d, w8_t, MD4C00, MD4S00);
  71   MD4_STEP (MD4_Fo, d, a, b, c, w9_t, MD4C00, MD4S01);
  72   MD4_STEP (MD4_Fo, c, d, a, b, wa_t, MD4C00, MD4S02);
  73   MD4_STEP (MD4_Fo, b, c, d, a, wb_t, MD4C00, MD4S03);
  74   MD4_STEP (MD4_Fo, a, b, c, d, wc_t, MD4C00, MD4S00);
  75   MD4_STEP (MD4_Fo, d, a, b, c, wd_t, MD4C00, MD4S01);
  76   MD4_STEP (MD4_Fo, c, d, a, b, we_t, MD4C00, MD4S02);
  77   MD4_STEP (MD4_Fo, b, c, d, a, wf_t, MD4C00, MD4S03);
  78
  79   MD4_STEP (MD4_Go, a, b, c, d, w0_t, MD4C01, MD4S10);
  80   MD4_STEP (MD4_Go, d, a, b, c, w4_t, MD4C01, MD4S11);
  81   MD4_STEP (MD4_Go, c, d, a, b, w8_t, MD4C01, MD4S12);
  82   MD4_STEP (MD4_Go, b, c, d, a, wc_t, MD4C01, MD4S13);
  83   MD4_STEP (MD4_Go, a, b, c, d, w1_t, MD4C01, MD4S10);
  84   MD4_STEP (MD4_Go, d, a, b, c, w5_t, MD4C01, MD4S11);
  85   MD4_STEP (MD4_Go, c, d, a, b, w9_t, MD4C01, MD4S12);
  86   MD4_STEP (MD4_Go, b, c, d, a, wd_t, MD4C01, MD4S13);
  87   MD4_STEP (MD4_Go, a, b, c, d, w2_t, MD4C01, MD4S10);
  88   MD4_STEP (MD4_Go, d, a, b, c, w6_t, MD4C01, MD4S11);
  89   MD4_STEP (MD4_Go, c, d, a, b, wa_t, MD4C01, MD4S12);
  90   MD4_STEP (MD4_Go, b, c, d, a, we_t, MD4C01, MD4S13);
  91   MD4_STEP (MD4_Go, a, b, c, d, w3_t, MD4C01, MD4S10);
  92   MD4_STEP (MD4_Go, d, a, b, c, w7_t, MD4C01, MD4S11);
  93   MD4_STEP (MD4_Go, c, d, a, b, wb_t, MD4C01, MD4S12);
  94   MD4_STEP (MD4_Go, b, c, d, a, wf_t, MD4C01, MD4S13);
  95
  96   MD4_STEP (MD4_H , a, b, c, d, w0_t, MD4C02, MD4S20);
  97   MD4_STEP (MD4_H , d, a, b, c, w8_t, MD4C02, MD4S21);
  98   MD4_STEP (MD4_H , c, d, a, b, w4_t, MD4C02, MD4S22);
  99   MD4_STEP (MD4_H , b, c, d, a, wc_t, MD4C02, MD4S23);
 100   MD4_STEP (MD4_H , a, b, c, d, w2_t, MD4C02, MD4S20);
 101   MD4_STEP (MD4_H , d, a, b, c, wa_t, MD4C02, MD4S21);
 102   MD4_STEP (MD4_H , c, d, a, b, w6_t, MD4C02, MD4S22);
 103   MD4_STEP (MD4_H , b, c, d, a, we_t, MD4C02, MD4S23);
 104   MD4_STEP (MD4_H , a, b, c, d, w1_t, MD4C02, MD4S20);
 105   MD4_STEP (MD4_H , d, a, b, c, w9_t, MD4C02, MD4S21);
 106   MD4_STEP (MD4_H , c, d, a, b, w5_t, MD4C02, MD4S22);
 107   MD4_STEP (MD4_H , b, c, d, a, wd_t, MD4C02, MD4S23);
 108   MD4_STEP (MD4_H , a, b, c, d, w3_t, MD4C02, MD4S20);
 109   MD4_STEP (MD4_H , d, a, b, c, wb_t, MD4C02, MD4S21);
 110   MD4_STEP (MD4_H , c, d, a, b, w7_t, MD4C02, MD4S22);
 111   MD4_STEP (MD4_H , b, c, d, a, wf_t, MD4C02, MD4S23);
 112
 113   digest[0] += a;
 114   digest[1] += b;
 115   digest[2] += c;
 116   digest[3] += d;
 117 }
 118
 119 __device__ static void md5_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[4])
 120 {
 121   u32x a = digest[0];
 122   u32x b = digest[1];
 123   u32x c = digest[2];
 124   u32x d = digest[3];
 125
 126   u32x w0_t = w0[0];
 127   u32x w1_t = w0[1];
 128   u32x w2_t = w0[2];
 129   u32x w3_t = w0[3];
 130   u32x w4_t = w1[0];
 131   u32x w5_t = w1[1];
 132   u32x w6_t = w1[2];
 133   u32x w7_t = w1[3];
 134   u32x w8_t = w2[0];
 135   u32x w9_t = w2[1];
 136   u32x wa_t = w2[2];
 137   u32x wb_t = w2[3];
 138   u32x wc_t = w3[0];
 139   u32x wd_t = w3[1];
 140   u32x we_t = w3[2];
 141   u32x wf_t = w3[3];
 142
 143   MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
 144   MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
 145   MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
 146   MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
 147   MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
 148   MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
 149   MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
 150   MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
 151   MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
 152   MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
 153   MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
 154   MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
 155   MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
 156   MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
 157   MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
 158   MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
 159
 160   MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
 161   MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
 162   MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
 163   MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
 164   MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
 165   MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
 166   MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
 167   MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
 168   MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
 169   MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
 170   MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
 171   MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
 172   MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
 173   MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
 174   MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
 175   MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
 176
 177   MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
 178   MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
 179   MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
 180   MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
 181   MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
 182   MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
 183   MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
 184   MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
 185   MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
 186   MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
 187   MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
 188   MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
 189   MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
 190   MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
 191   MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
 192   MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
 193
 194   MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
 195   MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
 196   MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
 197   MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
 198   MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
 199   MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
 200   MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
 201   MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
 202   MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
 203   MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
 204   MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
 205   MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
 206   MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
 207   MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
 208   MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
 209   MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
 210
 211   digest[0] += a;
 212   digest[1] += b;
 213   digest[2] += c;
 214   digest[3] += d;
 215 }
 216
 217 __device__ static void hmac_md5_pad (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[4], u32x opad[4])
 218 {
 219   w0[0] = w0[0] ^ 0x36363636;
 220   w0[1] = w0[1] ^ 0x36363636;
 221   w0[2] = w0[2] ^ 0x36363636;
 222   w0[3] = w0[3] ^ 0x36363636;
 223   w1[0] = w1[0] ^ 0x36363636;
 224   w1[1] = w1[1] ^ 0x36363636;
 225   w1[2] = w1[2] ^ 0x36363636;
 226   w1[3] = w1[3] ^ 0x36363636;
 227   w2[0] = w2[0] ^ 0x36363636;
 228   w2[1] = w2[1] ^ 0x36363636;
 229   w2[2] = w2[2] ^ 0x36363636;
 230   w2[3] = w2[3] ^ 0x36363636;
 231   w3[0] = w3[0] ^ 0x36363636;
 232   w3[1] = w3[1] ^ 0x36363636;
 233   w3[2] = w3[2] ^ 0x36363636;
 234   w3[3] = w3[3] ^ 0x36363636;
 235
 236   ipad[0] = MD5M_A;
 237   ipad[1] = MD5M_B;
 238   ipad[2] = MD5M_C;
 239   ipad[3] = MD5M_D;
 240
 241   md5_transform (w0, w1, w2, w3, ipad);
 242
 243   w0[0] = w0[0] ^ 0x6a6a6a6a;
 244   w0[1] = w0[1] ^ 0x6a6a6a6a;
 245   w0[2] = w0[2] ^ 0x6a6a6a6a;
 246   w0[3] = w0[3] ^ 0x6a6a6a6a;
 247   w1[0] = w1[0] ^ 0x6a6a6a6a;
 248   w1[1] = w1[1] ^ 0x6a6a6a6a;
 249   w1[2] = w1[2] ^ 0x6a6a6a6a;
 250   w1[3] = w1[3] ^ 0x6a6a6a6a;
 251   w2[0] = w2[0] ^ 0x6a6a6a6a;
 252   w2[1] = w2[1] ^ 0x6a6a6a6a;
 253   w2[2] = w2[2] ^ 0x6a6a6a6a;
 254   w2[3] = w2[3] ^ 0x6a6a6a6a;
 255   w3[0] = w3[0] ^ 0x6a6a6a6a;
 256   w3[1] = w3[1] ^ 0x6a6a6a6a;
 257   w3[2] = w3[2] ^ 0x6a6a6a6a;
 258   w3[3] = w3[3] ^ 0x6a6a6a6a;
 259
 260   opad[0] = MD5M_A;
 261   opad[1] = MD5M_B;
 262   opad[2] = MD5M_C;
 263   opad[3] = MD5M_D;
 264
 265   md5_transform (w0, w1, w2, w3, opad);
 266 }
 267
 268 __device__ static void hmac_md5_run (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[4], u32x opad[4], u32x digest[4])
 269 {
 270   digest[0] = ipad[0];
 271   digest[1] = ipad[1];
 272   digest[2] = ipad[2];
 273   digest[3] = ipad[3];
 274
 275   md5_transform (w0, w1, w2, w3, digest);
 276
 277   w0[0] = digest[0];
 278   w0[1] = digest[1];
 279   w0[2] = digest[2];
 280   w0[3] = digest[3];
 281   w1[0] = 0x80;
 282   w1[1] = 0;
 283   w1[2] = 0;
 284   w1[3] = 0;
 285   w2[0] = 0;
 286   w2[1] = 0;
 287   w2[2] = 0;
 288   w2[3] = 0;
 289   w3[0] = 0;
 290   w3[1] = 0;
 291   w3[2] = (64 + 16) * 8;
 292   w3[3] = 0;
 293
 294   digest[0] = opad[0];
 295   digest[1] = opad[1];
 296   digest[2] = opad[2];
 297   digest[3] = opad[3];
 298
 299   md5_transform (w0, w1, w2, w3, digest);
 300 }
 301
 302 __device__ __constant__ bf_t c_bfs[1024];
 303
 304 __device__ static void m05600m (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 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 netntlm_t *netntlm_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, u32 s_userdomain_buf[64], u32 s_chall_buf[256])
 305 {
 306   /**
 307    * modifier
 308    */
 309
 310   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 311   const u32 lid = threadIdx.x;
 312
 313   /**
 314    * prepare
 315    */
 316
 317   const u32 userdomain_len = netntlm_bufs[salt_pos].user_len
 318                             + netntlm_bufs[salt_pos].domain_len;
 319
 320   const u32 chall_len = netntlm_bufs[salt_pos].srvchall_len
 321                        + netntlm_bufs[salt_pos].clichall_len;
 322
 323   /**
 324    * loop
 325    */
 326
 327   u32x w0l = w0[0];
 328
 329   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 330   {
 331     const u32 w0r = c_bfs[il_pos].i;
 332
 333     w0[0] = w0l | w0r;
 334
 335     u32x digest[4];
 336
 337     digest[0] = MD4M_A;
 338     digest[1] = MD4M_B;
 339     digest[2] = MD4M_C;
 340     digest[3] = MD4M_D;
 341
 342     md4_transform (w0, w1, w2, w3, digest);
 343
 344     u32x w0_t[4];
 345     u32x w1_t[4];
 346     u32x w2_t[4];
 347     u32x w3_t[4];
 348
 349     w0_t[0] = digest[0];
 350     w0_t[1] = digest[1];
 351     w0_t[2] = digest[2];
 352     w0_t[3] = digest[3];
 353     w1_t[0] = 0;
 354     w1_t[1] = 0;
 355     w1_t[2] = 0;
 356     w1_t[3] = 0;
 357     w2_t[0] = 0;
 358     w2_t[1] = 0;
 359     w2_t[2] = 0;
 360     w2_t[3] = 0;
 361     w3_t[0] = 0;
 362     w3_t[1] = 0;
 363     w3_t[2] = 0;
 364     w3_t[3] = 0;
 365
 366     digest[0] = MD5M_A;
 367     digest[1] = MD5M_B;
 368     digest[2] = MD5M_C;
 369     digest[3] = MD5M_D;
 370
 371     u32x ipad[4];
 372     u32x opad[4];
 373
 374     hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
 375
 376     int left;
 377     int off;
 378
 379     for (left = userdomain_len, off = 0; left >= 56; left -= 64, off += 16)
 380     {
 381       w0_t[0] = s_userdomain_buf[off +  0];
 382       w0_t[1] = s_userdomain_buf[off +  1];
 383       w0_t[2] = s_userdomain_buf[off +  2];
 384       w0_t[3] = s_userdomain_buf[off +  3];
 385       w1_t[0] = s_userdomain_buf[off +  4];
 386       w1_t[1] = s_userdomain_buf[off +  5];
 387       w1_t[2] = s_userdomain_buf[off +  6];
 388       w1_t[3] = s_userdomain_buf[off +  7];
 389       w2_t[0] = s_userdomain_buf[off +  8];
 390       w2_t[1] = s_userdomain_buf[off +  9];
 391       w2_t[2] = s_userdomain_buf[off + 10];
 392       w2_t[3] = s_userdomain_buf[off + 11];
 393       w3_t[0] = s_userdomain_buf[off + 12];
 394       w3_t[1] = s_userdomain_buf[off + 13];
 395       w3_t[2] = s_userdomain_buf[off + 14];
 396       w3_t[3] = s_userdomain_buf[off + 15];
 397
 398       md5_transform (w0_t, w1_t, w2_t, w3_t, ipad);
 399     }
 400
 401     w0_t[0] = s_userdomain_buf[off +  0];
 402     w0_t[1] = s_userdomain_buf[off +  1];
 403     w0_t[2] = s_userdomain_buf[off +  2];
 404     w0_t[3] = s_userdomain_buf[off +  3];
 405     w1_t[0] = s_userdomain_buf[off +  4];
 406     w1_t[1] = s_userdomain_buf[off +  5];
 407     w1_t[2] = s_userdomain_buf[off +  6];
 408     w1_t[3] = s_userdomain_buf[off +  7];
 409     w2_t[0] = s_userdomain_buf[off +  8];
 410     w2_t[1] = s_userdomain_buf[off +  9];
 411     w2_t[2] = s_userdomain_buf[off + 10];
 412     w2_t[3] = s_userdomain_buf[off + 11];
 413     w3_t[0] = s_userdomain_buf[off + 12];
 414     w3_t[1] = s_userdomain_buf[off + 13];
 415     w3_t[2] = (64 + userdomain_len) * 8;
 416     w3_t[3] = 0;
 417
 418     hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
 419
 420     w0_t[0] = digest[0];
 421     w0_t[1] = digest[1];
 422     w0_t[2] = digest[2];
 423     w0_t[3] = digest[3];
 424     w1_t[0] = 0;
 425     w1_t[1] = 0;
 426     w1_t[2] = 0;
 427     w1_t[3] = 0;
 428     w2_t[0] = 0;
 429     w2_t[1] = 0;
 430     w2_t[2] = 0;
 431     w2_t[3] = 0;
 432     w3_t[0] = 0;
 433     w3_t[1] = 0;
 434     w3_t[2] = 0;
 435     w3_t[3] = 0;
 436
 437     digest[0] = MD5M_A;
 438     digest[1] = MD5M_B;
 439     digest[2] = MD5M_C;
 440     digest[3] = MD5M_D;
 441
 442     hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
 443
 444     for (left = chall_len, off = 0; left >= 56; left -= 64, off += 16)
 445     {
 446       w0_t[0] = s_chall_buf[off +  0];
 447       w0_t[1] = s_chall_buf[off +  1];
 448       w0_t[2] = s_chall_buf[off +  2];
 449       w0_t[3] = s_chall_buf[off +  3];
 450       w1_t[0] = s_chall_buf[off +  4];
 451       w1_t[1] = s_chall_buf[off +  5];
 452       w1_t[2] = s_chall_buf[off +  6];
 453       w1_t[3] = s_chall_buf[off +  7];
 454       w2_t[0] = s_chall_buf[off +  8];
 455       w2_t[1] = s_chall_buf[off +  9];
 456       w2_t[2] = s_chall_buf[off + 10];
 457       w2_t[3] = s_chall_buf[off + 11];
 458       w3_t[0] = s_chall_buf[off + 12];
 459       w3_t[1] = s_chall_buf[off + 13];
 460       w3_t[2] = s_chall_buf[off + 14];
 461       w3_t[3] = s_chall_buf[off + 15];
 462
 463       md5_transform (w0_t, w1_t, w2_t, w3_t, ipad);
 464     }
 465
 466     w0_t[0] = s_chall_buf[off +  0];
 467     w0_t[1] = s_chall_buf[off +  1];
 468     w0_t[2] = s_chall_buf[off +  2];
 469     w0_t[3] = s_chall_buf[off +  3];
 470     w1_t[0] = s_chall_buf[off +  4];
 471     w1_t[1] = s_chall_buf[off +  5];
 472     w1_t[2] = s_chall_buf[off +  6];
 473     w1_t[3] = s_chall_buf[off +  7];
 474     w2_t[0] = s_chall_buf[off +  8];
 475     w2_t[1] = s_chall_buf[off +  9];
 476     w2_t[2] = s_chall_buf[off + 10];
 477     w2_t[3] = s_chall_buf[off + 11];
 478     w3_t[0] = s_chall_buf[off + 12];
 479     w3_t[1] = s_chall_buf[off + 13];
 480     w3_t[2] = (64 + chall_len) * 8;
 481     w3_t[3] = 0;
 482
 483     hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
 484
 485     const u32x r0 = digest[0];
 486     const u32x r1 = digest[3];
 487     const u32x r2 = digest[2];
 488     const u32x r3 = digest[1];
 489
 490     #include VECT_COMPARE_M
 491   }
 492 }
 493
 494 __device__ static void m05600s (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 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 netntlm_t *netntlm_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, u32 s_userdomain_buf[64], u32 s_chall_buf[256])
 495 {
 496   /**
 497    * modifier
 498    */
 499
 500   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 501   const u32 lid = threadIdx.x;
 502
 503   /**
 504    * digest
 505    */
 506
 507   const u32 search[4] =
 508   {
 509     digests_buf[digests_offset].digest_buf[DGST_R0],
 510     digests_buf[digests_offset].digest_buf[DGST_R1],
 511     digests_buf[digests_offset].digest_buf[DGST_R2],
 512     digests_buf[digests_offset].digest_buf[DGST_R3]
 513   };
 514
 515   /**
 516    * prepare
 517    */
 518
 519   const u32 userdomain_len = netntlm_bufs[salt_pos].user_len
 520                             + netntlm_bufs[salt_pos].domain_len;
 521
 522   const u32 chall_len = netntlm_bufs[salt_pos].srvchall_len
 523                        + netntlm_bufs[salt_pos].clichall_len;
 524
 525   /**
 526    * loop
 527    */
 528
 529   u32x w0l = w0[0];
 530
 531   for (u32 il_pos = 0; il_pos < bfs_cnt; il_pos++)
 532   {
 533     const u32 w0r = c_bfs[il_pos].i;
 534
 535     w0[0] = w0l | w0r;
 536
 537     u32x digest[4];
 538
 539     digest[0] = MD4M_A;
 540     digest[1] = MD4M_B;
 541     digest[2] = MD4M_C;
 542     digest[3] = MD4M_D;
 543
 544     md4_transform (w0, w1, w2, w3, digest);
 545
 546     u32x w0_t[4];
 547     u32x w1_t[4];
 548     u32x w2_t[4];
 549     u32x w3_t[4];
 550
 551     w0_t[0] = digest[0];
 552     w0_t[1] = digest[1];
 553     w0_t[2] = digest[2];
 554     w0_t[3] = digest[3];
 555     w1_t[0] = 0;
 556     w1_t[1] = 0;
 557     w1_t[2] = 0;
 558     w1_t[3] = 0;
 559     w2_t[0] = 0;
 560     w2_t[1] = 0;
 561     w2_t[2] = 0;
 562     w2_t[3] = 0;
 563     w3_t[0] = 0;
 564     w3_t[1] = 0;
 565     w3_t[2] = 0;
 566     w3_t[3] = 0;
 567
 568     digest[0] = MD5M_A;
 569     digest[1] = MD5M_B;
 570     digest[2] = MD5M_C;
 571     digest[3] = MD5M_D;
 572
 573     u32x ipad[4];
 574     u32x opad[4];
 575
 576     hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
 577
 578     int left;
 579     int off;
 580
 581     for (left = userdomain_len, off = 0; left >= 56; left -= 64, off += 16)
 582     {
 583       w0_t[0] = s_userdomain_buf[off +  0];
 584       w0_t[1] = s_userdomain_buf[off +  1];
 585       w0_t[2] = s_userdomain_buf[off +  2];
 586       w0_t[3] = s_userdomain_buf[off +  3];
 587       w1_t[0] = s_userdomain_buf[off +  4];
 588       w1_t[1] = s_userdomain_buf[off +  5];
 589       w1_t[2] = s_userdomain_buf[off +  6];
 590       w1_t[3] = s_userdomain_buf[off +  7];
 591       w2_t[0] = s_userdomain_buf[off +  8];
 592       w2_t[1] = s_userdomain_buf[off +  9];
 593       w2_t[2] = s_userdomain_buf[off + 10];
 594       w2_t[3] = s_userdomain_buf[off + 11];
 595       w3_t[0] = s_userdomain_buf[off + 12];
 596       w3_t[1] = s_userdomain_buf[off + 13];
 597       w3_t[2] = s_userdomain_buf[off + 14];
 598       w3_t[3] = s_userdomain_buf[off + 15];
 599
 600       md5_transform (w0_t, w1_t, w2_t, w3_t, ipad);
 601     }
 602
 603     w0_t[0] = s_userdomain_buf[off +  0];
 604     w0_t[1] = s_userdomain_buf[off +  1];
 605     w0_t[2] = s_userdomain_buf[off +  2];
 606     w0_t[3] = s_userdomain_buf[off +  3];
 607     w1_t[0] = s_userdomain_buf[off +  4];
 608     w1_t[1] = s_userdomain_buf[off +  5];
 609     w1_t[2] = s_userdomain_buf[off +  6];
 610     w1_t[3] = s_userdomain_buf[off +  7];
 611     w2_t[0] = s_userdomain_buf[off +  8];
 612     w2_t[1] = s_userdomain_buf[off +  9];
 613     w2_t[2] = s_userdomain_buf[off + 10];
 614     w2_t[3] = s_userdomain_buf[off + 11];
 615     w3_t[0] = s_userdomain_buf[off + 12];
 616     w3_t[1] = s_userdomain_buf[off + 13];
 617     w3_t[2] = (64 + userdomain_len) * 8;
 618     w3_t[3] = 0;
 619
 620     hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
 621
 622     w0_t[0] = digest[0];
 623     w0_t[1] = digest[1];
 624     w0_t[2] = digest[2];
 625     w0_t[3] = digest[3];
 626     w1_t[0] = 0;
 627     w1_t[1] = 0;
 628     w1_t[2] = 0;
 629     w1_t[3] = 0;
 630     w2_t[0] = 0;
 631     w2_t[1] = 0;
 632     w2_t[2] = 0;
 633     w2_t[3] = 0;
 634     w3_t[0] = 0;
 635     w3_t[1] = 0;
 636     w3_t[2] = 0;
 637     w3_t[3] = 0;
 638
 639     digest[0] = MD5M_A;
 640     digest[1] = MD5M_B;
 641     digest[2] = MD5M_C;
 642     digest[3] = MD5M_D;
 643
 644     hmac_md5_pad (w0_t, w1_t, w2_t, w3_t, ipad, opad);
 645
 646     for (left = chall_len, off = 0; left >= 56; left -= 64, off += 16)
 647     {
 648       w0_t[0] = s_chall_buf[off +  0];
 649       w0_t[1] = s_chall_buf[off +  1];
 650       w0_t[2] = s_chall_buf[off +  2];
 651       w0_t[3] = s_chall_buf[off +  3];
 652       w1_t[0] = s_chall_buf[off +  4];
 653       w1_t[1] = s_chall_buf[off +  5];
 654       w1_t[2] = s_chall_buf[off +  6];
 655       w1_t[3] = s_chall_buf[off +  7];
 656       w2_t[0] = s_chall_buf[off +  8];
 657       w2_t[1] = s_chall_buf[off +  9];
 658       w2_t[2] = s_chall_buf[off + 10];
 659       w2_t[3] = s_chall_buf[off + 11];
 660       w3_t[0] = s_chall_buf[off + 12];
 661       w3_t[1] = s_chall_buf[off + 13];
 662       w3_t[2] = s_chall_buf[off + 14];
 663       w3_t[3] = s_chall_buf[off + 15];
 664
 665       md5_transform (w0_t, w1_t, w2_t, w3_t, ipad);
 666     }
 667
 668     w0_t[0] = s_chall_buf[off +  0];
 669     w0_t[1] = s_chall_buf[off +  1];
 670     w0_t[2] = s_chall_buf[off +  2];
 671     w0_t[3] = s_chall_buf[off +  3];
 672     w1_t[0] = s_chall_buf[off +  4];
 673     w1_t[1] = s_chall_buf[off +  5];
 674     w1_t[2] = s_chall_buf[off +  6];
 675     w1_t[3] = s_chall_buf[off +  7];
 676     w2_t[0] = s_chall_buf[off +  8];
 677     w2_t[1] = s_chall_buf[off +  9];
 678     w2_t[2] = s_chall_buf[off + 10];
 679     w2_t[3] = s_chall_buf[off + 11];
 680     w3_t[0] = s_chall_buf[off + 12];
 681     w3_t[1] = s_chall_buf[off + 13];
 682     w3_t[2] = (64 + chall_len) * 8;
 683     w3_t[3] = 0;
 684
 685     hmac_md5_run (w0_t, w1_t, w2_t, w3_t, ipad, opad, digest);
 686
 687     const u32x r0 = digest[0];
 688     const u32x r1 = digest[3];
 689     const u32x r2 = digest[2];
 690     const u32x r3 = digest[1];
 691
 692     #include VECT_COMPARE_S
 693   }
 694 }
 695
 696 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 697 {
 698   /**
 699    * base
 700    */
 701
 702   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 703   const u32 lid = threadIdx.x;
 704
 705   u32x w0[4];
 706
 707   w0[0] = pws[gid].i[ 0];
 708   w0[1] = pws[gid].i[ 1];
 709   w0[2] = pws[gid].i[ 2];
 710   w0[3] = pws[gid].i[ 3];
 711
 712   u32x w1[4];
 713
 714   w1[0] = 0;
 715   w1[1] = 0;
 716   w1[2] = 0;
 717   w1[3] = 0;
 718
 719   u32x w2[4];
 720
 721   w2[0] = 0;
 722   w2[1] = 0;
 723   w2[2] = 0;
 724   w2[3] = 0;
 725
 726   u32x w3[4];
 727
 728   w3[0] = 0;
 729   w3[1] = 0;
 730   w3[2] = pws[gid].i[14];
 731   w3[3] = 0;
 732
 733   const u32 pw_len = pws[gid].pw_len;
 734
 735   /**
 736    * salt
 737    */
 738
 739   __shared__ u32 s_userdomain_buf[64];
 740   __shared__ u32 s_chall_buf[256];
 741
 742   if (lid < 64)
 743   {
 744     s_userdomain_buf[lid] = netntlm_bufs[salt_pos].userdomain_buf[lid];
 745   }
 746
 747   s_chall_buf[lid] = netntlm_bufs[salt_pos].chall_buf[lid];
 748
 749   __syncthreads ();
 750
 751   if (gid >= gid_max) return;
 752
 753   /**
 754    * main
 755    */
 756
 757   m05600m (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, netntlm_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, s_userdomain_buf, s_chall_buf);
 758 }
 759
 760 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 761 {
 762   /**
 763    * base
 764    */
 765
 766   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 767   const u32 lid = threadIdx.x;
 768
 769   u32x w0[4];
 770
 771   w0[0] = pws[gid].i[ 0];
 772   w0[1] = pws[gid].i[ 1];
 773   w0[2] = pws[gid].i[ 2];
 774   w0[3] = pws[gid].i[ 3];
 775
 776   u32x w1[4];
 777
 778   w1[0] = pws[gid].i[ 4];
 779   w1[1] = pws[gid].i[ 5];
 780   w1[2] = pws[gid].i[ 6];
 781   w1[3] = pws[gid].i[ 7];
 782
 783   u32x w2[4];
 784
 785   w2[0] = 0;
 786   w2[1] = 0;
 787   w2[2] = 0;
 788   w2[3] = 0;
 789
 790   u32x w3[4];
 791
 792   w3[0] = 0;
 793   w3[1] = 0;
 794   w3[2] = pws[gid].i[14];
 795   w3[3] = 0;
 796
 797   const u32 pw_len = pws[gid].pw_len;
 798
 799   /**
 800    * salt
 801    */
 802
 803   __shared__ u32 s_userdomain_buf[64];
 804   __shared__ u32 s_chall_buf[256];
 805
 806   if (lid < 64)
 807   {
 808     s_userdomain_buf[lid] = netntlm_bufs[salt_pos].userdomain_buf[lid];
 809   }
 810
 811   s_chall_buf[lid] = netntlm_bufs[salt_pos].chall_buf[lid];
 812
 813   __syncthreads ();
 814
 815   if (gid >= gid_max) return;
 816
 817   /**
 818    * main
 819    */
 820
 821   m05600m (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, netntlm_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, s_userdomain_buf, s_chall_buf);
 822 }
 823
 824 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 825 {
 826 }
 827
 828 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 829 {
 830   /**
 831    * base
 832    */
 833
 834   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 835   const u32 lid = threadIdx.x;
 836
 837   u32x w0[4];
 838
 839   w0[0] = pws[gid].i[ 0];
 840   w0[1] = pws[gid].i[ 1];
 841   w0[2] = pws[gid].i[ 2];
 842   w0[3] = pws[gid].i[ 3];
 843
 844   u32x w1[4];
 845
 846   w1[0] = 0;
 847   w1[1] = 0;
 848   w1[2] = 0;
 849   w1[3] = 0;
 850
 851   u32x w2[4];
 852
 853   w2[0] = 0;
 854   w2[1] = 0;
 855   w2[2] = 0;
 856   w2[3] = 0;
 857
 858   u32x w3[4];
 859
 860   w3[0] = 0;
 861   w3[1] = 0;
 862   w3[2] = pws[gid].i[14];
 863   w3[3] = 0;
 864
 865   const u32 pw_len = pws[gid].pw_len;
 866
 867   /**
 868    * salt
 869    */
 870
 871   __shared__ u32 s_userdomain_buf[64];
 872   __shared__ u32 s_chall_buf[256];
 873
 874   if (lid < 64)
 875   {
 876     s_userdomain_buf[lid] = netntlm_bufs[salt_pos].userdomain_buf[lid];
 877   }
 878
 879   s_chall_buf[lid] = netntlm_bufs[salt_pos].chall_buf[lid];
 880
 881   __syncthreads ();
 882
 883   if (gid >= gid_max) return;
 884
 885   /**
 886    * main
 887    */
 888
 889   m05600s (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, netntlm_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, s_userdomain_buf, s_chall_buf);
 890 }
 891
 892 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 893 {
 894   /**
 895    * base
 896    */
 897
 898   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 899   const u32 lid = threadIdx.x;
 900
 901   u32x w0[4];
 902
 903   w0[0] = pws[gid].i[ 0];
 904   w0[1] = pws[gid].i[ 1];
 905   w0[2] = pws[gid].i[ 2];
 906   w0[3] = pws[gid].i[ 3];
 907
 908   u32x w1[4];
 909
 910   w1[0] = pws[gid].i[ 4];
 911   w1[1] = pws[gid].i[ 5];
 912   w1[2] = pws[gid].i[ 6];
 913   w1[3] = pws[gid].i[ 7];
 914
 915   u32x w2[4];
 916
 917   w2[0] = 0;
 918   w2[1] = 0;
 919   w2[2] = 0;
 920   w2[3] = 0;
 921
 922   u32x w3[4];
 923
 924   w3[0] = 0;
 925   w3[1] = 0;
 926   w3[2] = pws[gid].i[14];
 927   w3[3] = 0;
 928
 929   const u32 pw_len = pws[gid].pw_len;
 930
 931   /**
 932    * salt
 933    */
 934
 935   __shared__ u32 s_userdomain_buf[64];
 936   __shared__ u32 s_chall_buf[256];
 937
 938   if (lid < 64)
 939   {
 940     s_userdomain_buf[lid] = netntlm_bufs[salt_pos].userdomain_buf[lid];
 941   }
 942
 943   s_chall_buf[lid] = netntlm_bufs[salt_pos].chall_buf[lid];
 944
 945   __syncthreads ();
 946
 947   if (gid >= gid_max) return;
 948
 949   /**
 950    * main
 951    */
 952
 953   m05600s (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, netntlm_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, s_userdomain_buf, s_chall_buf);
 954 }
 955
 956 extern "C" __global__ void __launch_bounds__ (256, 1) m05600_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 netntlm_t *netntlm_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)
 957 {
 958 }