nv/m02500.cu

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