nv/m02100.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _DCC2_
   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 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   MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
  48   MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
  49   MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
  50   MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
  51   MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
  52   MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
  53   MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
  54   MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
  55   MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
  56   MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
  57   MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
  58   MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
  59   MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
  60   MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
  61   MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
  62   MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
  63
  64   MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
  65   MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
  66   MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
  67   MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
  68   MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
  69   MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
  70   MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
  71   MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
  72   MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
  73   MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
  74   MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
  75   MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
  76   MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
  77   MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
  78   MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
  79   MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
  80
  81   MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
  82   MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
  83   MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
  84   MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
  85   MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
  86   MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
  87   MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
  88   MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
  89   MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
  90   MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
  91   MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
  92   MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
  93   MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
  94   MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
  95   MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
  96   MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
  97
  98   digest[0] += a;
  99   digest[1] += b;
 100   digest[2] += c;
 101   digest[3] += d;
 102 }
 103
 104 __device__ static void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
 105 {
 106   u32x A = digest[0];
 107   u32x B = digest[1];
 108   u32x C = digest[2];
 109   u32x D = digest[3];
 110   u32x E = digest[4];
 111
 112   u32x w0_t = w0[0];
 113   u32x w1_t = w0[1];
 114   u32x w2_t = w0[2];
 115   u32x w3_t = w0[3];
 116   u32x w4_t = w1[0];
 117   u32x w5_t = w1[1];
 118   u32x w6_t = w1[2];
 119   u32x w7_t = w1[3];
 120   u32x w8_t = w2[0];
 121   u32x w9_t = w2[1];
 122   u32x wa_t = w2[2];
 123   u32x wb_t = w2[3];
 124   u32x wc_t = w3[0];
 125   u32x wd_t = w3[1];
 126   u32x we_t = w3[2];
 127   u32x wf_t = w3[3];
 128
 129   #undef K
 130   #define K SHA1C00
 131
 132   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
 133   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
 134   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
 135   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
 136   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
 137   SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
 138   SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
 139   SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
 140   SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
 141   SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
 142   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
 143   SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
 144   SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
 145   SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
 146   SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
 147   SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
 148   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
 149   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
 150   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
 151   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
 152
 153   #undef K
 154   #define K SHA1C01
 155
 156   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
 157   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
 158   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
 159   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
 160   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
 161   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
 162   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
 163   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
 164   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
 165   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
 166   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
 167   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
 168   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
 169   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
 170   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
 171   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
 172   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
 173   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
 174   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
 175   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
 176
 177   #undef K
 178   #define K SHA1C02
 179
 180   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
 181   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
 182   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
 183   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
 184   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
 185   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
 186   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
 187   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
 188   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
 189   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
 190   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
 191   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
 192   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
 193   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
 194   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
 195   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
 196   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
 197   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
 198   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
 199   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
 200
 201   #undef K
 202   #define K SHA1C03
 203
 204   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
 205   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
 206   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
 207   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
 208   w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
 209   w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
 210   w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
 211   w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
 212   w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
 213   w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
 214   w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
 215   w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
 216   w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
 217   w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
 218   wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
 219   wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
 220   wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
 221   wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
 222   we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
 223   wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
 224
 225   digest[0] += A;
 226   digest[1] += B;
 227   digest[2] += C;
 228   digest[3] += D;
 229   digest[4] += E;
 230 }
 231
 232 __device__ static void hmac_sha1_pad (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5])
 233 {
 234   w0[0] = w0[0] ^ 0x36363636;
 235   w0[1] = w0[1] ^ 0x36363636;
 236   w0[2] = w0[2] ^ 0x36363636;
 237   w0[3] = w0[3] ^ 0x36363636;
 238   w1[0] = w1[0] ^ 0x36363636;
 239   w1[1] = w1[1] ^ 0x36363636;
 240   w1[2] = w1[2] ^ 0x36363636;
 241   w1[3] = w1[3] ^ 0x36363636;
 242   w2[0] = w2[0] ^ 0x36363636;
 243   w2[1] = w2[1] ^ 0x36363636;
 244   w2[2] = w2[2] ^ 0x36363636;
 245   w2[3] = w2[3] ^ 0x36363636;
 246   w3[0] = w3[0] ^ 0x36363636;
 247   w3[1] = w3[1] ^ 0x36363636;
 248   w3[2] = w3[2] ^ 0x36363636;
 249   w3[3] = w3[3] ^ 0x36363636;
 250
 251   ipad[0] = SHA1M_A;
 252   ipad[1] = SHA1M_B;
 253   ipad[2] = SHA1M_C;
 254   ipad[3] = SHA1M_D;
 255   ipad[4] = SHA1M_E;
 256
 257   sha1_transform (w0, w1, w2, w3, ipad);
 258
 259   w0[0] = w0[0] ^ 0x6a6a6a6a;
 260   w0[1] = w0[1] ^ 0x6a6a6a6a;
 261   w0[2] = w0[2] ^ 0x6a6a6a6a;
 262   w0[3] = w0[3] ^ 0x6a6a6a6a;
 263   w1[0] = w1[0] ^ 0x6a6a6a6a;
 264   w1[1] = w1[1] ^ 0x6a6a6a6a;
 265   w1[2] = w1[2] ^ 0x6a6a6a6a;
 266   w1[3] = w1[3] ^ 0x6a6a6a6a;
 267   w2[0] = w2[0] ^ 0x6a6a6a6a;
 268   w2[1] = w2[1] ^ 0x6a6a6a6a;
 269   w2[2] = w2[2] ^ 0x6a6a6a6a;
 270   w2[3] = w2[3] ^ 0x6a6a6a6a;
 271   w3[0] = w3[0] ^ 0x6a6a6a6a;
 272   w3[1] = w3[1] ^ 0x6a6a6a6a;
 273   w3[2] = w3[2] ^ 0x6a6a6a6a;
 274   w3[3] = w3[3] ^ 0x6a6a6a6a;
 275
 276   opad[0] = SHA1M_A;
 277   opad[1] = SHA1M_B;
 278   opad[2] = SHA1M_C;
 279   opad[3] = SHA1M_D;
 280   opad[4] = SHA1M_E;
 281
 282   sha1_transform (w0, w1, w2, w3, opad);
 283 }
 284
 285 __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])
 286 {
 287   digest[0] = ipad[0];
 288   digest[1] = ipad[1];
 289   digest[2] = ipad[2];
 290   digest[3] = ipad[3];
 291   digest[4] = ipad[4];
 292
 293   sha1_transform (w0, w1, w2, w3, digest);
 294
 295   w0[0] = digest[0];
 296   w0[1] = digest[1];
 297   w0[2] = digest[2];
 298   w0[3] = digest[3];
 299   w1[0] = digest[4];
 300   w1[1] = 0x80000000;
 301   w1[2] = 0;
 302   w1[3] = 0;
 303   w2[0] = 0;
 304   w2[1] = 0;
 305   w2[2] = 0;
 306   w2[3] = 0;
 307   w3[0] = 0;
 308   w3[1] = 0;
 309   w3[2] = 0;
 310   w3[3] = (64 + 20) * 8;
 311
 312   digest[0] = opad[0];
 313   digest[1] = opad[1];
 314   digest[2] = opad[2];
 315   digest[3] = opad[3];
 316   digest[4] = opad[4];
 317
 318   sha1_transform (w0, w1, w2, w3, digest);
 319 }
 320
 321 extern "C" __global__ void __launch_bounds__ (256, 1) m02100_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, dcc2_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 void *esalt_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)
 322 {
 323   /**
 324    * base
 325    */
 326
 327   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 328
 329   if (gid >= gid_max) return;
 330
 331   u32x w0[4];
 332
 333   w0[0] = pws[gid].i[ 0];
 334   w0[1] = pws[gid].i[ 1];
 335   w0[2] = pws[gid].i[ 2];
 336   w0[3] = pws[gid].i[ 3];
 337
 338   u32x w1[4];
 339
 340   w1[0] = 0;
 341   w1[1] = 0;
 342   w1[2] = 0;
 343   w1[3] = 0;
 344
 345   u32x w2[4];
 346
 347   w2[0] = 0;
 348   w2[1] = 0;
 349   w2[2] = 0;
 350   w2[3] = 0;
 351
 352   u32x w3[4];
 353
 354   w3[0] = 0;
 355   w3[1] = 0;
 356   w3[2] = 0;
 357   w3[3] = 0;
 358
 359   const u32 pw_len = pws[gid].pw_len;
 360
 361   /**
 362    * salt
 363    */
 364
 365   u32 salt_len = salt_bufs[salt_pos].salt_len;
 366
 367   u32 salt_buf0[4];
 368   u32 salt_buf1[4];
 369   u32 salt_buf2[4];
 370
 371   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
 372   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
 373   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
 374   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
 375   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
 376   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
 377   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
 378   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
 379   salt_buf2[0] = salt_bufs[salt_pos].salt_buf[8];
 380   salt_buf2[1] = salt_bufs[salt_pos].salt_buf[9];
 381   salt_buf2[2] = 0;
 382   salt_buf2[3] = 0;
 383
 384   /**
 385    * generate dcc
 386    */
 387
 388   append_0x80_1 (w0, pw_len);
 389
 390   make_unicode (w0, w0, w1);
 391
 392   w3[2] = pw_len * 2 * 8;
 393
 394   u32x digest_md4[4];
 395
 396   digest_md4[0] = MD4M_A;
 397   digest_md4[1] = MD4M_B;
 398   digest_md4[2] = MD4M_C;
 399   digest_md4[3] = MD4M_D;
 400
 401   md4_transform (w0, w1, w2, w3, digest_md4);
 402
 403   w0[0] = digest_md4[0];
 404   w0[1] = digest_md4[1];
 405   w0[2] = digest_md4[2];
 406   w0[3] = digest_md4[3];
 407   w1[0] = salt_buf0[0];
 408   w1[1] = salt_buf0[1];
 409   w1[2] = salt_buf0[2];
 410   w1[3] = salt_buf0[3];
 411   w2[0] = salt_buf1[0];
 412   w2[1] = salt_buf1[1];
 413   w2[2] = salt_buf1[2];
 414   w2[3] = salt_buf1[3];
 415   w3[0] = salt_buf2[0];
 416   w3[1] = salt_buf2[1];
 417   w3[2] = (16 + salt_len) * 8;
 418   w3[3] = 0;
 419
 420   append_0x80_4 (w0, w1, w2, w3, 16 + salt_len);
 421
 422   digest_md4[0] = MD4M_A;
 423   digest_md4[1] = MD4M_B;
 424   digest_md4[2] = MD4M_C;
 425   digest_md4[3] = MD4M_D;
 426
 427   md4_transform (w0, w1, w2, w3, digest_md4);
 428
 429   /**
 430    * pads
 431    */
 432
 433   w0[0] = swap_workaround (digest_md4[0]);
 434   w0[1] = swap_workaround (digest_md4[1]);
 435   w0[2] = swap_workaround (digest_md4[2]);
 436   w0[3] = swap_workaround (digest_md4[3]);
 437   w1[0] = 0;
 438   w1[1] = 0;
 439   w1[2] = 0;
 440   w1[3] = 0;
 441   w2[0] = 0;
 442   w2[1] = 0;
 443   w2[2] = 0;
 444   w2[3] = 0;
 445   w3[0] = 0;
 446   w3[1] = 0;
 447   w3[2] = 0;
 448   w3[3] = 0;
 449
 450   u32x ipad[5];
 451   u32x opad[5];
 452
 453   hmac_sha1_pad (w0, w1, w2, w3, ipad, opad);
 454
 455   tmps[gid].ipad[0] = ipad[0];
 456   tmps[gid].ipad[1] = ipad[1];
 457   tmps[gid].ipad[2] = ipad[2];
 458   tmps[gid].ipad[3] = ipad[3];
 459   tmps[gid].ipad[4] = ipad[4];
 460
 461   tmps[gid].opad[0] = opad[0];
 462   tmps[gid].opad[1] = opad[1];
 463   tmps[gid].opad[2] = opad[2];
 464   tmps[gid].opad[3] = opad[3];
 465   tmps[gid].opad[4] = opad[4];
 466
 467   /**
 468    * hmac1
 469    */
 470
 471   w0[0] = salt_buf0[0];
 472   w0[1] = salt_buf0[1];
 473   w0[2] = salt_buf0[2];
 474   w0[3] = salt_buf0[3];
 475   w1[0] = salt_buf1[0];
 476   w1[1] = salt_buf1[1];
 477   w1[2] = salt_buf1[2];
 478   w1[3] = salt_buf1[3];
 479   w2[0] = salt_buf2[0];
 480   w2[1] = salt_buf2[1];
 481   w2[2] = 0;
 482   w2[3] = 0;
 483   w3[0] = 0;
 484   w3[1] = 0;
 485   w3[2] = 0;
 486   w3[3] = (64 + salt_len + 4) * 8;
 487
 488   append_0x01_4 (w0, w1, w2, w3, salt_len + 3);
 489   append_0x80_4 (w0, w1, w2, w3, salt_len + 4);
 490
 491   w0[0] = swap_workaround (w0[0]);
 492   w0[1] = swap_workaround (w0[1]);
 493   w0[2] = swap_workaround (w0[2]);
 494   w0[3] = swap_workaround (w0[3]);
 495   w1[0] = swap_workaround (w1[0]);
 496   w1[1] = swap_workaround (w1[1]);
 497   w1[2] = swap_workaround (w1[2]);
 498   w1[3] = swap_workaround (w1[3]);
 499   w2[0] = swap_workaround (w2[0]);
 500   w2[1] = swap_workaround (w2[1]);
 501   w2[2] = swap_workaround (w2[2]);
 502   w2[3] = swap_workaround (w2[3]);
 503   w3[0] = swap_workaround (w3[0]);
 504   w3[1] = swap_workaround (w3[1]);
 505
 506   u32x digest[5];
 507
 508   hmac_sha1_run (w0, w1, w2, w3, ipad, opad, digest);
 509
 510   tmps[gid].dgst[0] = digest[0];
 511   tmps[gid].dgst[1] = digest[1];
 512   tmps[gid].dgst[2] = digest[2];
 513   tmps[gid].dgst[3] = digest[3];
 514   tmps[gid].dgst[4] = digest[4];
 515
 516   tmps[gid].out[0] = digest[0];
 517   tmps[gid].out[1] = digest[1];
 518   tmps[gid].out[2] = digest[2];
 519   tmps[gid].out[3] = digest[3];
 520 }
 521
 522 extern "C" __global__ void __launch_bounds__ (256, 1) m02100_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, dcc2_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 void *esalt_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)
 523 {
 524   /**
 525    * base
 526    */
 527
 528   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 529
 530   if (gid >= gid_max) return;
 531
 532   u32x ipad[5];
 533   u32x opad[5];
 534
 535   ipad[0] = tmps[gid].ipad[0];
 536   ipad[1] = tmps[gid].ipad[1];
 537   ipad[2] = tmps[gid].ipad[2];
 538   ipad[3] = tmps[gid].ipad[3];
 539   ipad[4] = tmps[gid].ipad[4];
 540
 541   opad[0] = tmps[gid].opad[0];
 542   opad[1] = tmps[gid].opad[1];
 543   opad[2] = tmps[gid].opad[2];
 544   opad[3] = tmps[gid].opad[3];
 545   opad[4] = tmps[gid].opad[4];
 546
 547   /**
 548    * iter1
 549    */
 550
 551   u32x dgst[5];
 552   u32x out[4];
 553
 554   dgst[0] = tmps[gid].dgst[0];
 555   dgst[1] = tmps[gid].dgst[1];
 556   dgst[2] = tmps[gid].dgst[2];
 557   dgst[3] = tmps[gid].dgst[3];
 558   dgst[4] = tmps[gid].dgst[4];
 559
 560   out[0] = tmps[gid].out[0];
 561   out[1] = tmps[gid].out[1];
 562   out[2] = tmps[gid].out[2];
 563   out[3] = tmps[gid].out[3];
 564
 565   for (u32 i = 0; i < loop_cnt; i++)
 566   {
 567     u32x w0[4];
 568     u32x w1[4];
 569     u32x w2[4];
 570     u32x w3[4];
 571
 572     w0[0] = dgst[0];
 573     w0[1] = dgst[1];
 574     w0[2] = dgst[2];
 575     w0[3] = dgst[3];
 576     w1[0] = dgst[4];
 577     w1[1] = 0x80000000;
 578     w1[2] = 0;
 579     w1[3] = 0;
 580     w2[0] = 0;
 581     w2[1] = 0;
 582     w2[2] = 0;
 583     w2[3] = 0;
 584     w3[0] = 0;
 585     w3[1] = 0;
 586     w3[2] = 0;
 587     w3[3] = (64 + 20) * 8;
 588
 589     hmac_sha1_run (w0, w1, w2, w3, ipad, opad, dgst);
 590
 591     out[0] ^= dgst[0];
 592     out[1] ^= dgst[1];
 593     out[2] ^= dgst[2];
 594     out[3] ^= dgst[3];
 595   }
 596
 597   tmps[gid].dgst[0] = dgst[0];
 598   tmps[gid].dgst[1] = dgst[1];
 599   tmps[gid].dgst[2] = dgst[2];
 600   tmps[gid].dgst[3] = dgst[3];
 601   tmps[gid].dgst[4] = dgst[4];
 602
 603   tmps[gid].out[0] = out[0];
 604   tmps[gid].out[1] = out[1];
 605   tmps[gid].out[2] = out[2];
 606   tmps[gid].out[3] = out[3];
 607 }
 608
 609 extern "C" __global__ void __launch_bounds__ (256, 1) m02100_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, dcc2_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 void *esalt_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)
 610 {
 611   /**
 612    * modifier
 613    */
 614
 615   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 616
 617   if (gid >= gid_max) return;
 618
 619   const u32 lid = threadIdx.x;
 620
 621   const u32x r0 = tmps[gid].out[DGST_R0];
 622   const u32x r1 = tmps[gid].out[DGST_R1];
 623   const u32x r2 = tmps[gid].out[DGST_R2];
 624   const u32x r3 = tmps[gid].out[DGST_R3];
 625
 626   #define il_pos 0
 627
 628   #include VECT_COMPARE_M
 629 }