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