nv/m05500_a1.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _MD4_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  30 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  31 #endif
  32
  33 #ifdef  VECT_SIZE2
  34 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  35 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  36 #endif
  37
  38 #define PERM_OP(a,b,tt,n,m) \
  39 {                           \
  40   tt = a >> n;              \
  41   tt = tt ^ b;              \
  42   tt = tt & m;              \
  43   b = b ^ tt;               \
  44   tt = tt << n;             \
  45   a = a ^ tt;               \
  46 }
  47
  48 #define HPERM_OP(a,tt,n,m)  \
  49 {                           \
  50   tt = a << (16 + n);       \
  51   tt = tt ^ a;              \
  52   tt = tt & m;              \
  53   a  = a ^ tt;              \
  54   tt = tt >> (16 + n);      \
  55   a  = a ^ tt;              \
  56 }
  57
  58 __device__ __constant__ u32 c_SPtrans[8][64] =
  59 {
  60   /* nibble 0 */
  61   0x02080800, 0x00080000, 0x02000002, 0x02080802,
  62   0x02000000, 0x00080802, 0x00080002, 0x02000002,
  63   0x00080802, 0x02080800, 0x02080000, 0x00000802,
  64   0x02000802, 0x02000000, 0x00000000, 0x00080002,
  65   0x00080000, 0x00000002, 0x02000800, 0x00080800,
  66   0x02080802, 0x02080000, 0x00000802, 0x02000800,
  67   0x00000002, 0x00000800, 0x00080800, 0x02080002,
  68   0x00000800, 0x02000802, 0x02080002, 0x00000000,
  69   0x00000000, 0x02080802, 0x02000800, 0x00080002,
  70   0x02080800, 0x00080000, 0x00000802, 0x02000800,
  71   0x02080002, 0x00000800, 0x00080800, 0x02000002,
  72   0x00080802, 0x00000002, 0x02000002, 0x02080000,
  73   0x02080802, 0x00080800, 0x02080000, 0x02000802,
  74   0x02000000, 0x00000802, 0x00080002, 0x00000000,
  75   0x00080000, 0x02000000, 0x02000802, 0x02080800,
  76   0x00000002, 0x02080002, 0x00000800, 0x00080802,
  77   /* nibble 1 */
  78   0x40108010, 0x00000000, 0x00108000, 0x40100000,
  79   0x40000010, 0x00008010, 0x40008000, 0x00108000,
  80   0x00008000, 0x40100010, 0x00000010, 0x40008000,
  81   0x00100010, 0x40108000, 0x40100000, 0x00000010,
  82   0x00100000, 0x40008010, 0x40100010, 0x00008000,
  83   0x00108010, 0x40000000, 0x00000000, 0x00100010,
  84   0x40008010, 0x00108010, 0x40108000, 0x40000010,
  85   0x40000000, 0x00100000, 0x00008010, 0x40108010,
  86   0x00100010, 0x40108000, 0x40008000, 0x00108010,
  87   0x40108010, 0x00100010, 0x40000010, 0x00000000,
  88   0x40000000, 0x00008010, 0x00100000, 0x40100010,
  89   0x00008000, 0x40000000, 0x00108010, 0x40008010,
  90   0x40108000, 0x00008000, 0x00000000, 0x40000010,
  91   0x00000010, 0x40108010, 0x00108000, 0x40100000,
  92   0x40100010, 0x00100000, 0x00008010, 0x40008000,
  93   0x40008010, 0x00000010, 0x40100000, 0x00108000,
  94   /* nibble 2 */
  95   0x04000001, 0x04040100, 0x00000100, 0x04000101,
  96   0x00040001, 0x04000000, 0x04000101, 0x00040100,
  97   0x04000100, 0x00040000, 0x04040000, 0x00000001,
  98   0x04040101, 0x00000101, 0x00000001, 0x04040001,
  99   0x00000000, 0x00040001, 0x04040100, 0x00000100,
 100   0x00000101, 0x04040101, 0x00040000, 0x04000001,
 101   0x04040001, 0x04000100, 0x00040101, 0x04040000,
 102   0x00040100, 0x00000000, 0x04000000, 0x00040101,
 103   0x04040100, 0x00000100, 0x00000001, 0x00040000,
 104   0x00000101, 0x00040001, 0x04040000, 0x04000101,
 105   0x00000000, 0x04040100, 0x00040100, 0x04040001,
 106   0x00040001, 0x04000000, 0x04040101, 0x00000001,
 107   0x00040101, 0x04000001, 0x04000000, 0x04040101,
 108   0x00040000, 0x04000100, 0x04000101, 0x00040100,
 109   0x04000100, 0x00000000, 0x04040001, 0x00000101,
 110   0x04000001, 0x00040101, 0x00000100, 0x04040000,
 111   /* nibble 3 */
 112   0x00401008, 0x10001000, 0x00000008, 0x10401008,
 113   0x00000000, 0x10400000, 0x10001008, 0x00400008,
 114   0x10401000, 0x10000008, 0x10000000, 0x00001008,
 115   0x10000008, 0x00401008, 0x00400000, 0x10000000,
 116   0x10400008, 0x00401000, 0x00001000, 0x00000008,
 117   0x00401000, 0x10001008, 0x10400000, 0x00001000,
 118   0x00001008, 0x00000000, 0x00400008, 0x10401000,
 119   0x10001000, 0x10400008, 0x10401008, 0x00400000,
 120   0x10400008, 0x00001008, 0x00400000, 0x10000008,
 121   0x00401000, 0x10001000, 0x00000008, 0x10400000,
 122   0x10001008, 0x00000000, 0x00001000, 0x00400008,
 123   0x00000000, 0x10400008, 0x10401000, 0x00001000,
 124   0x10000000, 0x10401008, 0x00401008, 0x00400000,
 125   0x10401008, 0x00000008, 0x10001000, 0x00401008,
 126   0x00400008, 0x00401000, 0x10400000, 0x10001008,
 127   0x00001008, 0x10000000, 0x10000008, 0x10401000,
 128   /* nibble 4 */
 129   0x08000000, 0x00010000, 0x00000400, 0x08010420,
 130   0x08010020, 0x08000400, 0x00010420, 0x08010000,
 131   0x00010000, 0x00000020, 0x08000020, 0x00010400,
 132   0x08000420, 0x08010020, 0x08010400, 0x00000000,
 133   0x00010400, 0x08000000, 0x00010020, 0x00000420,
 134   0x08000400, 0x00010420, 0x00000000, 0x08000020,
 135   0x00000020, 0x08000420, 0x08010420, 0x00010020,
 136   0x08010000, 0x00000400, 0x00000420, 0x08010400,
 137   0x08010400, 0x08000420, 0x00010020, 0x08010000,
 138   0x00010000, 0x00000020, 0x08000020, 0x08000400,
 139   0x08000000, 0x00010400, 0x08010420, 0x00000000,
 140   0x00010420, 0x08000000, 0x00000400, 0x00010020,
 141   0x08000420, 0x00000400, 0x00000000, 0x08010420,
 142   0x08010020, 0x08010400, 0x00000420, 0x00010000,
 143   0x00010400, 0x08010020, 0x08000400, 0x00000420,
 144   0x00000020, 0x00010420, 0x08010000, 0x08000020,
 145   /* nibble 5 */
 146   0x80000040, 0x00200040, 0x00000000, 0x80202000,
 147   0x00200040, 0x00002000, 0x80002040, 0x00200000,
 148   0x00002040, 0x80202040, 0x00202000, 0x80000000,
 149   0x80002000, 0x80000040, 0x80200000, 0x00202040,
 150   0x00200000, 0x80002040, 0x80200040, 0x00000000,
 151   0x00002000, 0x00000040, 0x80202000, 0x80200040,
 152   0x80202040, 0x80200000, 0x80000000, 0x00002040,
 153   0x00000040, 0x00202000, 0x00202040, 0x80002000,
 154   0x00002040, 0x80000000, 0x80002000, 0x00202040,
 155   0x80202000, 0x00200040, 0x00000000, 0x80002000,
 156   0x80000000, 0x00002000, 0x80200040, 0x00200000,
 157   0x00200040, 0x80202040, 0x00202000, 0x00000040,
 158   0x80202040, 0x00202000, 0x00200000, 0x80002040,
 159   0x80000040, 0x80200000, 0x00202040, 0x00000000,
 160   0x00002000, 0x80000040, 0x80002040, 0x80202000,
 161   0x80200000, 0x00002040, 0x00000040, 0x80200040,
 162   /* nibble 6 */
 163   0x00004000, 0x00000200, 0x01000200, 0x01000004,
 164   0x01004204, 0x00004004, 0x00004200, 0x00000000,
 165   0x01000000, 0x01000204, 0x00000204, 0x01004000,
 166   0x00000004, 0x01004200, 0x01004000, 0x00000204,
 167   0x01000204, 0x00004000, 0x00004004, 0x01004204,
 168   0x00000000, 0x01000200, 0x01000004, 0x00004200,
 169   0x01004004, 0x00004204, 0x01004200, 0x00000004,
 170   0x00004204, 0x01004004, 0x00000200, 0x01000000,
 171   0x00004204, 0x01004000, 0x01004004, 0x00000204,
 172   0x00004000, 0x00000200, 0x01000000, 0x01004004,
 173   0x01000204, 0x00004204, 0x00004200, 0x00000000,
 174   0x00000200, 0x01000004, 0x00000004, 0x01000200,
 175   0x00000000, 0x01000204, 0x01000200, 0x00004200,
 176   0x00000204, 0x00004000, 0x01004204, 0x01000000,
 177   0x01004200, 0x00000004, 0x00004004, 0x01004204,
 178   0x01000004, 0x01004200, 0x01004000, 0x00004004,
 179   /* nibble 7 */
 180   0x20800080, 0x20820000, 0x00020080, 0x00000000,
 181   0x20020000, 0x00800080, 0x20800000, 0x20820080,
 182   0x00000080, 0x20000000, 0x00820000, 0x00020080,
 183   0x00820080, 0x20020080, 0x20000080, 0x20800000,
 184   0x00020000, 0x00820080, 0x00800080, 0x20020000,
 185   0x20820080, 0x20000080, 0x00000000, 0x00820000,
 186   0x20000000, 0x00800000, 0x20020080, 0x20800080,
 187   0x00800000, 0x00020000, 0x20820000, 0x00000080,
 188   0x00800000, 0x00020000, 0x20000080, 0x20820080,
 189   0x00020080, 0x20000000, 0x00000000, 0x00820000,
 190   0x20800080, 0x20020080, 0x20020000, 0x00800080,
 191   0x20820000, 0x00000080, 0x00800080, 0x20020000,
 192   0x20820080, 0x00800000, 0x20800000, 0x20000080,
 193   0x00820000, 0x00020080, 0x20020080, 0x20800000,
 194   0x00000080, 0x20820000, 0x00820080, 0x00000000,
 195   0x20000000, 0x20800080, 0x00020000, 0x00820080,
 196 };
 197
 198 __device__ __constant__ u32 c_skb[8][64] =
 199 {
 200   /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
 201   0x00000000, 0x00000010, 0x20000000, 0x20000010,
 202   0x00010000, 0x00010010, 0x20010000, 0x20010010,
 203   0x00000800, 0x00000810, 0x20000800, 0x20000810,
 204   0x00010800, 0x00010810, 0x20010800, 0x20010810,
 205   0x00000020, 0x00000030, 0x20000020, 0x20000030,
 206   0x00010020, 0x00010030, 0x20010020, 0x20010030,
 207   0x00000820, 0x00000830, 0x20000820, 0x20000830,
 208   0x00010820, 0x00010830, 0x20010820, 0x20010830,
 209   0x00080000, 0x00080010, 0x20080000, 0x20080010,
 210   0x00090000, 0x00090010, 0x20090000, 0x20090010,
 211   0x00080800, 0x00080810, 0x20080800, 0x20080810,
 212   0x00090800, 0x00090810, 0x20090800, 0x20090810,
 213   0x00080020, 0x00080030, 0x20080020, 0x20080030,
 214   0x00090020, 0x00090030, 0x20090020, 0x20090030,
 215   0x00080820, 0x00080830, 0x20080820, 0x20080830,
 216   0x00090820, 0x00090830, 0x20090820, 0x20090830,
 217   /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
 218   0x00000000, 0x02000000, 0x00002000, 0x02002000,
 219   0x00200000, 0x02200000, 0x00202000, 0x02202000,
 220   0x00000004, 0x02000004, 0x00002004, 0x02002004,
 221   0x00200004, 0x02200004, 0x00202004, 0x02202004,
 222   0x00000400, 0x02000400, 0x00002400, 0x02002400,
 223   0x00200400, 0x02200400, 0x00202400, 0x02202400,
 224   0x00000404, 0x02000404, 0x00002404, 0x02002404,
 225   0x00200404, 0x02200404, 0x00202404, 0x02202404,
 226   0x10000000, 0x12000000, 0x10002000, 0x12002000,
 227   0x10200000, 0x12200000, 0x10202000, 0x12202000,
 228   0x10000004, 0x12000004, 0x10002004, 0x12002004,
 229   0x10200004, 0x12200004, 0x10202004, 0x12202004,
 230   0x10000400, 0x12000400, 0x10002400, 0x12002400,
 231   0x10200400, 0x12200400, 0x10202400, 0x12202400,
 232   0x10000404, 0x12000404, 0x10002404, 0x12002404,
 233   0x10200404, 0x12200404, 0x10202404, 0x12202404,
 234   /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
 235   0x00000000, 0x00000001, 0x00040000, 0x00040001,
 236   0x01000000, 0x01000001, 0x01040000, 0x01040001,
 237   0x00000002, 0x00000003, 0x00040002, 0x00040003,
 238   0x01000002, 0x01000003, 0x01040002, 0x01040003,
 239   0x00000200, 0x00000201, 0x00040200, 0x00040201,
 240   0x01000200, 0x01000201, 0x01040200, 0x01040201,
 241   0x00000202, 0x00000203, 0x00040202, 0x00040203,
 242   0x01000202, 0x01000203, 0x01040202, 0x01040203,
 243   0x08000000, 0x08000001, 0x08040000, 0x08040001,
 244   0x09000000, 0x09000001, 0x09040000, 0x09040001,
 245   0x08000002, 0x08000003, 0x08040002, 0x08040003,
 246   0x09000002, 0x09000003, 0x09040002, 0x09040003,
 247   0x08000200, 0x08000201, 0x08040200, 0x08040201,
 248   0x09000200, 0x09000201, 0x09040200, 0x09040201,
 249   0x08000202, 0x08000203, 0x08040202, 0x08040203,
 250   0x09000202, 0x09000203, 0x09040202, 0x09040203,
 251   /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
 252   0x00000000, 0x00100000, 0x00000100, 0x00100100,
 253   0x00000008, 0x00100008, 0x00000108, 0x00100108,
 254   0x00001000, 0x00101000, 0x00001100, 0x00101100,
 255   0x00001008, 0x00101008, 0x00001108, 0x00101108,
 256   0x04000000, 0x04100000, 0x04000100, 0x04100100,
 257   0x04000008, 0x04100008, 0x04000108, 0x04100108,
 258   0x04001000, 0x04101000, 0x04001100, 0x04101100,
 259   0x04001008, 0x04101008, 0x04001108, 0x04101108,
 260   0x00020000, 0x00120000, 0x00020100, 0x00120100,
 261   0x00020008, 0x00120008, 0x00020108, 0x00120108,
 262   0x00021000, 0x00121000, 0x00021100, 0x00121100,
 263   0x00021008, 0x00121008, 0x00021108, 0x00121108,
 264   0x04020000, 0x04120000, 0x04020100, 0x04120100,
 265   0x04020008, 0x04120008, 0x04020108, 0x04120108,
 266   0x04021000, 0x04121000, 0x04021100, 0x04121100,
 267   0x04021008, 0x04121008, 0x04021108, 0x04121108,
 268   /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
 269   0x00000000, 0x10000000, 0x00010000, 0x10010000,
 270   0x00000004, 0x10000004, 0x00010004, 0x10010004,
 271   0x20000000, 0x30000000, 0x20010000, 0x30010000,
 272   0x20000004, 0x30000004, 0x20010004, 0x30010004,
 273   0x00100000, 0x10100000, 0x00110000, 0x10110000,
 274   0x00100004, 0x10100004, 0x00110004, 0x10110004,
 275   0x20100000, 0x30100000, 0x20110000, 0x30110000,
 276   0x20100004, 0x30100004, 0x20110004, 0x30110004,
 277   0x00001000, 0x10001000, 0x00011000, 0x10011000,
 278   0x00001004, 0x10001004, 0x00011004, 0x10011004,
 279   0x20001000, 0x30001000, 0x20011000, 0x30011000,
 280   0x20001004, 0x30001004, 0x20011004, 0x30011004,
 281   0x00101000, 0x10101000, 0x00111000, 0x10111000,
 282   0x00101004, 0x10101004, 0x00111004, 0x10111004,
 283   0x20101000, 0x30101000, 0x20111000, 0x30111000,
 284   0x20101004, 0x30101004, 0x20111004, 0x30111004,
 285   /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
 286   0x00000000, 0x08000000, 0x00000008, 0x08000008,
 287   0x00000400, 0x08000400, 0x00000408, 0x08000408,
 288   0x00020000, 0x08020000, 0x00020008, 0x08020008,
 289   0x00020400, 0x08020400, 0x00020408, 0x08020408,
 290   0x00000001, 0x08000001, 0x00000009, 0x08000009,
 291   0x00000401, 0x08000401, 0x00000409, 0x08000409,
 292   0x00020001, 0x08020001, 0x00020009, 0x08020009,
 293   0x00020401, 0x08020401, 0x00020409, 0x08020409,
 294   0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
 295   0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
 296   0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
 297   0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
 298   0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
 299   0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
 300   0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
 301   0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
 302   /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
 303   0x00000000, 0x00000100, 0x00080000, 0x00080100,
 304   0x01000000, 0x01000100, 0x01080000, 0x01080100,
 305   0x00000010, 0x00000110, 0x00080010, 0x00080110,
 306   0x01000010, 0x01000110, 0x01080010, 0x01080110,
 307   0x00200000, 0x00200100, 0x00280000, 0x00280100,
 308   0x01200000, 0x01200100, 0x01280000, 0x01280100,
 309   0x00200010, 0x00200110, 0x00280010, 0x00280110,
 310   0x01200010, 0x01200110, 0x01280010, 0x01280110,
 311   0x00000200, 0x00000300, 0x00080200, 0x00080300,
 312   0x01000200, 0x01000300, 0x01080200, 0x01080300,
 313   0x00000210, 0x00000310, 0x00080210, 0x00080310,
 314   0x01000210, 0x01000310, 0x01080210, 0x01080310,
 315   0x00200200, 0x00200300, 0x00280200, 0x00280300,
 316   0x01200200, 0x01200300, 0x01280200, 0x01280300,
 317   0x00200210, 0x00200310, 0x00280210, 0x00280310,
 318   0x01200210, 0x01200310, 0x01280210, 0x01280310,
 319   /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
 320   0x00000000, 0x04000000, 0x00040000, 0x04040000,
 321   0x00000002, 0x04000002, 0x00040002, 0x04040002,
 322   0x00002000, 0x04002000, 0x00042000, 0x04042000,
 323   0x00002002, 0x04002002, 0x00042002, 0x04042002,
 324   0x00000020, 0x04000020, 0x00040020, 0x04040020,
 325   0x00000022, 0x04000022, 0x00040022, 0x04040022,
 326   0x00002020, 0x04002020, 0x00042020, 0x04042020,
 327   0x00002022, 0x04002022, 0x00042022, 0x04042022,
 328   0x00000800, 0x04000800, 0x00040800, 0x04040800,
 329   0x00000802, 0x04000802, 0x00040802, 0x04040802,
 330   0x00002800, 0x04002800, 0x00042800, 0x04042800,
 331   0x00002802, 0x04002802, 0x00042802, 0x04042802,
 332   0x00000820, 0x04000820, 0x00040820, 0x04040820,
 333   0x00000822, 0x04000822, 0x00040822, 0x04040822,
 334   0x00002820, 0x04002820, 0x00042820, 0x04042820,
 335   0x00002822, 0x04002822, 0x00042822, 0x04042822
 336 };
 337
 338 #ifdef VECT_SIZE1
 339 #define BOX(i,n,S) u32x ((S)[(n)][(i)])
 340 #endif
 341
 342 #ifdef VECT_SIZE2
 343 #define BOX(i,n,S) u32x ((S)[(n)][(i).x], (S)[(n)][(i).y])
 344 #endif
 345
 346 __device__ static void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], u32 s_SPtrans[8][64])
 347 {
 348   u32x tt;
 349
 350   u32x r = data[0];
 351   u32x l = data[1];
 352
 353   #pragma unroll 16
 354   for (u32 i = 0; i < 16; i++)
 355   {
 356     u32x u = Kc[i] ^ r;
 357     u32x t = Kd[i] ^ rotl32 (r, 28u);
 358
 359     l ^= BOX (((u >>  2) & 0x3f), 0, s_SPtrans)
 360        | BOX (((u >> 10) & 0x3f), 2, s_SPtrans)
 361        | BOX (((u >> 18) & 0x3f), 4, s_SPtrans)
 362        | BOX (((u >> 26) & 0x3f), 6, s_SPtrans)
 363        | BOX (((t >>  2) & 0x3f), 1, s_SPtrans)
 364        | BOX (((t >> 10) & 0x3f), 3, s_SPtrans)
 365        | BOX (((t >> 18) & 0x3f), 5, s_SPtrans)
 366        | BOX (((t >> 26) & 0x3f), 7, s_SPtrans);
 367
 368     tt = l;
 369     l  = r;
 370     r  = tt;
 371   }
 372
 373   iv[0] = l;
 374   iv[1] = r;
 375 }
 376
 377 __device__ static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], u32 s_skb[8][64])
 378 {
 379   u32x tt;
 380
 381   PERM_OP  (d, c, tt, 4, 0x0f0f0f0f);
 382   HPERM_OP (c,    tt, 2, 0xcccc0000);
 383   HPERM_OP (d,    tt, 2, 0xcccc0000);
 384   PERM_OP  (d, c, tt, 1, 0x55555555);
 385   PERM_OP  (c, d, tt, 8, 0x00ff00ff);
 386   PERM_OP  (d, c, tt, 1, 0x55555555);
 387
 388   d = ((d & 0x000000ff) << 16)
 389     | ((d & 0x0000ff00) <<  0)
 390     | ((d & 0x00ff0000) >> 16)
 391     | ((c & 0xf0000000) >>  4);
 392
 393   c = c & 0x0fffffff;
 394
 395   #pragma unroll 16
 396   for (u32 i = 0; i < 16; i++)
 397   {
 398     const u32 shifts3s0[16] = {  1,  1,  2,  2,  2,  2,  2,  2,  1,  2,  2,  2,  2,  2,  2,  1 };
 399     const u32 shifts3s1[16] = { 27, 27, 26, 26, 26, 26, 26, 26, 27, 26, 26, 26, 26, 26, 26, 27 };
 400
 401     c = c >> shifts3s0[i] | c << shifts3s1[i];
 402     d = d >> shifts3s0[i] | d << shifts3s1[i];
 403
 404     c = c & 0x0fffffff;
 405     d = d & 0x0fffffff;
 406
 407     u32x s = BOX ((( c >>  0) & 0x3f),  0, s_skb)
 408             | BOX ((((c >>  6) & 0x03)
 409                   | ((c >>  7) & 0x3c)), 1, s_skb)
 410             | BOX ((((c >> 13) & 0x0f)
 411                   | ((c >> 14) & 0x30)), 2, s_skb)
 412             | BOX ((((c >> 20) & 0x01)
 413                   | ((c >> 21) & 0x06)
 414                   | ((c >> 22) & 0x38)), 3, s_skb);
 415
 416     u32x t = BOX ((( d >>  0) & 0x3f),  4, s_skb)
 417             | BOX ((((d >>  7) & 0x03)
 418                   | ((d >>  8) & 0x3c)), 5, s_skb)
 419             | BOX ((((d >> 15) & 0x3f)), 6, s_skb)
 420             | BOX ((((d >> 21) & 0x0f)
 421                   | ((d >> 22) & 0x30)), 7, s_skb);
 422
 423     #if __CUDA_ARCH__ >= 200
 424     Kc[i] = __byte_perm (s, t, 0x5410);
 425     Kd[i] = __byte_perm (s, t, 0x7632);
 426     #else
 427     Kc[i] = ((t << 16) | (s & 0x0000ffff));
 428     Kd[i] = ((s >> 16) | (t & 0xffff0000));
 429     #endif
 430
 431     Kc[i] = rotl32 (Kc[i], 2u);
 432     Kd[i] = rotl32 (Kd[i], 2u);
 433   }
 434 }
 435
 436 __device__ static void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
 437 {
 438   u32x t[8];
 439
 440   t[0] = (w0 >>  0) & 0xff;
 441   t[1] = (w0 >>  8) & 0xff;
 442   t[2] = (w0 >> 16) & 0xff;
 443   t[3] = (w0 >> 24) & 0xff;
 444   t[4] = (w1 >>  0) & 0xff;
 445   t[5] = (w1 >>  8) & 0xff;
 446   t[6] = (w1 >> 16) & 0xff;
 447   t[7] = (w1 >> 24) & 0xff;
 448
 449   u32x k[8];
 450
 451   k[0] =               (t[0] >> 0);
 452   k[1] = (t[0] << 7) | (t[1] >> 1);
 453   k[2] = (t[1] << 6) | (t[2] >> 2);
 454   k[3] = (t[2] << 5) | (t[3] >> 3);
 455   k[4] = (t[3] << 4) | (t[4] >> 4);
 456   k[5] = (t[4] << 3) | (t[5] >> 5);
 457   k[6] = (t[5] << 2) | (t[6] >> 6);
 458   k[7] = (t[6] << 1);
 459
 460   out[0] = ((k[0] & 0xff) <<  0)
 461          | ((k[1] & 0xff) <<  8)
 462          | ((k[2] & 0xff) << 16)
 463          | ((k[3] & 0xff) << 24);
 464
 465   out[1] = ((k[4] & 0xff) <<  0)
 466          | ((k[5] & 0xff) <<  8)
 467          | ((k[6] & 0xff) << 16)
 468          | ((k[7] & 0xff) << 24);
 469 }
 470
 471 __device__ __constant__ comb_t c_combs[1024];
 472
 473 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 474 {
 475   /**
 476    * modifier
 477    */
 478
 479   const u32 lid = threadIdx.x;
 480
 481   /**
 482    * base
 483    */
 484
 485   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 486
 487   u32x wordl0[4];
 488
 489   wordl0[0] = pws[gid].i[ 0];
 490   wordl0[1] = pws[gid].i[ 1];
 491   wordl0[2] = pws[gid].i[ 2];
 492   wordl0[3] = pws[gid].i[ 3];
 493
 494   u32x wordl1[4];
 495
 496   wordl1[0] = pws[gid].i[ 4];
 497   wordl1[1] = pws[gid].i[ 5];
 498   wordl1[2] = pws[gid].i[ 6];
 499   wordl1[3] = pws[gid].i[ 7];
 500
 501   u32x wordl2[4];
 502
 503   wordl2[0] = 0;
 504   wordl2[1] = 0;
 505   wordl2[2] = 0;
 506   wordl2[3] = 0;
 507
 508   u32x wordl3[4];
 509
 510   wordl3[0] = 0;
 511   wordl3[1] = 0;
 512   wordl3[2] = 0;
 513   wordl3[3] = 0;
 514
 515   const u32 pw_l_len = pws[gid].pw_len;
 516
 517   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 518   {
 519     append_0x80_2 (wordl0, wordl1, pw_l_len);
 520
 521     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 522   }
 523
 524   /**
 525    * sbox, kbox
 526    */
 527
 528   __shared__ u32 s_SPtrans[8][64];
 529   __shared__ u32 s_skb[8][64];
 530
 531   if (lid < 64)
 532   {
 533     s_SPtrans[0][lid] = c_SPtrans[0][lid];
 534     s_SPtrans[1][lid] = c_SPtrans[1][lid];
 535     s_SPtrans[2][lid] = c_SPtrans[2][lid];
 536     s_SPtrans[3][lid] = c_SPtrans[3][lid];
 537     s_SPtrans[4][lid] = c_SPtrans[4][lid];
 538     s_SPtrans[5][lid] = c_SPtrans[5][lid];
 539     s_SPtrans[6][lid] = c_SPtrans[6][lid];
 540     s_SPtrans[7][lid] = c_SPtrans[7][lid];
 541
 542     s_skb[0][lid] = c_skb[0][lid];
 543     s_skb[1][lid] = c_skb[1][lid];
 544     s_skb[2][lid] = c_skb[2][lid];
 545     s_skb[3][lid] = c_skb[3][lid];
 546     s_skb[4][lid] = c_skb[4][lid];
 547     s_skb[5][lid] = c_skb[5][lid];
 548     s_skb[6][lid] = c_skb[6][lid];
 549     s_skb[7][lid] = c_skb[7][lid];
 550   }
 551
 552   __syncthreads ();
 553
 554   if (gid >= gid_max) return;
 555
 556   /**
 557    * salt
 558    */
 559
 560   const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
 561   const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
 562   const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
 563
 564   u32x data[2];
 565
 566   data[0] = s0;
 567   data[1] = s1;
 568
 569   /**
 570    * loop
 571    */
 572
 573   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 574   {
 575     const u32 pw_r_len = c_combs[il_pos].pw_len;
 576
 577     const u32 pw_len = pw_l_len + pw_r_len;
 578
 579     u32 wordr0[4];
 580
 581     wordr0[0] = c_combs[il_pos].i[0];
 582     wordr0[1] = c_combs[il_pos].i[1];
 583     wordr0[2] = c_combs[il_pos].i[2];
 584     wordr0[3] = c_combs[il_pos].i[3];
 585
 586     u32 wordr1[4];
 587
 588     wordr1[0] = c_combs[il_pos].i[4];
 589     wordr1[1] = c_combs[il_pos].i[5];
 590     wordr1[2] = c_combs[il_pos].i[6];
 591     wordr1[3] = c_combs[il_pos].i[7];
 592
 593     u32 wordr2[4];
 594
 595     wordr2[0] = 0;
 596     wordr2[1] = 0;
 597     wordr2[2] = 0;
 598     wordr2[3] = 0;
 599
 600     u32 wordr3[4];
 601
 602     wordr3[0] = 0;
 603     wordr3[1] = 0;
 604     wordr3[2] = 0;
 605     wordr3[3] = 0;
 606
 607     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 608     {
 609       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 610     }
 611
 612     u32x w0[4];
 613
 614     w0[0] = wordl0[0] | wordr0[0];
 615     w0[1] = wordl0[1] | wordr0[1];
 616     w0[2] = wordl0[2] | wordr0[2];
 617     w0[3] = wordl0[3] | wordr0[3];
 618
 619     u32x w1[4];
 620
 621     w1[0] = wordl1[0] | wordr1[0];
 622     w1[1] = wordl1[1] | wordr1[1];
 623     w1[2] = wordl1[2] | wordr1[2];
 624     w1[3] = wordl1[3] | wordr1[3];
 625
 626     u32x w2[4];
 627
 628     w2[0] = 0;
 629     w2[1] = 0;
 630     w2[2] = 0;
 631     w2[3] = 0;
 632
 633     u32x w3[4];
 634
 635     w3[0] = 0;
 636     w3[1] = 0;
 637     w3[2] = 0;
 638     w3[3] = 0;
 639
 640     u32x w0_t[4];
 641     u32x w1_t[4];
 642     u32x w2_t[4];
 643     u32x w3_t[4];
 644
 645     make_unicode (w0, w0_t, w1_t);
 646     make_unicode (w1, w2_t, w3_t);
 647
 648     w3_t[2] = pw_len * 8 * 2;
 649
 650     u32x a = MD4M_A;
 651     u32x b = MD4M_B;
 652     u32x c = MD4M_C;
 653     u32x d = MD4M_D;
 654
 655     MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
 656     MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
 657     MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
 658     MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
 659     MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
 660     MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
 661     MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
 662     MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
 663     MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
 664     MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
 665     MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
 666     MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
 667     MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
 668     MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
 669     MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
 670     MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
 671
 672     MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
 673     MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
 674     MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
 675     MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
 676     MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
 677     MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
 678     MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
 679     MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
 680     MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
 681     MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
 682     MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
 683     MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
 684     MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
 685     MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
 686     MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
 687     MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
 688
 689     MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
 690     MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
 691     MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
 692     MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
 693     MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
 694     MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
 695     MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
 696     MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
 697     MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
 698     MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
 699     MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
 700     MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
 701     MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
 702     MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
 703
 704     if (s2 != ((d + MD4M_D) >> 16)) continue;
 705
 706     MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
 707     MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
 708
 709     a += MD4M_A;
 710     b += MD4M_B;
 711     c += MD4M_C;
 712     d += MD4M_D;
 713
 714     /**
 715      * DES1
 716      */
 717
 718     u32x key[2];
 719
 720     transform_netntlmv1_key (a, b, key);
 721
 722     u32x Kc[16];
 723     u32x Kd[16];
 724
 725     _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
 726
 727     u32x iv1[2];
 728
 729     _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
 730
 731     /**
 732      * DES2
 733      */
 734
 735     transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
 736
 737     _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
 738
 739     u32x iv2[2];
 740
 741     _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
 742
 743     /**
 744      * compare
 745      */
 746
 747     const u32x r0 = iv1[0];
 748     const u32x r1 = iv1[1];
 749     const u32x r2 = iv2[0];
 750     const u32x r3 = iv2[1];
 751
 752     #include VECT_COMPARE_M
 753   }
 754 }
 755
 756 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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)
 757 {
 758 }
 759
 760 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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)
 761 {
 762 }
 763
 764 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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 combs_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 765 {
 766   /**
 767    * modifier
 768    */
 769
 770   const u32 lid = threadIdx.x;
 771
 772   /**
 773    * base
 774    */
 775
 776   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 777
 778   u32x wordl0[4];
 779
 780   wordl0[0] = pws[gid].i[ 0];
 781   wordl0[1] = pws[gid].i[ 1];
 782   wordl0[2] = pws[gid].i[ 2];
 783   wordl0[3] = pws[gid].i[ 3];
 784
 785   u32x wordl1[4];
 786
 787   wordl1[0] = pws[gid].i[ 4];
 788   wordl1[1] = pws[gid].i[ 5];
 789   wordl1[2] = pws[gid].i[ 6];
 790   wordl1[3] = pws[gid].i[ 7];
 791
 792   u32x wordl2[4];
 793
 794   wordl2[0] = 0;
 795   wordl2[1] = 0;
 796   wordl2[2] = 0;
 797   wordl2[3] = 0;
 798
 799   u32x wordl3[4];
 800
 801   wordl3[0] = 0;
 802   wordl3[1] = 0;
 803   wordl3[2] = 0;
 804   wordl3[3] = 0;
 805
 806   const u32 pw_l_len = pws[gid].pw_len;
 807
 808   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
 809   {
 810     append_0x80_2 (wordl0, wordl1, pw_l_len);
 811
 812     switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, c_combs[0].pw_len);
 813   }
 814
 815   /**
 816    * sbox, kbox
 817    */
 818
 819   __shared__ u32 s_SPtrans[8][64];
 820   __shared__ u32 s_skb[8][64];
 821
 822   if (lid < 64)
 823   {
 824     s_SPtrans[0][lid] = c_SPtrans[0][lid];
 825     s_SPtrans[1][lid] = c_SPtrans[1][lid];
 826     s_SPtrans[2][lid] = c_SPtrans[2][lid];
 827     s_SPtrans[3][lid] = c_SPtrans[3][lid];
 828     s_SPtrans[4][lid] = c_SPtrans[4][lid];
 829     s_SPtrans[5][lid] = c_SPtrans[5][lid];
 830     s_SPtrans[6][lid] = c_SPtrans[6][lid];
 831     s_SPtrans[7][lid] = c_SPtrans[7][lid];
 832
 833     s_skb[0][lid] = c_skb[0][lid];
 834     s_skb[1][lid] = c_skb[1][lid];
 835     s_skb[2][lid] = c_skb[2][lid];
 836     s_skb[3][lid] = c_skb[3][lid];
 837     s_skb[4][lid] = c_skb[4][lid];
 838     s_skb[5][lid] = c_skb[5][lid];
 839     s_skb[6][lid] = c_skb[6][lid];
 840     s_skb[7][lid] = c_skb[7][lid];
 841   }
 842
 843   __syncthreads ();
 844
 845   if (gid >= gid_max) return;
 846
 847   /**
 848    * salt
 849    */
 850
 851   const u32 s0 = salt_bufs[salt_pos].salt_buf[0];
 852   const u32 s1 = salt_bufs[salt_pos].salt_buf[1];
 853   const u32 s2 = salt_bufs[salt_pos].salt_buf[2];
 854
 855   u32x data[2];
 856
 857   data[0] = s0;
 858   data[1] = s1;
 859
 860   /**
 861    * digest
 862    */
 863
 864   const u32 search[4] =
 865   {
 866     digests_buf[digests_offset].digest_buf[DGST_R0],
 867     digests_buf[digests_offset].digest_buf[DGST_R1],
 868     digests_buf[digests_offset].digest_buf[DGST_R2],
 869     digests_buf[digests_offset].digest_buf[DGST_R3]
 870   };
 871
 872   /**
 873    * loop
 874    */
 875
 876   for (u32 il_pos = 0; il_pos < combs_cnt; il_pos++)
 877   {
 878     const u32 pw_r_len = c_combs[il_pos].pw_len;
 879
 880     const u32 pw_len = pw_l_len + pw_r_len;
 881
 882     u32 wordr0[4];
 883
 884     wordr0[0] = c_combs[il_pos].i[0];
 885     wordr0[1] = c_combs[il_pos].i[1];
 886     wordr0[2] = c_combs[il_pos].i[2];
 887     wordr0[3] = c_combs[il_pos].i[3];
 888
 889     u32 wordr1[4];
 890
 891     wordr1[0] = c_combs[il_pos].i[4];
 892     wordr1[1] = c_combs[il_pos].i[5];
 893     wordr1[2] = c_combs[il_pos].i[6];
 894     wordr1[3] = c_combs[il_pos].i[7];
 895
 896     u32 wordr2[4];
 897
 898     wordr2[0] = 0;
 899     wordr2[1] = 0;
 900     wordr2[2] = 0;
 901     wordr2[3] = 0;
 902
 903     u32 wordr3[4];
 904
 905     wordr3[0] = 0;
 906     wordr3[1] = 0;
 907     wordr3[2] = 0;
 908     wordr3[3] = 0;
 909
 910     if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
 911     {
 912       switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
 913     }
 914
 915     u32x w0[4];
 916
 917     w0[0] = wordl0[0] | wordr0[0];
 918     w0[1] = wordl0[1] | wordr0[1];
 919     w0[2] = wordl0[2] | wordr0[2];
 920     w0[3] = wordl0[3] | wordr0[3];
 921
 922     u32x w1[4];
 923
 924     w1[0] = wordl1[0] | wordr1[0];
 925     w1[1] = wordl1[1] | wordr1[1];
 926     w1[2] = wordl1[2] | wordr1[2];
 927     w1[3] = wordl1[3] | wordr1[3];
 928
 929     u32x w2[4];
 930
 931     w2[0] = 0;
 932     w2[1] = 0;
 933     w2[2] = 0;
 934     w2[3] = 0;
 935
 936     u32x w3[4];
 937
 938     w3[0] = 0;
 939     w3[1] = 0;
 940     w3[2] = 0;
 941     w3[3] = 0;
 942
 943     u32x w0_t[4];
 944     u32x w1_t[4];
 945     u32x w2_t[4];
 946     u32x w3_t[4];
 947
 948     make_unicode (w0, w0_t, w1_t);
 949     make_unicode (w1, w2_t, w3_t);
 950
 951     w3_t[2] = pw_len * 8 * 2;
 952
 953     u32x a = MD4M_A;
 954     u32x b = MD4M_B;
 955     u32x c = MD4M_C;
 956     u32x d = MD4M_D;
 957
 958     MD4_STEP (MD4_Fo, a, b, c, d, w0_t[0], MD4C00, MD4S00);
 959     MD4_STEP (MD4_Fo, d, a, b, c, w0_t[1], MD4C00, MD4S01);
 960     MD4_STEP (MD4_Fo, c, d, a, b, w0_t[2], MD4C00, MD4S02);
 961     MD4_STEP (MD4_Fo, b, c, d, a, w0_t[3], MD4C00, MD4S03);
 962     MD4_STEP (MD4_Fo, a, b, c, d, w1_t[0], MD4C00, MD4S00);
 963     MD4_STEP (MD4_Fo, d, a, b, c, w1_t[1], MD4C00, MD4S01);
 964     MD4_STEP (MD4_Fo, c, d, a, b, w1_t[2], MD4C00, MD4S02);
 965     MD4_STEP (MD4_Fo, b, c, d, a, w1_t[3], MD4C00, MD4S03);
 966     MD4_STEP (MD4_Fo, a, b, c, d, w2_t[0], MD4C00, MD4S00);
 967     MD4_STEP (MD4_Fo, d, a, b, c, w2_t[1], MD4C00, MD4S01);
 968     MD4_STEP (MD4_Fo, c, d, a, b, w2_t[2], MD4C00, MD4S02);
 969     MD4_STEP (MD4_Fo, b, c, d, a, w2_t[3], MD4C00, MD4S03);
 970     MD4_STEP (MD4_Fo, a, b, c, d, w3_t[0], MD4C00, MD4S00);
 971     MD4_STEP (MD4_Fo, d, a, b, c, w3_t[1], MD4C00, MD4S01);
 972     MD4_STEP (MD4_Fo, c, d, a, b, w3_t[2], MD4C00, MD4S02);
 973     MD4_STEP (MD4_Fo, b, c, d, a, w3_t[3], MD4C00, MD4S03);
 974
 975     MD4_STEP (MD4_Go, a, b, c, d, w0_t[0], MD4C01, MD4S10);
 976     MD4_STEP (MD4_Go, d, a, b, c, w1_t[0], MD4C01, MD4S11);
 977     MD4_STEP (MD4_Go, c, d, a, b, w2_t[0], MD4C01, MD4S12);
 978     MD4_STEP (MD4_Go, b, c, d, a, w3_t[0], MD4C01, MD4S13);
 979     MD4_STEP (MD4_Go, a, b, c, d, w0_t[1], MD4C01, MD4S10);
 980     MD4_STEP (MD4_Go, d, a, b, c, w1_t[1], MD4C01, MD4S11);
 981     MD4_STEP (MD4_Go, c, d, a, b, w2_t[1], MD4C01, MD4S12);
 982     MD4_STEP (MD4_Go, b, c, d, a, w3_t[1], MD4C01, MD4S13);
 983     MD4_STEP (MD4_Go, a, b, c, d, w0_t[2], MD4C01, MD4S10);
 984     MD4_STEP (MD4_Go, d, a, b, c, w1_t[2], MD4C01, MD4S11);
 985     MD4_STEP (MD4_Go, c, d, a, b, w2_t[2], MD4C01, MD4S12);
 986     MD4_STEP (MD4_Go, b, c, d, a, w3_t[2], MD4C01, MD4S13);
 987     MD4_STEP (MD4_Go, a, b, c, d, w0_t[3], MD4C01, MD4S10);
 988     MD4_STEP (MD4_Go, d, a, b, c, w1_t[3], MD4C01, MD4S11);
 989     MD4_STEP (MD4_Go, c, d, a, b, w2_t[3], MD4C01, MD4S12);
 990     MD4_STEP (MD4_Go, b, c, d, a, w3_t[3], MD4C01, MD4S13);
 991
 992     MD4_STEP (MD4_H , a, b, c, d, w0_t[0], MD4C02, MD4S20);
 993     MD4_STEP (MD4_H , d, a, b, c, w2_t[0], MD4C02, MD4S21);
 994     MD4_STEP (MD4_H , c, d, a, b, w1_t[0], MD4C02, MD4S22);
 995     MD4_STEP (MD4_H , b, c, d, a, w3_t[0], MD4C02, MD4S23);
 996     MD4_STEP (MD4_H , a, b, c, d, w0_t[2], MD4C02, MD4S20);
 997     MD4_STEP (MD4_H , d, a, b, c, w2_t[2], MD4C02, MD4S21);
 998     MD4_STEP (MD4_H , c, d, a, b, w1_t[2], MD4C02, MD4S22);
 999     MD4_STEP (MD4_H , b, c, d, a, w3_t[2], MD4C02, MD4S23);
1000     MD4_STEP (MD4_H , a, b, c, d, w0_t[1], MD4C02, MD4S20);
1001     MD4_STEP (MD4_H , d, a, b, c, w2_t[1], MD4C02, MD4S21);
1002     MD4_STEP (MD4_H , c, d, a, b, w1_t[1], MD4C02, MD4S22);
1003     MD4_STEP (MD4_H , b, c, d, a, w3_t[1], MD4C02, MD4S23);
1004     MD4_STEP (MD4_H , a, b, c, d, w0_t[3], MD4C02, MD4S20);
1005     MD4_STEP (MD4_H , d, a, b, c, w2_t[3], MD4C02, MD4S21);
1006
1007     if (s2 != ((d + MD4M_D) >> 16)) continue;
1008
1009     MD4_STEP (MD4_H , c, d, a, b, w1_t[3], MD4C02, MD4S22);
1010     MD4_STEP (MD4_H , b, c, d, a, w3_t[3], MD4C02, MD4S23);
1011
1012     a += MD4M_A;
1013     b += MD4M_B;
1014     c += MD4M_C;
1015     d += MD4M_D;
1016
1017     /**
1018      * DES1
1019      */
1020
1021     u32x key[2];
1022
1023     transform_netntlmv1_key (a, b, key);
1024
1025     u32x Kc[16];
1026     u32x Kd[16];
1027
1028     _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1029
1030     u32x iv1[2];
1031
1032     _des_crypt_encrypt (iv1, data, Kc, Kd, s_SPtrans);
1033
1034     /**
1035      * DES2
1036      */
1037
1038     transform_netntlmv1_key (((b >> 24) | (c << 8)), ((c >> 24) | (d << 8)), key);
1039
1040     _des_crypt_keysetup (key[0], key[1], Kc, Kd, s_skb);
1041
1042     u32x iv2[2];
1043
1044     _des_crypt_encrypt (iv2, data, Kc, Kd, s_SPtrans);
1045
1046     /**
1047      * compare
1048      */
1049
1050     const u32x r0 = iv1[0];
1051     const u32x r1 = iv1[1];
1052     const u32x r2 = iv2[0];
1053     const u32x r3 = iv2[1];
1054
1055     #include VECT_COMPARE_S
1056   }
1057 }
1058
1059 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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)
1060 {
1061 }
1062
1063 extern "C" __global__ void __launch_bounds__ (256, 1) m05500_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 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)
1064 {
1065 }