nv/m01500_a0.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _DES_
   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 #include "include/rp_gpu.h"
  28 #include "rp_nv.c"
  29
  30 #ifdef  VECT_SIZE1
  31 #define VECT_COMPARE_S "check_single_vect1_comp4.c"
  32 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  33 #endif
  34
  35 #ifdef  VECT_SIZE2
  36 #define VECT_COMPARE_S "check_single_vect2_comp4.c"
  37 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  38 #endif
  39
  40 #ifdef  VECT_SIZE4
  41 #define VECT_COMPARE_S "check_single_vect4_comp4.c"
  42 #define VECT_COMPARE_M "check_multi_vect4_comp4.c"
  43 #endif
  44
  45 __device__ __constant__ gpu_rule_t c_rules[1024];
  46
  47 #define PERM_OP(a,b,tt,n,m) \
  48 {                           \
  49   tt = a >> n;              \
  50   tt = tt ^ b;              \
  51   tt = tt & m;              \
  52   b = b ^ tt;               \
  53   tt = tt << n;             \
  54   a = a ^ tt;               \
  55 }
  56
  57 #define HPERM_OP(a,tt,n,m)  \
  58 {                           \
  59   tt = a << (16 + n);       \
  60   tt = tt ^ a;              \
  61   tt = tt & m;              \
  62   a  = a ^ tt;              \
  63   tt = tt >> (16 + n);      \
  64   a  = a ^ tt;              \
  65 }
  66
  67 __device__ __constant__ u32 c_SPtrans[8][64] =
  68 {
  69   /* nibble 0 */
  70   0x00820200, 0x00020000, 0x80800000, 0x80820200,
  71   0x00800000, 0x80020200, 0x80020000, 0x80800000,
  72   0x80020200, 0x00820200, 0x00820000, 0x80000200,
  73   0x80800200, 0x00800000, 0x00000000, 0x80020000,
  74   0x00020000, 0x80000000, 0x00800200, 0x00020200,
  75   0x80820200, 0x00820000, 0x80000200, 0x00800200,
  76   0x80000000, 0x00000200, 0x00020200, 0x80820000,
  77   0x00000200, 0x80800200, 0x80820000, 0x00000000,
  78   0x00000000, 0x80820200, 0x00800200, 0x80020000,
  79   0x00820200, 0x00020000, 0x80000200, 0x00800200,
  80   0x80820000, 0x00000200, 0x00020200, 0x80800000,
  81   0x80020200, 0x80000000, 0x80800000, 0x00820000,
  82   0x80820200, 0x00020200, 0x00820000, 0x80800200,
  83   0x00800000, 0x80000200, 0x80020000, 0x00000000,
  84   0x00020000, 0x00800000, 0x80800200, 0x00820200,
  85   0x80000000, 0x80820000, 0x00000200, 0x80020200,
  86   /* nibble 1 */
  87   0x10042004, 0x00000000, 0x00042000, 0x10040000,
  88   0x10000004, 0x00002004, 0x10002000, 0x00042000,
  89   0x00002000, 0x10040004, 0x00000004, 0x10002000,
  90   0x00040004, 0x10042000, 0x10040000, 0x00000004,
  91   0x00040000, 0x10002004, 0x10040004, 0x00002000,
  92   0x00042004, 0x10000000, 0x00000000, 0x00040004,
  93   0x10002004, 0x00042004, 0x10042000, 0x10000004,
  94   0x10000000, 0x00040000, 0x00002004, 0x10042004,
  95   0x00040004, 0x10042000, 0x10002000, 0x00042004,
  96   0x10042004, 0x00040004, 0x10000004, 0x00000000,
  97   0x10000000, 0x00002004, 0x00040000, 0x10040004,
  98   0x00002000, 0x10000000, 0x00042004, 0x10002004,
  99   0x10042000, 0x00002000, 0x00000000, 0x10000004,
 100   0x00000004, 0x10042004, 0x00042000, 0x10040000,
 101   0x10040004, 0x00040000, 0x00002004, 0x10002000,
 102   0x10002004, 0x00000004, 0x10040000, 0x00042000,
 103   /* nibble 2 */
 104   0x41000000, 0x01010040, 0x00000040, 0x41000040,
 105   0x40010000, 0x01000000, 0x41000040, 0x00010040,
 106   0x01000040, 0x00010000, 0x01010000, 0x40000000,
 107   0x41010040, 0x40000040, 0x40000000, 0x41010000,
 108   0x00000000, 0x40010000, 0x01010040, 0x00000040,
 109   0x40000040, 0x41010040, 0x00010000, 0x41000000,
 110   0x41010000, 0x01000040, 0x40010040, 0x01010000,
 111   0x00010040, 0x00000000, 0x01000000, 0x40010040,
 112   0x01010040, 0x00000040, 0x40000000, 0x00010000,
 113   0x40000040, 0x40010000, 0x01010000, 0x41000040,
 114   0x00000000, 0x01010040, 0x00010040, 0x41010000,
 115   0x40010000, 0x01000000, 0x41010040, 0x40000000,
 116   0x40010040, 0x41000000, 0x01000000, 0x41010040,
 117   0x00010000, 0x01000040, 0x41000040, 0x00010040,
 118   0x01000040, 0x00000000, 0x41010000, 0x40000040,
 119   0x41000000, 0x40010040, 0x00000040, 0x01010000,
 120   /* nibble 3 */
 121   0x00100402, 0x04000400, 0x00000002, 0x04100402,
 122   0x00000000, 0x04100000, 0x04000402, 0x00100002,
 123   0x04100400, 0x04000002, 0x04000000, 0x00000402,
 124   0x04000002, 0x00100402, 0x00100000, 0x04000000,
 125   0x04100002, 0x00100400, 0x00000400, 0x00000002,
 126   0x00100400, 0x04000402, 0x04100000, 0x00000400,
 127   0x00000402, 0x00000000, 0x00100002, 0x04100400,
 128   0x04000400, 0x04100002, 0x04100402, 0x00100000,
 129   0x04100002, 0x00000402, 0x00100000, 0x04000002,
 130   0x00100400, 0x04000400, 0x00000002, 0x04100000,
 131   0x04000402, 0x00000000, 0x00000400, 0x00100002,
 132   0x00000000, 0x04100002, 0x04100400, 0x00000400,
 133   0x04000000, 0x04100402, 0x00100402, 0x00100000,
 134   0x04100402, 0x00000002, 0x04000400, 0x00100402,
 135   0x00100002, 0x00100400, 0x04100000, 0x04000402,
 136   0x00000402, 0x04000000, 0x04000002, 0x04100400,
 137   /* nibble 4 */
 138   0x02000000, 0x00004000, 0x00000100, 0x02004108,
 139   0x02004008, 0x02000100, 0x00004108, 0x02004000,
 140   0x00004000, 0x00000008, 0x02000008, 0x00004100,
 141   0x02000108, 0x02004008, 0x02004100, 0x00000000,
 142   0x00004100, 0x02000000, 0x00004008, 0x00000108,
 143   0x02000100, 0x00004108, 0x00000000, 0x02000008,
 144   0x00000008, 0x02000108, 0x02004108, 0x00004008,
 145   0x02004000, 0x00000100, 0x00000108, 0x02004100,
 146   0x02004100, 0x02000108, 0x00004008, 0x02004000,
 147   0x00004000, 0x00000008, 0x02000008, 0x02000100,
 148   0x02000000, 0x00004100, 0x02004108, 0x00000000,
 149   0x00004108, 0x02000000, 0x00000100, 0x00004008,
 150   0x02000108, 0x00000100, 0x00000000, 0x02004108,
 151   0x02004008, 0x02004100, 0x00000108, 0x00004000,
 152   0x00004100, 0x02004008, 0x02000100, 0x00000108,
 153   0x00000008, 0x00004108, 0x02004000, 0x02000008,
 154   /* nibble 5 */
 155   0x20000010, 0x00080010, 0x00000000, 0x20080800,
 156   0x00080010, 0x00000800, 0x20000810, 0x00080000,
 157   0x00000810, 0x20080810, 0x00080800, 0x20000000,
 158   0x20000800, 0x20000010, 0x20080000, 0x00080810,
 159   0x00080000, 0x20000810, 0x20080010, 0x00000000,
 160   0x00000800, 0x00000010, 0x20080800, 0x20080010,
 161   0x20080810, 0x20080000, 0x20000000, 0x00000810,
 162   0x00000010, 0x00080800, 0x00080810, 0x20000800,
 163   0x00000810, 0x20000000, 0x20000800, 0x00080810,
 164   0x20080800, 0x00080010, 0x00000000, 0x20000800,
 165   0x20000000, 0x00000800, 0x20080010, 0x00080000,
 166   0x00080010, 0x20080810, 0x00080800, 0x00000010,
 167   0x20080810, 0x00080800, 0x00080000, 0x20000810,
 168   0x20000010, 0x20080000, 0x00080810, 0x00000000,
 169   0x00000800, 0x20000010, 0x20000810, 0x20080800,
 170   0x20080000, 0x00000810, 0x00000010, 0x20080010,
 171   /* nibble 6 */
 172   0x00001000, 0x00000080, 0x00400080, 0x00400001,
 173   0x00401081, 0x00001001, 0x00001080, 0x00000000,
 174   0x00400000, 0x00400081, 0x00000081, 0x00401000,
 175   0x00000001, 0x00401080, 0x00401000, 0x00000081,
 176   0x00400081, 0x00001000, 0x00001001, 0x00401081,
 177   0x00000000, 0x00400080, 0x00400001, 0x00001080,
 178   0x00401001, 0x00001081, 0x00401080, 0x00000001,
 179   0x00001081, 0x00401001, 0x00000080, 0x00400000,
 180   0x00001081, 0x00401000, 0x00401001, 0x00000081,
 181   0x00001000, 0x00000080, 0x00400000, 0x00401001,
 182   0x00400081, 0x00001081, 0x00001080, 0x00000000,
 183   0x00000080, 0x00400001, 0x00000001, 0x00400080,
 184   0x00000000, 0x00400081, 0x00400080, 0x00001080,
 185   0x00000081, 0x00001000, 0x00401081, 0x00400000,
 186   0x00401080, 0x00000001, 0x00001001, 0x00401081,
 187   0x00400001, 0x00401080, 0x00401000, 0x00001001,
 188   /* nibble 7 */
 189   0x08200020, 0x08208000, 0x00008020, 0x00000000,
 190   0x08008000, 0x00200020, 0x08200000, 0x08208020,
 191   0x00000020, 0x08000000, 0x00208000, 0x00008020,
 192   0x00208020, 0x08008020, 0x08000020, 0x08200000,
 193   0x00008000, 0x00208020, 0x00200020, 0x08008000,
 194   0x08208020, 0x08000020, 0x00000000, 0x00208000,
 195   0x08000000, 0x00200000, 0x08008020, 0x08200020,
 196   0x00200000, 0x00008000, 0x08208000, 0x00000020,
 197   0x00200000, 0x00008000, 0x08000020, 0x08208020,
 198   0x00008020, 0x08000000, 0x00000000, 0x00208000,
 199   0x08200020, 0x08008020, 0x08008000, 0x00200020,
 200   0x08208000, 0x00000020, 0x00200020, 0x08008000,
 201   0x08208020, 0x00200000, 0x08200000, 0x08000020,
 202   0x00208000, 0x00008020, 0x08008020, 0x08200000,
 203   0x00000020, 0x08208000, 0x00208020, 0x00000000,
 204   0x08000000, 0x08200020, 0x00008000, 0x00208020
 205 };
 206
 207 __device__ __constant__ u32 c_skb[8][64] =
 208 {
 209   /* for C bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
 210   0x00000000, 0x00000010, 0x20000000, 0x20000010,
 211   0x00010000, 0x00010010, 0x20010000, 0x20010010,
 212   0x00000800, 0x00000810, 0x20000800, 0x20000810,
 213   0x00010800, 0x00010810, 0x20010800, 0x20010810,
 214   0x00000020, 0x00000030, 0x20000020, 0x20000030,
 215   0x00010020, 0x00010030, 0x20010020, 0x20010030,
 216   0x00000820, 0x00000830, 0x20000820, 0x20000830,
 217   0x00010820, 0x00010830, 0x20010820, 0x20010830,
 218   0x00080000, 0x00080010, 0x20080000, 0x20080010,
 219   0x00090000, 0x00090010, 0x20090000, 0x20090010,
 220   0x00080800, 0x00080810, 0x20080800, 0x20080810,
 221   0x00090800, 0x00090810, 0x20090800, 0x20090810,
 222   0x00080020, 0x00080030, 0x20080020, 0x20080030,
 223   0x00090020, 0x00090030, 0x20090020, 0x20090030,
 224   0x00080820, 0x00080830, 0x20080820, 0x20080830,
 225   0x00090820, 0x00090830, 0x20090820, 0x20090830,
 226   /* for C bits (numbered as per FIPS 46) 7 8 10 11 12 13 */
 227   0x00000000, 0x02000000, 0x00002000, 0x02002000,
 228   0x00200000, 0x02200000, 0x00202000, 0x02202000,
 229   0x00000004, 0x02000004, 0x00002004, 0x02002004,
 230   0x00200004, 0x02200004, 0x00202004, 0x02202004,
 231   0x00000400, 0x02000400, 0x00002400, 0x02002400,
 232   0x00200400, 0x02200400, 0x00202400, 0x02202400,
 233   0x00000404, 0x02000404, 0x00002404, 0x02002404,
 234   0x00200404, 0x02200404, 0x00202404, 0x02202404,
 235   0x10000000, 0x12000000, 0x10002000, 0x12002000,
 236   0x10200000, 0x12200000, 0x10202000, 0x12202000,
 237   0x10000004, 0x12000004, 0x10002004, 0x12002004,
 238   0x10200004, 0x12200004, 0x10202004, 0x12202004,
 239   0x10000400, 0x12000400, 0x10002400, 0x12002400,
 240   0x10200400, 0x12200400, 0x10202400, 0x12202400,
 241   0x10000404, 0x12000404, 0x10002404, 0x12002404,
 242   0x10200404, 0x12200404, 0x10202404, 0x12202404,
 243   /* for C bits (numbered as per FIPS 46) 14 15 16 17 19 20 */
 244   0x00000000, 0x00000001, 0x00040000, 0x00040001,
 245   0x01000000, 0x01000001, 0x01040000, 0x01040001,
 246   0x00000002, 0x00000003, 0x00040002, 0x00040003,
 247   0x01000002, 0x01000003, 0x01040002, 0x01040003,
 248   0x00000200, 0x00000201, 0x00040200, 0x00040201,
 249   0x01000200, 0x01000201, 0x01040200, 0x01040201,
 250   0x00000202, 0x00000203, 0x00040202, 0x00040203,
 251   0x01000202, 0x01000203, 0x01040202, 0x01040203,
 252   0x08000000, 0x08000001, 0x08040000, 0x08040001,
 253   0x09000000, 0x09000001, 0x09040000, 0x09040001,
 254   0x08000002, 0x08000003, 0x08040002, 0x08040003,
 255   0x09000002, 0x09000003, 0x09040002, 0x09040003,
 256   0x08000200, 0x08000201, 0x08040200, 0x08040201,
 257   0x09000200, 0x09000201, 0x09040200, 0x09040201,
 258   0x08000202, 0x08000203, 0x08040202, 0x08040203,
 259   0x09000202, 0x09000203, 0x09040202, 0x09040203,
 260   /* for C bits (numbered as per FIPS 46) 21 23 24 26 27 28 */
 261   0x00000000, 0x00100000, 0x00000100, 0x00100100,
 262   0x00000008, 0x00100008, 0x00000108, 0x00100108,
 263   0x00001000, 0x00101000, 0x00001100, 0x00101100,
 264   0x00001008, 0x00101008, 0x00001108, 0x00101108,
 265   0x04000000, 0x04100000, 0x04000100, 0x04100100,
 266   0x04000008, 0x04100008, 0x04000108, 0x04100108,
 267   0x04001000, 0x04101000, 0x04001100, 0x04101100,
 268   0x04001008, 0x04101008, 0x04001108, 0x04101108,
 269   0x00020000, 0x00120000, 0x00020100, 0x00120100,
 270   0x00020008, 0x00120008, 0x00020108, 0x00120108,
 271   0x00021000, 0x00121000, 0x00021100, 0x00121100,
 272   0x00021008, 0x00121008, 0x00021108, 0x00121108,
 273   0x04020000, 0x04120000, 0x04020100, 0x04120100,
 274   0x04020008, 0x04120008, 0x04020108, 0x04120108,
 275   0x04021000, 0x04121000, 0x04021100, 0x04121100,
 276   0x04021008, 0x04121008, 0x04021108, 0x04121108,
 277   /* for D bits (numbered as per FIPS 46) 1 2 3 4 5 6 */
 278   0x00000000, 0x10000000, 0x00010000, 0x10010000,
 279   0x00000004, 0x10000004, 0x00010004, 0x10010004,
 280   0x20000000, 0x30000000, 0x20010000, 0x30010000,
 281   0x20000004, 0x30000004, 0x20010004, 0x30010004,
 282   0x00100000, 0x10100000, 0x00110000, 0x10110000,
 283   0x00100004, 0x10100004, 0x00110004, 0x10110004,
 284   0x20100000, 0x30100000, 0x20110000, 0x30110000,
 285   0x20100004, 0x30100004, 0x20110004, 0x30110004,
 286   0x00001000, 0x10001000, 0x00011000, 0x10011000,
 287   0x00001004, 0x10001004, 0x00011004, 0x10011004,
 288   0x20001000, 0x30001000, 0x20011000, 0x30011000,
 289   0x20001004, 0x30001004, 0x20011004, 0x30011004,
 290   0x00101000, 0x10101000, 0x00111000, 0x10111000,
 291   0x00101004, 0x10101004, 0x00111004, 0x10111004,
 292   0x20101000, 0x30101000, 0x20111000, 0x30111000,
 293   0x20101004, 0x30101004, 0x20111004, 0x30111004,
 294   /* for D bits (numbered as per FIPS 46) 8 9 11 12 13 14 */
 295   0x00000000, 0x08000000, 0x00000008, 0x08000008,
 296   0x00000400, 0x08000400, 0x00000408, 0x08000408,
 297   0x00020000, 0x08020000, 0x00020008, 0x08020008,
 298   0x00020400, 0x08020400, 0x00020408, 0x08020408,
 299   0x00000001, 0x08000001, 0x00000009, 0x08000009,
 300   0x00000401, 0x08000401, 0x00000409, 0x08000409,
 301   0x00020001, 0x08020001, 0x00020009, 0x08020009,
 302   0x00020401, 0x08020401, 0x00020409, 0x08020409,
 303   0x02000000, 0x0A000000, 0x02000008, 0x0A000008,
 304   0x02000400, 0x0A000400, 0x02000408, 0x0A000408,
 305   0x02020000, 0x0A020000, 0x02020008, 0x0A020008,
 306   0x02020400, 0x0A020400, 0x02020408, 0x0A020408,
 307   0x02000001, 0x0A000001, 0x02000009, 0x0A000009,
 308   0x02000401, 0x0A000401, 0x02000409, 0x0A000409,
 309   0x02020001, 0x0A020001, 0x02020009, 0x0A020009,
 310   0x02020401, 0x0A020401, 0x02020409, 0x0A020409,
 311   /* for D bits (numbered as per FIPS 46) 16 17 18 19 20 21 */
 312   0x00000000, 0x00000100, 0x00080000, 0x00080100,
 313   0x01000000, 0x01000100, 0x01080000, 0x01080100,
 314   0x00000010, 0x00000110, 0x00080010, 0x00080110,
 315   0x01000010, 0x01000110, 0x01080010, 0x01080110,
 316   0x00200000, 0x00200100, 0x00280000, 0x00280100,
 317   0x01200000, 0x01200100, 0x01280000, 0x01280100,
 318   0x00200010, 0x00200110, 0x00280010, 0x00280110,
 319   0x01200010, 0x01200110, 0x01280010, 0x01280110,
 320   0x00000200, 0x00000300, 0x00080200, 0x00080300,
 321   0x01000200, 0x01000300, 0x01080200, 0x01080300,
 322   0x00000210, 0x00000310, 0x00080210, 0x00080310,
 323   0x01000210, 0x01000310, 0x01080210, 0x01080310,
 324   0x00200200, 0x00200300, 0x00280200, 0x00280300,
 325   0x01200200, 0x01200300, 0x01280200, 0x01280300,
 326   0x00200210, 0x00200310, 0x00280210, 0x00280310,
 327   0x01200210, 0x01200310, 0x01280210, 0x01280310,
 328   /* for D bits (numbered as per FIPS 46) 22 23 24 25 27 28 */
 329   0x00000000, 0x04000000, 0x00040000, 0x04040000,
 330   0x00000002, 0x04000002, 0x00040002, 0x04040002,
 331   0x00002000, 0x04002000, 0x00042000, 0x04042000,
 332   0x00002002, 0x04002002, 0x00042002, 0x04042002,
 333   0x00000020, 0x04000020, 0x00040020, 0x04040020,
 334   0x00000022, 0x04000022, 0x00040022, 0x04040022,
 335   0x00002020, 0x04002020, 0x00042020, 0x04042020,
 336   0x00002022, 0x04002022, 0x00042022, 0x04042022,
 337   0x00000800, 0x04000800, 0x00040800, 0x04040800,
 338   0x00000802, 0x04000802, 0x00040802, 0x04040802,
 339   0x00002800, 0x04002800, 0x00042800, 0x04042800,
 340   0x00002802, 0x04002802, 0x00042802, 0x04042802,
 341   0x00000820, 0x04000820, 0x00040820, 0x04040820,
 342   0x00000822, 0x04000822, 0x00040822, 0x04040822,
 343   0x00002820, 0x04002820, 0x00042820, 0x04042820,
 344   0x00002822, 0x04002822, 0x00042822, 0x04042822
 345 };
 346
 347 #ifdef VECT_SIZE1
 348 #define BOX(i,n,S) u32x ((S)[(n)][(i)])
 349 #endif
 350
 351 #ifdef VECT_SIZE2
 352 #define BOX(i,n,S) u32x ((S)[(n)][(i).x], (S)[(n)][(i).y])
 353 #endif
 354
 355 __device__ static void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], u32 s_skb[8][64])
 356 {
 357   u32x tt;
 358
 359   PERM_OP  (d, c, tt, 4, 0x0f0f0f0f);
 360   HPERM_OP (c,    tt, 2, 0xcccc0000);
 361   HPERM_OP (d,    tt, 2, 0xcccc0000);
 362   PERM_OP  (d, c, tt, 1, 0x55555555);
 363   PERM_OP  (c, d, tt, 8, 0x00ff00ff);
 364   PERM_OP  (d, c, tt, 1, 0x55555555);
 365
 366   d = ((d & 0x000000ff) << 16)
 367     | ((d & 0x0000ff00) <<  0)
 368     | ((d & 0x00ff0000) >> 16)
 369     | ((c & 0xf0000000) >>  4);
 370
 371   c = c & 0x0fffffff;
 372
 373   #pragma unroll
 374   for (u32 i = 0; i < 16; i++)
 375   {
 376     if ((i < 2) || (i == 8) || (i == 15))
 377     {
 378       c = ((c >> 1) | (c << 27));
 379       d = ((d >> 1) | (d << 27));
 380     }
 381     else
 382     {
 383       c = ((c >> 2) | (c << 26));
 384       d = ((d >> 2) | (d << 26));
 385     }
 386
 387     c = c & 0x0fffffff;
 388     d = d & 0x0fffffff;
 389
 390     const u32x c00 = (c >>  0) & 0x0000003f;
 391     const u32x c06 = (c >>  6) & 0x00383003;
 392     const u32x c07 = (c >>  7) & 0x0000003c;
 393     const u32x c13 = (c >> 13) & 0x0000060f;
 394     const u32x c20 = (c >> 20) & 0x00000001;
 395
 396     u32x s = BOX (((c00 >>  0) & 0xff), 0, s_skb)
 397             | BOX (((c06 >>  0) & 0xff)
 398                   |((c07 >>  0) & 0xff), 1, s_skb)
 399             | BOX (((c13 >>  0) & 0xff)
 400                   |((c06 >>  8) & 0xff), 2, s_skb)
 401             | BOX (((c20 >>  0) & 0xff)
 402                   |((c13 >>  8) & 0xff)
 403                   |((c06 >> 16) & 0xff), 3, s_skb);
 404
 405     const u32x d00 = (d >>  0) & 0x00003c3f;
 406     const u32x d07 = (d >>  7) & 0x00003f03;
 407     const u32x d21 = (d >> 21) & 0x0000000f;
 408     const u32x d22 = (d >> 22) & 0x00000030;
 409
 410     u32x t = BOX (((d00 >>  0) & 0xff), 4, s_skb)
 411             | BOX (((d07 >>  0) & 0xff)
 412                   |((d00 >>  8) & 0xff), 5, s_skb)
 413             | BOX (((d07 >>  8) & 0xff), 6, s_skb)
 414             | BOX (((d21 >>  0) & 0xff)
 415                   |((d22 >>  0) & 0xff), 7, s_skb);
 416
 417     #if __CUDA_ARCH__ >= 200
 418     Kc[i] = __byte_perm (s, t, 0x5410);
 419     Kd[i] = __byte_perm (s, t, 0x7632);
 420     #else
 421     Kc[i] = ((t << 16) | (s & 0x0000ffff));
 422     Kd[i] = ((s >> 16) | (t & 0xffff0000));
 423     #endif
 424   }
 425 }
 426
 427 __device__ static void _des_crypt_encrypt (u32x iv[2], u32 mask, u32x Kc[16], u32x Kd[16], u32 s_SPtrans[8][64])
 428 {
 429   const u32 E1 = (mask >> 2) & 0x3f0;
 430
 431   const u32 E0 = mask & 0x3f;
 432
 433   u32x r = 0;
 434   u32x l = 0;
 435
 436   for (u32 i = 0; i < 25; i++)
 437   {
 438     #pragma unroll
 439     for (u32 j = 0; j < 16; j += 2)
 440     {
 441       u32x t;
 442       u32x u;
 443
 444       t = r ^ (r >> 16);
 445       u = t & E0;
 446       t = t & E1;
 447       u = u ^ (u << 16);
 448       u = u ^ r;
 449       u = u ^ Kc[j + 0];
 450       t = t ^ (t << 16);
 451       t = t ^ r;
 452       t = rotl32 (t, 28u);
 453       t = t ^ Kd[j + 0];
 454
 455       l ^= BOX (((u >>  0) & 0x3f), 0, s_SPtrans)
 456          | BOX (((u >>  8) & 0x3f), 2, s_SPtrans)
 457          | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
 458          | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
 459          | BOX (((t >>  0) & 0x3f), 1, s_SPtrans)
 460          | BOX (((t >>  8) & 0x3f), 3, s_SPtrans)
 461          | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
 462          | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
 463
 464       t = l ^ (l >> 16);
 465       u = t & E0;
 466       t = t & E1;
 467       u = u ^ (u << 16);
 468       u = u ^ l;
 469       u = u ^ Kc[j + 1];
 470       t = t ^ (t << 16);
 471       t = t ^ l;
 472       t = rotl32 (t, 28u);
 473       t = t ^ Kd[j + 1];
 474
 475       r ^= BOX (((u >>  0) & 0x3f), 0, s_SPtrans)
 476          | BOX (((u >>  8) & 0x3f), 2, s_SPtrans)
 477          | BOX (((u >> 16) & 0x3f), 4, s_SPtrans)
 478          | BOX (((u >> 24) & 0x3f), 6, s_SPtrans)
 479          | BOX (((t >>  0) & 0x3f), 1, s_SPtrans)
 480          | BOX (((t >>  8) & 0x3f), 3, s_SPtrans)
 481          | BOX (((t >> 16) & 0x3f), 5, s_SPtrans)
 482          | BOX (((t >> 24) & 0x3f), 7, s_SPtrans);
 483     }
 484
 485     u32x tt;
 486
 487     tt = l;
 488     l  = r;
 489     r  = tt;
 490   }
 491
 492   iv[0] = rotl32 (r, 31);
 493   iv[1] = rotl32 (l, 31);
 494 }
 495
 496 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 497 {
 498   /**
 499    * modifier
 500    */
 501
 502   const u32 lid = threadIdx.x;
 503
 504   /**
 505    * base
 506    */
 507
 508   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 509
 510   u32x pw_buf0[4];
 511
 512   pw_buf0[0] = pws[gid].i[ 0];
 513   pw_buf0[1] = pws[gid].i[ 1];
 514   pw_buf0[2] = 0;
 515   pw_buf0[3] = 0;
 516
 517   const u32 pw_len = pws[gid].pw_len;
 518
 519   /**
 520    * sbox, kbox
 521    */
 522
 523   __shared__ u32 s_skb[8][64];
 524   __shared__ u32 s_SPtrans[8][64];
 525
 526   if (lid < 64)
 527   {
 528     s_skb[0][lid] = c_skb[0][lid];
 529     s_skb[1][lid] = c_skb[1][lid];
 530     s_skb[2][lid] = c_skb[2][lid];
 531     s_skb[3][lid] = c_skb[3][lid];
 532     s_skb[4][lid] = c_skb[4][lid];
 533     s_skb[5][lid] = c_skb[5][lid];
 534     s_skb[6][lid] = c_skb[6][lid];
 535     s_skb[7][lid] = c_skb[7][lid];
 536
 537     s_SPtrans[0][lid] = c_SPtrans[0][lid];
 538     s_SPtrans[1][lid] = c_SPtrans[1][lid];
 539     s_SPtrans[2][lid] = c_SPtrans[2][lid];
 540     s_SPtrans[3][lid] = c_SPtrans[3][lid];
 541     s_SPtrans[4][lid] = c_SPtrans[4][lid];
 542     s_SPtrans[5][lid] = c_SPtrans[5][lid];
 543     s_SPtrans[6][lid] = c_SPtrans[6][lid];
 544     s_SPtrans[7][lid] = c_SPtrans[7][lid];
 545   }
 546
 547   __syncthreads ();
 548
 549   if (gid >= gid_max) return;
 550
 551   /**
 552    * salt
 553    */
 554
 555   const u32 mask = salt_bufs[salt_pos].salt_buf[0];
 556
 557   /**
 558    * main
 559    */
 560
 561   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 562   {
 563     u32x w0[4];
 564
 565     w0[0] = pw_buf0[0];
 566     w0[1] = pw_buf0[1];
 567     w0[2] = 0;
 568     w0[3] = 0;
 569
 570     u32x w1[4];
 571
 572     w1[0] = 0;
 573     w1[1] = 0;
 574     w1[2] = 0;
 575     w1[3] = 0;
 576
 577     u32x w2[4];
 578
 579     w2[0] = 0;
 580     w2[1] = 0;
 581     w2[2] = 0;
 582     w2[3] = 0;
 583
 584     u32x w3[4];
 585
 586     w3[0] = 0;
 587     w3[1] = 0;
 588     w3[2] = 0;
 589     w3[3] = 0;
 590
 591     u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 592
 593     out_len = (out_len >= 8) ? 8 : out_len;
 594
 595     u32x data[2];
 596
 597     data[0] = (w0[0] << 1) & 0xfefefefe;
 598     data[1] = (w0[1] << 1) & 0xfefefefe;
 599
 600     u32x Kc[16];
 601     u32x Kd[16];
 602
 603     _des_crypt_keysetup (data[0], data[1], Kc, Kd, s_skb);
 604
 605     u32x iv[2];
 606
 607     _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
 608
 609     const u32x r0 = iv[0];
 610     const u32x r1 = iv[1];
 611     const u32x r2 = 0;
 612     const u32x r3 = 0;
 613
 614     #include VECT_COMPARE_M
 615   }
 616 }
 617
 618 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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)
 619 {
 620 }
 621
 622 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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)
 623 {
 624 }
 625
 626 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
 627 {
 628   /**
 629    * modifier
 630    */
 631
 632   const u32 lid = threadIdx.x;
 633
 634   /**
 635    * base
 636    */
 637
 638   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
 639
 640   u32x pw_buf0[4];
 641
 642   pw_buf0[0] = pws[gid].i[ 0];
 643   pw_buf0[1] = pws[gid].i[ 1];
 644   pw_buf0[2] = 0;
 645   pw_buf0[3] = 0;
 646
 647   const u32 pw_len = pws[gid].pw_len;
 648
 649   /**
 650    * sbox, kbox
 651    */
 652
 653   __shared__ u32 s_skb[8][64];
 654   __shared__ u32 s_SPtrans[8][64];
 655
 656   if (lid < 64)
 657   {
 658     s_skb[0][lid] = c_skb[0][lid];
 659     s_skb[1][lid] = c_skb[1][lid];
 660     s_skb[2][lid] = c_skb[2][lid];
 661     s_skb[3][lid] = c_skb[3][lid];
 662     s_skb[4][lid] = c_skb[4][lid];
 663     s_skb[5][lid] = c_skb[5][lid];
 664     s_skb[6][lid] = c_skb[6][lid];
 665     s_skb[7][lid] = c_skb[7][lid];
 666
 667     s_SPtrans[0][lid] = c_SPtrans[0][lid];
 668     s_SPtrans[1][lid] = c_SPtrans[1][lid];
 669     s_SPtrans[2][lid] = c_SPtrans[2][lid];
 670     s_SPtrans[3][lid] = c_SPtrans[3][lid];
 671     s_SPtrans[4][lid] = c_SPtrans[4][lid];
 672     s_SPtrans[5][lid] = c_SPtrans[5][lid];
 673     s_SPtrans[6][lid] = c_SPtrans[6][lid];
 674     s_SPtrans[7][lid] = c_SPtrans[7][lid];
 675   }
 676
 677   __syncthreads ();
 678
 679   if (gid >= gid_max) return;
 680
 681   /**
 682    * salt
 683    */
 684
 685   const u32 mask = salt_bufs[salt_pos].salt_buf[0];
 686
 687   /**
 688    * digest
 689    */
 690
 691   const u32 search[4] =
 692   {
 693     digests_buf[digests_offset].digest_buf[DGST_R0],
 694     digests_buf[digests_offset].digest_buf[DGST_R1],
 695     digests_buf[digests_offset].digest_buf[DGST_R2],
 696     digests_buf[digests_offset].digest_buf[DGST_R3]
 697   };
 698
 699   /**
 700    * main
 701    */
 702
 703   for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++)
 704   {
 705     u32x w0[4];
 706
 707     w0[0] = pw_buf0[0];
 708     w0[1] = pw_buf0[1];
 709     w0[2] = 0;
 710     w0[3] = 0;
 711
 712     u32x w1[4];
 713
 714     w1[0] = 0;
 715     w1[1] = 0;
 716     w1[2] = 0;
 717     w1[3] = 0;
 718
 719     u32x w2[4];
 720
 721     w2[0] = 0;
 722     w2[1] = 0;
 723     w2[2] = 0;
 724     w2[3] = 0;
 725
 726     u32x w3[4];
 727
 728     w3[0] = 0;
 729     w3[1] = 0;
 730     w3[2] = 0;
 731     w3[3] = 0;
 732
 733     u32 out_len = apply_rules (c_rules[il_pos].cmds, w0, w1, pw_len);
 734
 735     out_len = (out_len >= 8) ? 8 : out_len;
 736
 737     u32x data[2];
 738
 739     data[0] = (w0[0] << 1) & 0xfefefefe;
 740     data[1] = (w0[1] << 1) & 0xfefefefe;
 741
 742     u32x Kc[16];
 743     u32x Kd[16];
 744
 745     _des_crypt_keysetup (data[0], data[1], Kc, Kd, s_skb);
 746
 747     u32x iv[2];
 748
 749     _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
 750
 751     const u32x r0 = iv[0];
 752     const u32x r1 = iv[1];
 753     const u32x r2 = 0;
 754     const u32x r3 = 0;
 755
 756     #include VECT_COMPARE_S
 757   }
 758 }
 759
 760 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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)
 761 {
 762 }
 763
 764 extern "C" __global__ void __launch_bounds__ (256, 1) m01500_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)
 765 {
 766 }