nv/m05500_a3.cu

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