nv/m08500_a3.cu

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