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