nv/m11300.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SHA512_
   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 __device__ __constant__ u64 k[80] =
  29 {
  30   SHA512C00, SHA512C01, SHA512C02, SHA512C03,
  31   SHA512C04, SHA512C05, SHA512C06, SHA512C07,
  32   SHA512C08, SHA512C09, SHA512C0a, SHA512C0b,
  33   SHA512C0c, SHA512C0d, SHA512C0e, SHA512C0f,
  34   SHA512C10, SHA512C11, SHA512C12, SHA512C13,
  35   SHA512C14, SHA512C15, SHA512C16, SHA512C17,
  36   SHA512C18, SHA512C19, SHA512C1a, SHA512C1b,
  37   SHA512C1c, SHA512C1d, SHA512C1e, SHA512C1f,
  38   SHA512C20, SHA512C21, SHA512C22, SHA512C23,
  39   SHA512C24, SHA512C25, SHA512C26, SHA512C27,
  40   SHA512C28, SHA512C29, SHA512C2a, SHA512C2b,
  41   SHA512C2c, SHA512C2d, SHA512C2e, SHA512C2f,
  42   SHA512C30, SHA512C31, SHA512C32, SHA512C33,
  43   SHA512C34, SHA512C35, SHA512C36, SHA512C37,
  44   SHA512C38, SHA512C39, SHA512C3a, SHA512C3b,
  45   SHA512C3c, SHA512C3d, SHA512C3e, SHA512C3f,
  46   SHA512C40, SHA512C41, SHA512C42, SHA512C43,
  47   SHA512C44, SHA512C45, SHA512C46, SHA512C47,
  48   SHA512C48, SHA512C49, SHA512C4a, SHA512C4b,
  49   SHA512C4c, SHA512C4d, SHA512C4e, SHA512C4f,
  50 };
  51
  52 #define ROUND_EXPAND()                            \
  53 {                                                 \
  54   w0_t = SHA512_EXPAND (we_t, w9_t, w1_t, w0_t);  \
  55   w1_t = SHA512_EXPAND (wf_t, wa_t, w2_t, w1_t);  \
  56   w2_t = SHA512_EXPAND (w0_t, wb_t, w3_t, w2_t);  \
  57   w3_t = SHA512_EXPAND (w1_t, wc_t, w4_t, w3_t);  \
  58   w4_t = SHA512_EXPAND (w2_t, wd_t, w5_t, w4_t);  \
  59   w5_t = SHA512_EXPAND (w3_t, we_t, w6_t, w5_t);  \
  60   w6_t = SHA512_EXPAND (w4_t, wf_t, w7_t, w6_t);  \
  61   w7_t = SHA512_EXPAND (w5_t, w0_t, w8_t, w7_t);  \
  62   w8_t = SHA512_EXPAND (w6_t, w1_t, w9_t, w8_t);  \
  63   w9_t = SHA512_EXPAND (w7_t, w2_t, wa_t, w9_t);  \
  64   wa_t = SHA512_EXPAND (w8_t, w3_t, wb_t, wa_t);  \
  65   wb_t = SHA512_EXPAND (w9_t, w4_t, wc_t, wb_t);  \
  66   wc_t = SHA512_EXPAND (wa_t, w5_t, wd_t, wc_t);  \
  67   wd_t = SHA512_EXPAND (wb_t, w6_t, we_t, wd_t);  \
  68   we_t = SHA512_EXPAND (wc_t, w7_t, wf_t, we_t);  \
  69   wf_t = SHA512_EXPAND (wd_t, w8_t, w0_t, wf_t);  \
  70 }
  71
  72 #define ROUND_STEP(i)                                                             \
  73 {                                                                                 \
  74   SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w0_t, k[i +  0]);  \
  75   SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w1_t, k[i +  1]);  \
  76   SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, w2_t, k[i +  2]);  \
  77   SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, w3_t, k[i +  3]);  \
  78   SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, w4_t, k[i +  4]);  \
  79   SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, w5_t, k[i +  5]);  \
  80   SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, w6_t, k[i +  6]);  \
  81   SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, w7_t, k[i +  7]);  \
  82   SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w8_t, k[i +  8]);  \
  83   SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w9_t, k[i +  9]);  \
  84   SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, wa_t, k[i + 10]);  \
  85   SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, wb_t, k[i + 11]);  \
  86   SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, wc_t, k[i + 12]);  \
  87   SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, wd_t, k[i + 13]);  \
  88   SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, we_t, k[i + 14]);  \
  89   SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, wf_t, k[i + 15]);  \
  90 }
  91
  92 __device__ static void sha512_transform (const u64 w[16], u64 dgst[8])
  93 {
  94   u64 a = dgst[0];
  95   u64 b = dgst[1];
  96   u64 c = dgst[2];
  97   u64 d = dgst[3];
  98   u64 e = dgst[4];
  99   u64 f = dgst[5];
 100   u64 g = dgst[6];
 101   u64 h = dgst[7];
 102
 103   u64 w0_t = w[ 0];
 104   u64 w1_t = w[ 1];
 105   u64 w2_t = w[ 2];
 106   u64 w3_t = w[ 3];
 107   u64 w4_t = w[ 4];
 108   u64 w5_t = w[ 5];
 109   u64 w6_t = w[ 6];
 110   u64 w7_t = w[ 7];
 111   u64 w8_t = w[ 8];
 112   u64 w9_t = w[ 9];
 113   u64 wa_t = w[10];
 114   u64 wb_t = w[11];
 115   u64 wc_t = w[12];
 116   u64 wd_t = w[13];
 117   u64 we_t = w[14];
 118   u64 wf_t = w[15];
 119
 120   ROUND_STEP (0);
 121
 122   for (int i = 16; i < 80; i += 16)
 123   {
 124     ROUND_EXPAND (); ROUND_STEP (i);
 125   }
 126
 127   dgst[0] += a;
 128   dgst[1] += b;
 129   dgst[2] += c;
 130   dgst[3] += d;
 131   dgst[4] += e;
 132   dgst[5] += f;
 133   dgst[6] += g;
 134   dgst[7] += h;
 135 }
 136
 137 __device__ __constant__ u32 te0[256] =
 138 {
 139   0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
 140   0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
 141   0x60303050, 0x02010103, 0xce6767a9, 0x562b2b7d,
 142   0xe7fefe19, 0xb5d7d762, 0x4dababe6, 0xec76769a,
 143   0x8fcaca45, 0x1f82829d, 0x89c9c940, 0xfa7d7d87,
 144   0xeffafa15, 0xb25959eb, 0x8e4747c9, 0xfbf0f00b,
 145   0x41adadec, 0xb3d4d467, 0x5fa2a2fd, 0x45afafea,
 146   0x239c9cbf, 0x53a4a4f7, 0xe4727296, 0x9bc0c05b,
 147   0x75b7b7c2, 0xe1fdfd1c, 0x3d9393ae, 0x4c26266a,
 148   0x6c36365a, 0x7e3f3f41, 0xf5f7f702, 0x83cccc4f,
 149   0x6834345c, 0x51a5a5f4, 0xd1e5e534, 0xf9f1f108,
 150   0xe2717193, 0xabd8d873, 0x62313153, 0x2a15153f,
 151   0x0804040c, 0x95c7c752, 0x46232365, 0x9dc3c35e,
 152   0x30181828, 0x379696a1, 0x0a05050f, 0x2f9a9ab5,
 153   0x0e070709, 0x24121236, 0x1b80809b, 0xdfe2e23d,
 154   0xcdebeb26, 0x4e272769, 0x7fb2b2cd, 0xea75759f,
 155   0x1209091b, 0x1d83839e, 0x582c2c74, 0x341a1a2e,
 156   0x361b1b2d, 0xdc6e6eb2, 0xb45a5aee, 0x5ba0a0fb,
 157   0xa45252f6, 0x763b3b4d, 0xb7d6d661, 0x7db3b3ce,
 158   0x5229297b, 0xdde3e33e, 0x5e2f2f71, 0x13848497,
 159   0xa65353f5, 0xb9d1d168, 0x00000000, 0xc1eded2c,
 160   0x40202060, 0xe3fcfc1f, 0x79b1b1c8, 0xb65b5bed,
 161   0xd46a6abe, 0x8dcbcb46, 0x67bebed9, 0x7239394b,
 162   0x944a4ade, 0x984c4cd4, 0xb05858e8, 0x85cfcf4a,
 163   0xbbd0d06b, 0xc5efef2a, 0x4faaaae5, 0xedfbfb16,
 164   0x864343c5, 0x9a4d4dd7, 0x66333355, 0x11858594,
 165   0x8a4545cf, 0xe9f9f910, 0x04020206, 0xfe7f7f81,
 166   0xa05050f0, 0x783c3c44, 0x259f9fba, 0x4ba8a8e3,
 167   0xa25151f3, 0x5da3a3fe, 0x804040c0, 0x058f8f8a,
 168   0x3f9292ad, 0x219d9dbc, 0x70383848, 0xf1f5f504,
 169   0x63bcbcdf, 0x77b6b6c1, 0xafdada75, 0x42212163,
 170   0x20101030, 0xe5ffff1a, 0xfdf3f30e, 0xbfd2d26d,
 171   0x81cdcd4c, 0x180c0c14, 0x26131335, 0xc3ecec2f,
 172   0xbe5f5fe1, 0x359797a2, 0x884444cc, 0x2e171739,
 173   0x93c4c457, 0x55a7a7f2, 0xfc7e7e82, 0x7a3d3d47,
 174   0xc86464ac, 0xba5d5de7, 0x3219192b, 0xe6737395,
 175   0xc06060a0, 0x19818198, 0x9e4f4fd1, 0xa3dcdc7f,
 176   0x44222266, 0x542a2a7e, 0x3b9090ab, 0x0b888883,
 177   0x8c4646ca, 0xc7eeee29, 0x6bb8b8d3, 0x2814143c,
 178   0xa7dede79, 0xbc5e5ee2, 0x160b0b1d, 0xaddbdb76,
 179   0xdbe0e03b, 0x64323256, 0x743a3a4e, 0x140a0a1e,
 180   0x924949db, 0x0c06060a, 0x4824246c, 0xb85c5ce4,
 181   0x9fc2c25d, 0xbdd3d36e, 0x43acacef, 0xc46262a6,
 182   0x399191a8, 0x319595a4, 0xd3e4e437, 0xf279798b,
 183   0xd5e7e732, 0x8bc8c843, 0x6e373759, 0xda6d6db7,
 184   0x018d8d8c, 0xb1d5d564, 0x9c4e4ed2, 0x49a9a9e0,
 185   0xd86c6cb4, 0xac5656fa, 0xf3f4f407, 0xcfeaea25,
 186   0xca6565af, 0xf47a7a8e, 0x47aeaee9, 0x10080818,
 187   0x6fbabad5, 0xf0787888, 0x4a25256f, 0x5c2e2e72,
 188   0x381c1c24, 0x57a6a6f1, 0x73b4b4c7, 0x97c6c651,
 189   0xcbe8e823, 0xa1dddd7c, 0xe874749c, 0x3e1f1f21,
 190   0x964b4bdd, 0x61bdbddc, 0x0d8b8b86, 0x0f8a8a85,
 191   0xe0707090, 0x7c3e3e42, 0x71b5b5c4, 0xcc6666aa,
 192   0x904848d8, 0x06030305, 0xf7f6f601, 0x1c0e0e12,
 193   0xc26161a3, 0x6a35355f, 0xae5757f9, 0x69b9b9d0,
 194   0x17868691, 0x99c1c158, 0x3a1d1d27, 0x279e9eb9,
 195   0xd9e1e138, 0xebf8f813, 0x2b9898b3, 0x22111133,
 196   0xd26969bb, 0xa9d9d970, 0x078e8e89, 0x339494a7,
 197   0x2d9b9bb6, 0x3c1e1e22, 0x15878792, 0xc9e9e920,
 198   0x87cece49, 0xaa5555ff, 0x50282878, 0xa5dfdf7a,
 199   0x038c8c8f, 0x59a1a1f8, 0x09898980, 0x1a0d0d17,
 200   0x65bfbfda, 0xd7e6e631, 0x844242c6, 0xd06868b8,
 201   0x824141c3, 0x299999b0, 0x5a2d2d77, 0x1e0f0f11,
 202   0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
 203 };
 204
 205 __device__ __constant__ u32 te1[256] =
 206 {
 207   0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
 208   0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
 209   0x50603030, 0x03020101, 0xa9ce6767, 0x7d562b2b,
 210   0x19e7fefe, 0x62b5d7d7, 0xe64dabab, 0x9aec7676,
 211   0x458fcaca, 0x9d1f8282, 0x4089c9c9, 0x87fa7d7d,
 212   0x15effafa, 0xebb25959, 0xc98e4747, 0x0bfbf0f0,
 213   0xec41adad, 0x67b3d4d4, 0xfd5fa2a2, 0xea45afaf,
 214   0xbf239c9c, 0xf753a4a4, 0x96e47272, 0x5b9bc0c0,
 215   0xc275b7b7, 0x1ce1fdfd, 0xae3d9393, 0x6a4c2626,
 216   0x5a6c3636, 0x417e3f3f, 0x02f5f7f7, 0x4f83cccc,
 217   0x5c683434, 0xf451a5a5, 0x34d1e5e5, 0x08f9f1f1,
 218   0x93e27171, 0x73abd8d8, 0x53623131, 0x3f2a1515,
 219   0x0c080404, 0x5295c7c7, 0x65462323, 0x5e9dc3c3,
 220   0x28301818, 0xa1379696, 0x0f0a0505, 0xb52f9a9a,
 221   0x090e0707, 0x36241212, 0x9b1b8080, 0x3ddfe2e2,
 222   0x26cdebeb, 0x694e2727, 0xcd7fb2b2, 0x9fea7575,
 223   0x1b120909, 0x9e1d8383, 0x74582c2c, 0x2e341a1a,
 224   0x2d361b1b, 0xb2dc6e6e, 0xeeb45a5a, 0xfb5ba0a0,
 225   0xf6a45252, 0x4d763b3b, 0x61b7d6d6, 0xce7db3b3,
 226   0x7b522929, 0x3edde3e3, 0x715e2f2f, 0x97138484,
 227   0xf5a65353, 0x68b9d1d1, 0x00000000, 0x2cc1eded,
 228   0x60402020, 0x1fe3fcfc, 0xc879b1b1, 0xedb65b5b,
 229   0xbed46a6a, 0x468dcbcb, 0xd967bebe, 0x4b723939,
 230   0xde944a4a, 0xd4984c4c, 0xe8b05858, 0x4a85cfcf,
 231   0x6bbbd0d0, 0x2ac5efef, 0xe54faaaa, 0x16edfbfb,
 232   0xc5864343, 0xd79a4d4d, 0x55663333, 0x94118585,
 233   0xcf8a4545, 0x10e9f9f9, 0x06040202, 0x81fe7f7f,
 234   0xf0a05050, 0x44783c3c, 0xba259f9f, 0xe34ba8a8,
 235   0xf3a25151, 0xfe5da3a3, 0xc0804040, 0x8a058f8f,
 236   0xad3f9292, 0xbc219d9d, 0x48703838, 0x04f1f5f5,
 237   0xdf63bcbc, 0xc177b6b6, 0x75afdada, 0x63422121,
 238   0x30201010, 0x1ae5ffff, 0x0efdf3f3, 0x6dbfd2d2,
 239   0x4c81cdcd, 0x14180c0c, 0x35261313, 0x2fc3ecec,
 240   0xe1be5f5f, 0xa2359797, 0xcc884444, 0x392e1717,
 241   0x5793c4c4, 0xf255a7a7, 0x82fc7e7e, 0x477a3d3d,
 242   0xacc86464, 0xe7ba5d5d, 0x2b321919, 0x95e67373,
 243   0xa0c06060, 0x98198181, 0xd19e4f4f, 0x7fa3dcdc,
 244   0x66442222, 0x7e542a2a, 0xab3b9090, 0x830b8888,
 245   0xca8c4646, 0x29c7eeee, 0xd36bb8b8, 0x3c281414,
 246   0x79a7dede, 0xe2bc5e5e, 0x1d160b0b, 0x76addbdb,
 247   0x3bdbe0e0, 0x56643232, 0x4e743a3a, 0x1e140a0a,
 248   0xdb924949, 0x0a0c0606, 0x6c482424, 0xe4b85c5c,
 249   0x5d9fc2c2, 0x6ebdd3d3, 0xef43acac, 0xa6c46262,
 250   0xa8399191, 0xa4319595, 0x37d3e4e4, 0x8bf27979,
 251   0x32d5e7e7, 0x438bc8c8, 0x596e3737, 0xb7da6d6d,
 252   0x8c018d8d, 0x64b1d5d5, 0xd29c4e4e, 0xe049a9a9,
 253   0xb4d86c6c, 0xfaac5656, 0x07f3f4f4, 0x25cfeaea,
 254   0xafca6565, 0x8ef47a7a, 0xe947aeae, 0x18100808,
 255   0xd56fbaba, 0x88f07878, 0x6f4a2525, 0x725c2e2e,
 256   0x24381c1c, 0xf157a6a6, 0xc773b4b4, 0x5197c6c6,
 257   0x23cbe8e8, 0x7ca1dddd, 0x9ce87474, 0x213e1f1f,
 258   0xdd964b4b, 0xdc61bdbd, 0x860d8b8b, 0x850f8a8a,
 259   0x90e07070, 0x427c3e3e, 0xc471b5b5, 0xaacc6666,
 260   0xd8904848, 0x05060303, 0x01f7f6f6, 0x121c0e0e,
 261   0xa3c26161, 0x5f6a3535, 0xf9ae5757, 0xd069b9b9,
 262   0x91178686, 0x5899c1c1, 0x273a1d1d, 0xb9279e9e,
 263   0x38d9e1e1, 0x13ebf8f8, 0xb32b9898, 0x33221111,
 264   0xbbd26969, 0x70a9d9d9, 0x89078e8e, 0xa7339494,
 265   0xb62d9b9b, 0x223c1e1e, 0x92158787, 0x20c9e9e9,
 266   0x4987cece, 0xffaa5555, 0x78502828, 0x7aa5dfdf,
 267   0x8f038c8c, 0xf859a1a1, 0x80098989, 0x171a0d0d,
 268   0xda65bfbf, 0x31d7e6e6, 0xc6844242, 0xb8d06868,
 269   0xc3824141, 0xb0299999, 0x775a2d2d, 0x111e0f0f,
 270   0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
 271 };
 272
 273 __device__ __constant__ u32 te2[256] =
 274 {
 275   0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
 276   0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
 277   0x30506030, 0x01030201, 0x67a9ce67, 0x2b7d562b,
 278   0xfe19e7fe, 0xd762b5d7, 0xabe64dab, 0x769aec76,
 279   0xca458fca, 0x829d1f82, 0xc94089c9, 0x7d87fa7d,
 280   0xfa15effa, 0x59ebb259, 0x47c98e47, 0xf00bfbf0,
 281   0xadec41ad, 0xd467b3d4, 0xa2fd5fa2, 0xafea45af,
 282   0x9cbf239c, 0xa4f753a4, 0x7296e472, 0xc05b9bc0,
 283   0xb7c275b7, 0xfd1ce1fd, 0x93ae3d93, 0x266a4c26,
 284   0x365a6c36, 0x3f417e3f, 0xf702f5f7, 0xcc4f83cc,
 285   0x345c6834, 0xa5f451a5, 0xe534d1e5, 0xf108f9f1,
 286   0x7193e271, 0xd873abd8, 0x31536231, 0x153f2a15,
 287   0x040c0804, 0xc75295c7, 0x23654623, 0xc35e9dc3,
 288   0x18283018, 0x96a13796, 0x050f0a05, 0x9ab52f9a,
 289   0x07090e07, 0x12362412, 0x809b1b80, 0xe23ddfe2,
 290   0xeb26cdeb, 0x27694e27, 0xb2cd7fb2, 0x759fea75,
 291   0x091b1209, 0x839e1d83, 0x2c74582c, 0x1a2e341a,
 292   0x1b2d361b, 0x6eb2dc6e, 0x5aeeb45a, 0xa0fb5ba0,
 293   0x52f6a452, 0x3b4d763b, 0xd661b7d6, 0xb3ce7db3,
 294   0x297b5229, 0xe33edde3, 0x2f715e2f, 0x84971384,
 295   0x53f5a653, 0xd168b9d1, 0x00000000, 0xed2cc1ed,
 296   0x20604020, 0xfc1fe3fc, 0xb1c879b1, 0x5bedb65b,
 297   0x6abed46a, 0xcb468dcb, 0xbed967be, 0x394b7239,
 298   0x4ade944a, 0x4cd4984c, 0x58e8b058, 0xcf4a85cf,
 299   0xd06bbbd0, 0xef2ac5ef, 0xaae54faa, 0xfb16edfb,
 300   0x43c58643, 0x4dd79a4d, 0x33556633, 0x85941185,
 301   0x45cf8a45, 0xf910e9f9, 0x02060402, 0x7f81fe7f,
 302   0x50f0a050, 0x3c44783c, 0x9fba259f, 0xa8e34ba8,
 303   0x51f3a251, 0xa3fe5da3, 0x40c08040, 0x8f8a058f,
 304   0x92ad3f92, 0x9dbc219d, 0x38487038, 0xf504f1f5,
 305   0xbcdf63bc, 0xb6c177b6, 0xda75afda, 0x21634221,
 306   0x10302010, 0xff1ae5ff, 0xf30efdf3, 0xd26dbfd2,
 307   0xcd4c81cd, 0x0c14180c, 0x13352613, 0xec2fc3ec,
 308   0x5fe1be5f, 0x97a23597, 0x44cc8844, 0x17392e17,
 309   0xc45793c4, 0xa7f255a7, 0x7e82fc7e, 0x3d477a3d,
 310   0x64acc864, 0x5de7ba5d, 0x192b3219, 0x7395e673,
 311   0x60a0c060, 0x81981981, 0x4fd19e4f, 0xdc7fa3dc,
 312   0x22664422, 0x2a7e542a, 0x90ab3b90, 0x88830b88,
 313   0x46ca8c46, 0xee29c7ee, 0xb8d36bb8, 0x143c2814,
 314   0xde79a7de, 0x5ee2bc5e, 0x0b1d160b, 0xdb76addb,
 315   0xe03bdbe0, 0x32566432, 0x3a4e743a, 0x0a1e140a,
 316   0x49db9249, 0x060a0c06, 0x246c4824, 0x5ce4b85c,
 317   0xc25d9fc2, 0xd36ebdd3, 0xacef43ac, 0x62a6c462,
 318   0x91a83991, 0x95a43195, 0xe437d3e4, 0x798bf279,
 319   0xe732d5e7, 0xc8438bc8, 0x37596e37, 0x6db7da6d,
 320   0x8d8c018d, 0xd564b1d5, 0x4ed29c4e, 0xa9e049a9,
 321   0x6cb4d86c, 0x56faac56, 0xf407f3f4, 0xea25cfea,
 322   0x65afca65, 0x7a8ef47a, 0xaee947ae, 0x08181008,
 323   0xbad56fba, 0x7888f078, 0x256f4a25, 0x2e725c2e,
 324   0x1c24381c, 0xa6f157a6, 0xb4c773b4, 0xc65197c6,
 325   0xe823cbe8, 0xdd7ca1dd, 0x749ce874, 0x1f213e1f,
 326   0x4bdd964b, 0xbddc61bd, 0x8b860d8b, 0x8a850f8a,
 327   0x7090e070, 0x3e427c3e, 0xb5c471b5, 0x66aacc66,
 328   0x48d89048, 0x03050603, 0xf601f7f6, 0x0e121c0e,
 329   0x61a3c261, 0x355f6a35, 0x57f9ae57, 0xb9d069b9,
 330   0x86911786, 0xc15899c1, 0x1d273a1d, 0x9eb9279e,
 331   0xe138d9e1, 0xf813ebf8, 0x98b32b98, 0x11332211,
 332   0x69bbd269, 0xd970a9d9, 0x8e89078e, 0x94a73394,
 333   0x9bb62d9b, 0x1e223c1e, 0x87921587, 0xe920c9e9,
 334   0xce4987ce, 0x55ffaa55, 0x28785028, 0xdf7aa5df,
 335   0x8c8f038c, 0xa1f859a1, 0x89800989, 0x0d171a0d,
 336   0xbfda65bf, 0xe631d7e6, 0x42c68442, 0x68b8d068,
 337   0x41c38241, 0x99b02999, 0x2d775a2d, 0x0f111e0f,
 338   0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
 339 };
 340
 341 __device__ __constant__ u32 te3[256] =
 342 {
 343   0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
 344   0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
 345   0x30305060, 0x01010302, 0x6767a9ce, 0x2b2b7d56,
 346   0xfefe19e7, 0xd7d762b5, 0xababe64d, 0x76769aec,
 347   0xcaca458f, 0x82829d1f, 0xc9c94089, 0x7d7d87fa,
 348   0xfafa15ef, 0x5959ebb2, 0x4747c98e, 0xf0f00bfb,
 349   0xadadec41, 0xd4d467b3, 0xa2a2fd5f, 0xafafea45,
 350   0x9c9cbf23, 0xa4a4f753, 0x727296e4, 0xc0c05b9b,
 351   0xb7b7c275, 0xfdfd1ce1, 0x9393ae3d, 0x26266a4c,
 352   0x36365a6c, 0x3f3f417e, 0xf7f702f5, 0xcccc4f83,
 353   0x34345c68, 0xa5a5f451, 0xe5e534d1, 0xf1f108f9,
 354   0x717193e2, 0xd8d873ab, 0x31315362, 0x15153f2a,
 355   0x04040c08, 0xc7c75295, 0x23236546, 0xc3c35e9d,
 356   0x18182830, 0x9696a137, 0x05050f0a, 0x9a9ab52f,
 357   0x0707090e, 0x12123624, 0x80809b1b, 0xe2e23ddf,
 358   0xebeb26cd, 0x2727694e, 0xb2b2cd7f, 0x75759fea,
 359   0x09091b12, 0x83839e1d, 0x2c2c7458, 0x1a1a2e34,
 360   0x1b1b2d36, 0x6e6eb2dc, 0x5a5aeeb4, 0xa0a0fb5b,
 361   0x5252f6a4, 0x3b3b4d76, 0xd6d661b7, 0xb3b3ce7d,
 362   0x29297b52, 0xe3e33edd, 0x2f2f715e, 0x84849713,
 363   0x5353f5a6, 0xd1d168b9, 0x00000000, 0xeded2cc1,
 364   0x20206040, 0xfcfc1fe3, 0xb1b1c879, 0x5b5bedb6,
 365   0x6a6abed4, 0xcbcb468d, 0xbebed967, 0x39394b72,
 366   0x4a4ade94, 0x4c4cd498, 0x5858e8b0, 0xcfcf4a85,
 367   0xd0d06bbb, 0xefef2ac5, 0xaaaae54f, 0xfbfb16ed,
 368   0x4343c586, 0x4d4dd79a, 0x33335566, 0x85859411,
 369   0x4545cf8a, 0xf9f910e9, 0x02020604, 0x7f7f81fe,
 370   0x5050f0a0, 0x3c3c4478, 0x9f9fba25, 0xa8a8e34b,
 371   0x5151f3a2, 0xa3a3fe5d, 0x4040c080, 0x8f8f8a05,
 372   0x9292ad3f, 0x9d9dbc21, 0x38384870, 0xf5f504f1,
 373   0xbcbcdf63, 0xb6b6c177, 0xdada75af, 0x21216342,
 374   0x10103020, 0xffff1ae5, 0xf3f30efd, 0xd2d26dbf,
 375   0xcdcd4c81, 0x0c0c1418, 0x13133526, 0xecec2fc3,
 376   0x5f5fe1be, 0x9797a235, 0x4444cc88, 0x1717392e,
 377   0xc4c45793, 0xa7a7f255, 0x7e7e82fc, 0x3d3d477a,
 378   0x6464acc8, 0x5d5de7ba, 0x19192b32, 0x737395e6,
 379   0x6060a0c0, 0x81819819, 0x4f4fd19e, 0xdcdc7fa3,
 380   0x22226644, 0x2a2a7e54, 0x9090ab3b, 0x8888830b,
 381   0x4646ca8c, 0xeeee29c7, 0xb8b8d36b, 0x14143c28,
 382   0xdede79a7, 0x5e5ee2bc, 0x0b0b1d16, 0xdbdb76ad,
 383   0xe0e03bdb, 0x32325664, 0x3a3a4e74, 0x0a0a1e14,
 384   0x4949db92, 0x06060a0c, 0x24246c48, 0x5c5ce4b8,
 385   0xc2c25d9f, 0xd3d36ebd, 0xacacef43, 0x6262a6c4,
 386   0x9191a839, 0x9595a431, 0xe4e437d3, 0x79798bf2,
 387   0xe7e732d5, 0xc8c8438b, 0x3737596e, 0x6d6db7da,
 388   0x8d8d8c01, 0xd5d564b1, 0x4e4ed29c, 0xa9a9e049,
 389   0x6c6cb4d8, 0x5656faac, 0xf4f407f3, 0xeaea25cf,
 390   0x6565afca, 0x7a7a8ef4, 0xaeaee947, 0x08081810,
 391   0xbabad56f, 0x787888f0, 0x25256f4a, 0x2e2e725c,
 392   0x1c1c2438, 0xa6a6f157, 0xb4b4c773, 0xc6c65197,
 393   0xe8e823cb, 0xdddd7ca1, 0x74749ce8, 0x1f1f213e,
 394   0x4b4bdd96, 0xbdbddc61, 0x8b8b860d, 0x8a8a850f,
 395   0x707090e0, 0x3e3e427c, 0xb5b5c471, 0x6666aacc,
 396   0x4848d890, 0x03030506, 0xf6f601f7, 0x0e0e121c,
 397   0x6161a3c2, 0x35355f6a, 0x5757f9ae, 0xb9b9d069,
 398   0x86869117, 0xc1c15899, 0x1d1d273a, 0x9e9eb927,
 399   0xe1e138d9, 0xf8f813eb, 0x9898b32b, 0x11113322,
 400   0x6969bbd2, 0xd9d970a9, 0x8e8e8907, 0x9494a733,
 401   0x9b9bb62d, 0x1e1e223c, 0x87879215, 0xe9e920c9,
 402   0xcece4987, 0x5555ffaa, 0x28287850, 0xdfdf7aa5,
 403   0x8c8c8f03, 0xa1a1f859, 0x89898009, 0x0d0d171a,
 404   0xbfbfda65, 0xe6e631d7, 0x4242c684, 0x6868b8d0,
 405   0x4141c382, 0x9999b029, 0x2d2d775a, 0x0f0f111e,
 406   0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
 407 };
 408
 409 __device__ __constant__ u32 te4[256] =
 410 {
 411   0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
 412   0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
 413   0x30303030, 0x01010101, 0x67676767, 0x2b2b2b2b,
 414   0xfefefefe, 0xd7d7d7d7, 0xabababab, 0x76767676,
 415   0xcacacaca, 0x82828282, 0xc9c9c9c9, 0x7d7d7d7d,
 416   0xfafafafa, 0x59595959, 0x47474747, 0xf0f0f0f0,
 417   0xadadadad, 0xd4d4d4d4, 0xa2a2a2a2, 0xafafafaf,
 418   0x9c9c9c9c, 0xa4a4a4a4, 0x72727272, 0xc0c0c0c0,
 419   0xb7b7b7b7, 0xfdfdfdfd, 0x93939393, 0x26262626,
 420   0x36363636, 0x3f3f3f3f, 0xf7f7f7f7, 0xcccccccc,
 421   0x34343434, 0xa5a5a5a5, 0xe5e5e5e5, 0xf1f1f1f1,
 422   0x71717171, 0xd8d8d8d8, 0x31313131, 0x15151515,
 423   0x04040404, 0xc7c7c7c7, 0x23232323, 0xc3c3c3c3,
 424   0x18181818, 0x96969696, 0x05050505, 0x9a9a9a9a,
 425   0x07070707, 0x12121212, 0x80808080, 0xe2e2e2e2,
 426   0xebebebeb, 0x27272727, 0xb2b2b2b2, 0x75757575,
 427   0x09090909, 0x83838383, 0x2c2c2c2c, 0x1a1a1a1a,
 428   0x1b1b1b1b, 0x6e6e6e6e, 0x5a5a5a5a, 0xa0a0a0a0,
 429   0x52525252, 0x3b3b3b3b, 0xd6d6d6d6, 0xb3b3b3b3,
 430   0x29292929, 0xe3e3e3e3, 0x2f2f2f2f, 0x84848484,
 431   0x53535353, 0xd1d1d1d1, 0x00000000, 0xedededed,
 432   0x20202020, 0xfcfcfcfc, 0xb1b1b1b1, 0x5b5b5b5b,
 433   0x6a6a6a6a, 0xcbcbcbcb, 0xbebebebe, 0x39393939,
 434   0x4a4a4a4a, 0x4c4c4c4c, 0x58585858, 0xcfcfcfcf,
 435   0xd0d0d0d0, 0xefefefef, 0xaaaaaaaa, 0xfbfbfbfb,
 436   0x43434343, 0x4d4d4d4d, 0x33333333, 0x85858585,
 437   0x45454545, 0xf9f9f9f9, 0x02020202, 0x7f7f7f7f,
 438   0x50505050, 0x3c3c3c3c, 0x9f9f9f9f, 0xa8a8a8a8,
 439   0x51515151, 0xa3a3a3a3, 0x40404040, 0x8f8f8f8f,
 440   0x92929292, 0x9d9d9d9d, 0x38383838, 0xf5f5f5f5,
 441   0xbcbcbcbc, 0xb6b6b6b6, 0xdadadada, 0x21212121,
 442   0x10101010, 0xffffffff, 0xf3f3f3f3, 0xd2d2d2d2,
 443   0xcdcdcdcd, 0x0c0c0c0c, 0x13131313, 0xecececec,
 444   0x5f5f5f5f, 0x97979797, 0x44444444, 0x17171717,
 445   0xc4c4c4c4, 0xa7a7a7a7, 0x7e7e7e7e, 0x3d3d3d3d,
 446   0x64646464, 0x5d5d5d5d, 0x19191919, 0x73737373,
 447   0x60606060, 0x81818181, 0x4f4f4f4f, 0xdcdcdcdc,
 448   0x22222222, 0x2a2a2a2a, 0x90909090, 0x88888888,
 449   0x46464646, 0xeeeeeeee, 0xb8b8b8b8, 0x14141414,
 450   0xdededede, 0x5e5e5e5e, 0x0b0b0b0b, 0xdbdbdbdb,
 451   0xe0e0e0e0, 0x32323232, 0x3a3a3a3a, 0x0a0a0a0a,
 452   0x49494949, 0x06060606, 0x24242424, 0x5c5c5c5c,
 453   0xc2c2c2c2, 0xd3d3d3d3, 0xacacacac, 0x62626262,
 454   0x91919191, 0x95959595, 0xe4e4e4e4, 0x79797979,
 455   0xe7e7e7e7, 0xc8c8c8c8, 0x37373737, 0x6d6d6d6d,
 456   0x8d8d8d8d, 0xd5d5d5d5, 0x4e4e4e4e, 0xa9a9a9a9,
 457   0x6c6c6c6c, 0x56565656, 0xf4f4f4f4, 0xeaeaeaea,
 458   0x65656565, 0x7a7a7a7a, 0xaeaeaeae, 0x08080808,
 459   0xbabababa, 0x78787878, 0x25252525, 0x2e2e2e2e,
 460   0x1c1c1c1c, 0xa6a6a6a6, 0xb4b4b4b4, 0xc6c6c6c6,
 461   0xe8e8e8e8, 0xdddddddd, 0x74747474, 0x1f1f1f1f,
 462   0x4b4b4b4b, 0xbdbdbdbd, 0x8b8b8b8b, 0x8a8a8a8a,
 463   0x70707070, 0x3e3e3e3e, 0xb5b5b5b5, 0x66666666,
 464   0x48484848, 0x03030303, 0xf6f6f6f6, 0x0e0e0e0e,
 465   0x61616161, 0x35353535, 0x57575757, 0xb9b9b9b9,
 466   0x86868686, 0xc1c1c1c1, 0x1d1d1d1d, 0x9e9e9e9e,
 467   0xe1e1e1e1, 0xf8f8f8f8, 0x98989898, 0x11111111,
 468   0x69696969, 0xd9d9d9d9, 0x8e8e8e8e, 0x94949494,
 469   0x9b9b9b9b, 0x1e1e1e1e, 0x87878787, 0xe9e9e9e9,
 470   0xcececece, 0x55555555, 0x28282828, 0xdfdfdfdf,
 471   0x8c8c8c8c, 0xa1a1a1a1, 0x89898989, 0x0d0d0d0d,
 472   0xbfbfbfbf, 0xe6e6e6e6, 0x42424242, 0x68686868,
 473   0x41414141, 0x99999999, 0x2d2d2d2d, 0x0f0f0f0f,
 474   0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
 475 };
 476
 477 __device__ __constant__ u32 td0[256] =
 478 {
 479   0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
 480   0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
 481   0x2030fa55, 0xad766df6, 0x88cc7691, 0xf5024c25,
 482   0x4fe5d7fc, 0xc52acbd7, 0x26354480, 0xb562a38f,
 483   0xdeb15a49, 0x25ba1b67, 0x45ea0e98, 0x5dfec0e1,
 484   0xc32f7502, 0x814cf012, 0x8d4697a3, 0x6bd3f9c6,
 485   0x038f5fe7, 0x15929c95, 0xbf6d7aeb, 0x955259da,
 486   0xd4be832d, 0x587421d3, 0x49e06929, 0x8ec9c844,
 487   0x75c2896a, 0xf48e7978, 0x99583e6b, 0x27b971dd,
 488   0xbee14fb6, 0xf088ad17, 0xc920ac66, 0x7dce3ab4,
 489   0x63df4a18, 0xe51a3182, 0x97513360, 0x62537f45,
 490   0xb16477e0, 0xbb6bae84, 0xfe81a01c, 0xf9082b94,
 491   0x70486858, 0x8f45fd19, 0x94de6c87, 0x527bf8b7,
 492   0xab73d323, 0x724b02e2, 0xe31f8f57, 0x6655ab2a,
 493   0xb2eb2807, 0x2fb5c203, 0x86c57b9a, 0xd33708a5,
 494   0x302887f2, 0x23bfa5b2, 0x02036aba, 0xed16825c,
 495   0x8acf1c2b, 0xa779b492, 0xf307f2f0, 0x4e69e2a1,
 496   0x65daf4cd, 0x0605bed5, 0xd134621f, 0xc4a6fe8a,
 497   0x342e539d, 0xa2f355a0, 0x058ae132, 0xa4f6eb75,
 498   0x0b83ec39, 0x4060efaa, 0x5e719f06, 0xbd6e1051,
 499   0x3e218af9, 0x96dd063d, 0xdd3e05ae, 0x4de6bd46,
 500   0x91548db5, 0x71c45d05, 0x0406d46f, 0x605015ff,
 501   0x1998fb24, 0xd6bde997, 0x894043cc, 0x67d99e77,
 502   0xb0e842bd, 0x07898b88, 0xe7195b38, 0x79c8eedb,
 503   0xa17c0a47, 0x7c420fe9, 0xf8841ec9, 0x00000000,
 504   0x09808683, 0x322bed48, 0x1e1170ac, 0x6c5a724e,
 505   0xfd0efffb, 0x0f853856, 0x3daed51e, 0x362d3927,
 506   0x0a0fd964, 0x685ca621, 0x9b5b54d1, 0x24362e3a,
 507   0x0c0a67b1, 0x9357e70f, 0xb4ee96d2, 0x1b9b919e,
 508   0x80c0c54f, 0x61dc20a2, 0x5a774b69, 0x1c121a16,
 509   0xe293ba0a, 0xc0a02ae5, 0x3c22e043, 0x121b171d,
 510   0x0e090d0b, 0xf28bc7ad, 0x2db6a8b9, 0x141ea9c8,
 511   0x57f11985, 0xaf75074c, 0xee99ddbb, 0xa37f60fd,
 512   0xf701269f, 0x5c72f5bc, 0x44663bc5, 0x5bfb7e34,
 513   0x8b432976, 0xcb23c6dc, 0xb6edfc68, 0xb8e4f163,
 514   0xd731dcca, 0x42638510, 0x13972240, 0x84c61120,
 515   0x854a247d, 0xd2bb3df8, 0xaef93211, 0xc729a16d,
 516   0x1d9e2f4b, 0xdcb230f3, 0x0d8652ec, 0x77c1e3d0,
 517   0x2bb3166c, 0xa970b999, 0x119448fa, 0x47e96422,
 518   0xa8fc8cc4, 0xa0f03f1a, 0x567d2cd8, 0x223390ef,
 519   0x87494ec7, 0xd938d1c1, 0x8ccaa2fe, 0x98d40b36,
 520   0xa6f581cf, 0xa57ade28, 0xdab78e26, 0x3fadbfa4,
 521   0x2c3a9de4, 0x5078920d, 0x6a5fcc9b, 0x547e4662,
 522   0xf68d13c2, 0x90d8b8e8, 0x2e39f75e, 0x82c3aff5,
 523   0x9f5d80be, 0x69d0937c, 0x6fd52da9, 0xcf2512b3,
 524   0xc8ac993b, 0x10187da7, 0xe89c636e, 0xdb3bbb7b,
 525   0xcd267809, 0x6e5918f4, 0xec9ab701, 0x834f9aa8,
 526   0xe6956e65, 0xaaffe67e, 0x21bccf08, 0xef15e8e6,
 527   0xbae79bd9, 0x4a6f36ce, 0xea9f09d4, 0x29b07cd6,
 528   0x31a4b2af, 0x2a3f2331, 0xc6a59430, 0x35a266c0,
 529   0x744ebc37, 0xfc82caa6, 0xe090d0b0, 0x33a7d815,
 530   0xf104984a, 0x41ecdaf7, 0x7fcd500e, 0x1791f62f,
 531   0x764dd68d, 0x43efb04d, 0xccaa4d54, 0xe49604df,
 532   0x9ed1b5e3, 0x4c6a881b, 0xc12c1fb8, 0x4665517f,
 533   0x9d5eea04, 0x018c355d, 0xfa877473, 0xfb0b412e,
 534   0xb3671d5a, 0x92dbd252, 0xe9105633, 0x6dd64713,
 535   0x9ad7618c, 0x37a10c7a, 0x59f8148e, 0xeb133c89,
 536   0xcea927ee, 0xb761c935, 0xe11ce5ed, 0x7a47b13c,
 537   0x9cd2df59, 0x55f2733f, 0x1814ce79, 0x73c737bf,
 538   0x53f7cdea, 0x5ffdaa5b, 0xdf3d6f14, 0x7844db86,
 539   0xcaaff381, 0xb968c43e, 0x3824342c, 0xc2a3405f,
 540   0x161dc372, 0xbce2250c, 0x283c498b, 0xff0d9541,
 541   0x39a80171, 0x080cb3de, 0xd8b4e49c, 0x6456c190,
 542   0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
 543 };
 544
 545 __device__ __constant__ u32 td1[256] =
 546 {
 547   0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
 548   0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
 549   0x552030fa, 0xf6ad766d, 0x9188cc76, 0x25f5024c,
 550   0xfc4fe5d7, 0xd7c52acb, 0x80263544, 0x8fb562a3,
 551   0x49deb15a, 0x6725ba1b, 0x9845ea0e, 0xe15dfec0,
 552   0x02c32f75, 0x12814cf0, 0xa38d4697, 0xc66bd3f9,
 553   0xe7038f5f, 0x9515929c, 0xebbf6d7a, 0xda955259,
 554   0x2dd4be83, 0xd3587421, 0x2949e069, 0x448ec9c8,
 555   0x6a75c289, 0x78f48e79, 0x6b99583e, 0xdd27b971,
 556   0xb6bee14f, 0x17f088ad, 0x66c920ac, 0xb47dce3a,
 557   0x1863df4a, 0x82e51a31, 0x60975133, 0x4562537f,
 558   0xe0b16477, 0x84bb6bae, 0x1cfe81a0, 0x94f9082b,
 559   0x58704868, 0x198f45fd, 0x8794de6c, 0xb7527bf8,
 560   0x23ab73d3, 0xe2724b02, 0x57e31f8f, 0x2a6655ab,
 561   0x07b2eb28, 0x032fb5c2, 0x9a86c57b, 0xa5d33708,
 562   0xf2302887, 0xb223bfa5, 0xba02036a, 0x5ced1682,
 563   0x2b8acf1c, 0x92a779b4, 0xf0f307f2, 0xa14e69e2,
 564   0xcd65daf4, 0xd50605be, 0x1fd13462, 0x8ac4a6fe,
 565   0x9d342e53, 0xa0a2f355, 0x32058ae1, 0x75a4f6eb,
 566   0x390b83ec, 0xaa4060ef, 0x065e719f, 0x51bd6e10,
 567   0xf93e218a, 0x3d96dd06, 0xaedd3e05, 0x464de6bd,
 568   0xb591548d, 0x0571c45d, 0x6f0406d4, 0xff605015,
 569   0x241998fb, 0x97d6bde9, 0xcc894043, 0x7767d99e,
 570   0xbdb0e842, 0x8807898b, 0x38e7195b, 0xdb79c8ee,
 571   0x47a17c0a, 0xe97c420f, 0xc9f8841e, 0x00000000,
 572   0x83098086, 0x48322bed, 0xac1e1170, 0x4e6c5a72,
 573   0xfbfd0eff, 0x560f8538, 0x1e3daed5, 0x27362d39,
 574   0x640a0fd9, 0x21685ca6, 0xd19b5b54, 0x3a24362e,
 575   0xb10c0a67, 0x0f9357e7, 0xd2b4ee96, 0x9e1b9b91,
 576   0x4f80c0c5, 0xa261dc20, 0x695a774b, 0x161c121a,
 577   0x0ae293ba, 0xe5c0a02a, 0x433c22e0, 0x1d121b17,
 578   0x0b0e090d, 0xadf28bc7, 0xb92db6a8, 0xc8141ea9,
 579   0x8557f119, 0x4caf7507, 0xbbee99dd, 0xfda37f60,
 580   0x9ff70126, 0xbc5c72f5, 0xc544663b, 0x345bfb7e,
 581   0x768b4329, 0xdccb23c6, 0x68b6edfc, 0x63b8e4f1,
 582   0xcad731dc, 0x10426385, 0x40139722, 0x2084c611,
 583   0x7d854a24, 0xf8d2bb3d, 0x11aef932, 0x6dc729a1,
 584   0x4b1d9e2f, 0xf3dcb230, 0xec0d8652, 0xd077c1e3,
 585   0x6c2bb316, 0x99a970b9, 0xfa119448, 0x2247e964,
 586   0xc4a8fc8c, 0x1aa0f03f, 0xd8567d2c, 0xef223390,
 587   0xc787494e, 0xc1d938d1, 0xfe8ccaa2, 0x3698d40b,
 588   0xcfa6f581, 0x28a57ade, 0x26dab78e, 0xa43fadbf,
 589   0xe42c3a9d, 0x0d507892, 0x9b6a5fcc, 0x62547e46,
 590   0xc2f68d13, 0xe890d8b8, 0x5e2e39f7, 0xf582c3af,
 591   0xbe9f5d80, 0x7c69d093, 0xa96fd52d, 0xb3cf2512,
 592   0x3bc8ac99, 0xa710187d, 0x6ee89c63, 0x7bdb3bbb,
 593   0x09cd2678, 0xf46e5918, 0x01ec9ab7, 0xa8834f9a,
 594   0x65e6956e, 0x7eaaffe6, 0x0821bccf, 0xe6ef15e8,
 595   0xd9bae79b, 0xce4a6f36, 0xd4ea9f09, 0xd629b07c,
 596   0xaf31a4b2, 0x312a3f23, 0x30c6a594, 0xc035a266,
 597   0x37744ebc, 0xa6fc82ca, 0xb0e090d0, 0x1533a7d8,
 598   0x4af10498, 0xf741ecda, 0x0e7fcd50, 0x2f1791f6,
 599   0x8d764dd6, 0x4d43efb0, 0x54ccaa4d, 0xdfe49604,
 600   0xe39ed1b5, 0x1b4c6a88, 0xb8c12c1f, 0x7f466551,
 601   0x049d5eea, 0x5d018c35, 0x73fa8774, 0x2efb0b41,
 602   0x5ab3671d, 0x5292dbd2, 0x33e91056, 0x136dd647,
 603   0x8c9ad761, 0x7a37a10c, 0x8e59f814, 0x89eb133c,
 604   0xeecea927, 0x35b761c9, 0xede11ce5, 0x3c7a47b1,
 605   0x599cd2df, 0x3f55f273, 0x791814ce, 0xbf73c737,
 606   0xea53f7cd, 0x5b5ffdaa, 0x14df3d6f, 0x867844db,
 607   0x81caaff3, 0x3eb968c4, 0x2c382434, 0x5fc2a340,
 608   0x72161dc3, 0x0cbce225, 0x8b283c49, 0x41ff0d95,
 609   0x7139a801, 0xde080cb3, 0x9cd8b4e4, 0x906456c1,
 610   0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
 611 };
 612
 613 __device__ __constant__ u32 td2[256] =
 614 {
 615   0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
 616   0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
 617   0xfa552030, 0x6df6ad76, 0x769188cc, 0x4c25f502,
 618   0xd7fc4fe5, 0xcbd7c52a, 0x44802635, 0xa38fb562,
 619   0x5a49deb1, 0x1b6725ba, 0x0e9845ea, 0xc0e15dfe,
 620   0x7502c32f, 0xf012814c, 0x97a38d46, 0xf9c66bd3,
 621   0x5fe7038f, 0x9c951592, 0x7aebbf6d, 0x59da9552,
 622   0x832dd4be, 0x21d35874, 0x692949e0, 0xc8448ec9,
 623   0x896a75c2, 0x7978f48e, 0x3e6b9958, 0x71dd27b9,
 624   0x4fb6bee1, 0xad17f088, 0xac66c920, 0x3ab47dce,
 625   0x4a1863df, 0x3182e51a, 0x33609751, 0x7f456253,
 626   0x77e0b164, 0xae84bb6b, 0xa01cfe81, 0x2b94f908,
 627   0x68587048, 0xfd198f45, 0x6c8794de, 0xf8b7527b,
 628   0xd323ab73, 0x02e2724b, 0x8f57e31f, 0xab2a6655,
 629   0x2807b2eb, 0xc2032fb5, 0x7b9a86c5, 0x08a5d337,
 630   0x87f23028, 0xa5b223bf, 0x6aba0203, 0x825ced16,
 631   0x1c2b8acf, 0xb492a779, 0xf2f0f307, 0xe2a14e69,
 632   0xf4cd65da, 0xbed50605, 0x621fd134, 0xfe8ac4a6,
 633   0x539d342e, 0x55a0a2f3, 0xe132058a, 0xeb75a4f6,
 634   0xec390b83, 0xefaa4060, 0x9f065e71, 0x1051bd6e,
 635   0x8af93e21, 0x063d96dd, 0x05aedd3e, 0xbd464de6,
 636   0x8db59154, 0x5d0571c4, 0xd46f0406, 0x15ff6050,
 637   0xfb241998, 0xe997d6bd, 0x43cc8940, 0x9e7767d9,
 638   0x42bdb0e8, 0x8b880789, 0x5b38e719, 0xeedb79c8,
 639   0x0a47a17c, 0x0fe97c42, 0x1ec9f884, 0x00000000,
 640   0x86830980, 0xed48322b, 0x70ac1e11, 0x724e6c5a,
 641   0xfffbfd0e, 0x38560f85, 0xd51e3dae, 0x3927362d,
 642   0xd9640a0f, 0xa621685c, 0x54d19b5b, 0x2e3a2436,
 643   0x67b10c0a, 0xe70f9357, 0x96d2b4ee, 0x919e1b9b,
 644   0xc54f80c0, 0x20a261dc, 0x4b695a77, 0x1a161c12,
 645   0xba0ae293, 0x2ae5c0a0, 0xe0433c22, 0x171d121b,
 646   0x0d0b0e09, 0xc7adf28b, 0xa8b92db6, 0xa9c8141e,
 647   0x198557f1, 0x074caf75, 0xddbbee99, 0x60fda37f,
 648   0x269ff701, 0xf5bc5c72, 0x3bc54466, 0x7e345bfb,
 649   0x29768b43, 0xc6dccb23, 0xfc68b6ed, 0xf163b8e4,
 650   0xdccad731, 0x85104263, 0x22401397, 0x112084c6,
 651   0x247d854a, 0x3df8d2bb, 0x3211aef9, 0xa16dc729,
 652   0x2f4b1d9e, 0x30f3dcb2, 0x52ec0d86, 0xe3d077c1,
 653   0x166c2bb3, 0xb999a970, 0x48fa1194, 0x642247e9,
 654   0x8cc4a8fc, 0x3f1aa0f0, 0x2cd8567d, 0x90ef2233,
 655   0x4ec78749, 0xd1c1d938, 0xa2fe8cca, 0x0b3698d4,
 656   0x81cfa6f5, 0xde28a57a, 0x8e26dab7, 0xbfa43fad,
 657   0x9de42c3a, 0x920d5078, 0xcc9b6a5f, 0x4662547e,
 658   0x13c2f68d, 0xb8e890d8, 0xf75e2e39, 0xaff582c3,
 659   0x80be9f5d, 0x937c69d0, 0x2da96fd5, 0x12b3cf25,
 660   0x993bc8ac, 0x7da71018, 0x636ee89c, 0xbb7bdb3b,
 661   0x7809cd26, 0x18f46e59, 0xb701ec9a, 0x9aa8834f,
 662   0x6e65e695, 0xe67eaaff, 0xcf0821bc, 0xe8e6ef15,
 663   0x9bd9bae7, 0x36ce4a6f, 0x09d4ea9f, 0x7cd629b0,
 664   0xb2af31a4, 0x23312a3f, 0x9430c6a5, 0x66c035a2,
 665   0xbc37744e, 0xcaa6fc82, 0xd0b0e090, 0xd81533a7,
 666   0x984af104, 0xdaf741ec, 0x500e7fcd, 0xf62f1791,
 667   0xd68d764d, 0xb04d43ef, 0x4d54ccaa, 0x04dfe496,
 668   0xb5e39ed1, 0x881b4c6a, 0x1fb8c12c, 0x517f4665,
 669   0xea049d5e, 0x355d018c, 0x7473fa87, 0x412efb0b,
 670   0x1d5ab367, 0xd25292db, 0x5633e910, 0x47136dd6,
 671   0x618c9ad7, 0x0c7a37a1, 0x148e59f8, 0x3c89eb13,
 672   0x27eecea9, 0xc935b761, 0xe5ede11c, 0xb13c7a47,
 673   0xdf599cd2, 0x733f55f2, 0xce791814, 0x37bf73c7,
 674   0xcdea53f7, 0xaa5b5ffd, 0x6f14df3d, 0xdb867844,
 675   0xf381caaf, 0xc43eb968, 0x342c3824, 0x405fc2a3,
 676   0xc372161d, 0x250cbce2, 0x498b283c, 0x9541ff0d,
 677   0x017139a8, 0xb3de080c, 0xe49cd8b4, 0xc1906456,
 678   0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
 679 };
 680
 681 __device__ __constant__ u32 td3[256] =
 682 {
 683   0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
 684   0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
 685   0x30fa5520, 0x766df6ad, 0xcc769188, 0x024c25f5,
 686   0xe5d7fc4f, 0x2acbd7c5, 0x35448026, 0x62a38fb5,
 687   0xb15a49de, 0xba1b6725, 0xea0e9845, 0xfec0e15d,
 688   0x2f7502c3, 0x4cf01281, 0x4697a38d, 0xd3f9c66b,
 689   0x8f5fe703, 0x929c9515, 0x6d7aebbf, 0x5259da95,
 690   0xbe832dd4, 0x7421d358, 0xe0692949, 0xc9c8448e,
 691   0xc2896a75, 0x8e7978f4, 0x583e6b99, 0xb971dd27,
 692   0xe14fb6be, 0x88ad17f0, 0x20ac66c9, 0xce3ab47d,
 693   0xdf4a1863, 0x1a3182e5, 0x51336097, 0x537f4562,
 694   0x6477e0b1, 0x6bae84bb, 0x81a01cfe, 0x082b94f9,
 695   0x48685870, 0x45fd198f, 0xde6c8794, 0x7bf8b752,
 696   0x73d323ab, 0x4b02e272, 0x1f8f57e3, 0x55ab2a66,
 697   0xeb2807b2, 0xb5c2032f, 0xc57b9a86, 0x3708a5d3,
 698   0x2887f230, 0xbfa5b223, 0x036aba02, 0x16825ced,
 699   0xcf1c2b8a, 0x79b492a7, 0x07f2f0f3, 0x69e2a14e,
 700   0xdaf4cd65, 0x05bed506, 0x34621fd1, 0xa6fe8ac4,
 701   0x2e539d34, 0xf355a0a2, 0x8ae13205, 0xf6eb75a4,
 702   0x83ec390b, 0x60efaa40, 0x719f065e, 0x6e1051bd,
 703   0x218af93e, 0xdd063d96, 0x3e05aedd, 0xe6bd464d,
 704   0x548db591, 0xc45d0571, 0x06d46f04, 0x5015ff60,
 705   0x98fb2419, 0xbde997d6, 0x4043cc89, 0xd99e7767,
 706   0xe842bdb0, 0x898b8807, 0x195b38e7, 0xc8eedb79,
 707   0x7c0a47a1, 0x420fe97c, 0x841ec9f8, 0x00000000,
 708   0x80868309, 0x2bed4832, 0x1170ac1e, 0x5a724e6c,
 709   0x0efffbfd, 0x8538560f, 0xaed51e3d, 0x2d392736,
 710   0x0fd9640a, 0x5ca62168, 0x5b54d19b, 0x362e3a24,
 711   0x0a67b10c, 0x57e70f93, 0xee96d2b4, 0x9b919e1b,
 712   0xc0c54f80, 0xdc20a261, 0x774b695a, 0x121a161c,
 713   0x93ba0ae2, 0xa02ae5c0, 0x22e0433c, 0x1b171d12,
 714   0x090d0b0e, 0x8bc7adf2, 0xb6a8b92d, 0x1ea9c814,
 715   0xf1198557, 0x75074caf, 0x99ddbbee, 0x7f60fda3,
 716   0x01269ff7, 0x72f5bc5c, 0x663bc544, 0xfb7e345b,
 717   0x4329768b, 0x23c6dccb, 0xedfc68b6, 0xe4f163b8,
 718   0x31dccad7, 0x63851042, 0x97224013, 0xc6112084,
 719   0x4a247d85, 0xbb3df8d2, 0xf93211ae, 0x29a16dc7,
 720   0x9e2f4b1d, 0xb230f3dc, 0x8652ec0d, 0xc1e3d077,
 721   0xb3166c2b, 0x70b999a9, 0x9448fa11, 0xe9642247,
 722   0xfc8cc4a8, 0xf03f1aa0, 0x7d2cd856, 0x3390ef22,
 723   0x494ec787, 0x38d1c1d9, 0xcaa2fe8c, 0xd40b3698,
 724   0xf581cfa6, 0x7ade28a5, 0xb78e26da, 0xadbfa43f,
 725   0x3a9de42c, 0x78920d50, 0x5fcc9b6a, 0x7e466254,
 726   0x8d13c2f6, 0xd8b8e890, 0x39f75e2e, 0xc3aff582,
 727   0x5d80be9f, 0xd0937c69, 0xd52da96f, 0x2512b3cf,
 728   0xac993bc8, 0x187da710, 0x9c636ee8, 0x3bbb7bdb,
 729   0x267809cd, 0x5918f46e, 0x9ab701ec, 0x4f9aa883,
 730   0x956e65e6, 0xffe67eaa, 0xbccf0821, 0x15e8e6ef,
 731   0xe79bd9ba, 0x6f36ce4a, 0x9f09d4ea, 0xb07cd629,
 732   0xa4b2af31, 0x3f23312a, 0xa59430c6, 0xa266c035,
 733   0x4ebc3774, 0x82caa6fc, 0x90d0b0e0, 0xa7d81533,
 734   0x04984af1, 0xecdaf741, 0xcd500e7f, 0x91f62f17,
 735   0x4dd68d76, 0xefb04d43, 0xaa4d54cc, 0x9604dfe4,
 736   0xd1b5e39e, 0x6a881b4c, 0x2c1fb8c1, 0x65517f46,
 737   0x5eea049d, 0x8c355d01, 0x877473fa, 0x0b412efb,
 738   0x671d5ab3, 0xdbd25292, 0x105633e9, 0xd647136d,
 739   0xd7618c9a, 0xa10c7a37, 0xf8148e59, 0x133c89eb,
 740   0xa927eece, 0x61c935b7, 0x1ce5ede1, 0x47b13c7a,
 741   0xd2df599c, 0xf2733f55, 0x14ce7918, 0xc737bf73,
 742   0xf7cdea53, 0xfdaa5b5f, 0x3d6f14df, 0x44db8678,
 743   0xaff381ca, 0x68c43eb9, 0x24342c38, 0xa3405fc2,
 744   0x1dc37216, 0xe2250cbc, 0x3c498b28, 0x0d9541ff,
 745   0xa8017139, 0x0cb3de08, 0xb4e49cd8, 0x56c19064,
 746   0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
 747 };
 748
 749 __device__ __constant__ u32 td4[256] =
 750 {
 751   0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
 752   0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
 753   0xbfbfbfbf, 0x40404040, 0xa3a3a3a3, 0x9e9e9e9e,
 754   0x81818181, 0xf3f3f3f3, 0xd7d7d7d7, 0xfbfbfbfb,
 755   0x7c7c7c7c, 0xe3e3e3e3, 0x39393939, 0x82828282,
 756   0x9b9b9b9b, 0x2f2f2f2f, 0xffffffff, 0x87878787,
 757   0x34343434, 0x8e8e8e8e, 0x43434343, 0x44444444,
 758   0xc4c4c4c4, 0xdededede, 0xe9e9e9e9, 0xcbcbcbcb,
 759   0x54545454, 0x7b7b7b7b, 0x94949494, 0x32323232,
 760   0xa6a6a6a6, 0xc2c2c2c2, 0x23232323, 0x3d3d3d3d,
 761   0xeeeeeeee, 0x4c4c4c4c, 0x95959595, 0x0b0b0b0b,
 762   0x42424242, 0xfafafafa, 0xc3c3c3c3, 0x4e4e4e4e,
 763   0x08080808, 0x2e2e2e2e, 0xa1a1a1a1, 0x66666666,
 764   0x28282828, 0xd9d9d9d9, 0x24242424, 0xb2b2b2b2,
 765   0x76767676, 0x5b5b5b5b, 0xa2a2a2a2, 0x49494949,
 766   0x6d6d6d6d, 0x8b8b8b8b, 0xd1d1d1d1, 0x25252525,
 767   0x72727272, 0xf8f8f8f8, 0xf6f6f6f6, 0x64646464,
 768   0x86868686, 0x68686868, 0x98989898, 0x16161616,
 769   0xd4d4d4d4, 0xa4a4a4a4, 0x5c5c5c5c, 0xcccccccc,
 770   0x5d5d5d5d, 0x65656565, 0xb6b6b6b6, 0x92929292,
 771   0x6c6c6c6c, 0x70707070, 0x48484848, 0x50505050,
 772   0xfdfdfdfd, 0xedededed, 0xb9b9b9b9, 0xdadadada,
 773   0x5e5e5e5e, 0x15151515, 0x46464646, 0x57575757,
 774   0xa7a7a7a7, 0x8d8d8d8d, 0x9d9d9d9d, 0x84848484,
 775   0x90909090, 0xd8d8d8d8, 0xabababab, 0x00000000,
 776   0x8c8c8c8c, 0xbcbcbcbc, 0xd3d3d3d3, 0x0a0a0a0a,
 777   0xf7f7f7f7, 0xe4e4e4e4, 0x58585858, 0x05050505,
 778   0xb8b8b8b8, 0xb3b3b3b3, 0x45454545, 0x06060606,
 779   0xd0d0d0d0, 0x2c2c2c2c, 0x1e1e1e1e, 0x8f8f8f8f,
 780   0xcacacaca, 0x3f3f3f3f, 0x0f0f0f0f, 0x02020202,
 781   0xc1c1c1c1, 0xafafafaf, 0xbdbdbdbd, 0x03030303,
 782   0x01010101, 0x13131313, 0x8a8a8a8a, 0x6b6b6b6b,
 783   0x3a3a3a3a, 0x91919191, 0x11111111, 0x41414141,
 784   0x4f4f4f4f, 0x67676767, 0xdcdcdcdc, 0xeaeaeaea,
 785   0x97979797, 0xf2f2f2f2, 0xcfcfcfcf, 0xcececece,
 786   0xf0f0f0f0, 0xb4b4b4b4, 0xe6e6e6e6, 0x73737373,
 787   0x96969696, 0xacacacac, 0x74747474, 0x22222222,
 788   0xe7e7e7e7, 0xadadadad, 0x35353535, 0x85858585,
 789   0xe2e2e2e2, 0xf9f9f9f9, 0x37373737, 0xe8e8e8e8,
 790   0x1c1c1c1c, 0x75757575, 0xdfdfdfdf, 0x6e6e6e6e,
 791   0x47474747, 0xf1f1f1f1, 0x1a1a1a1a, 0x71717171,
 792   0x1d1d1d1d, 0x29292929, 0xc5c5c5c5, 0x89898989,
 793   0x6f6f6f6f, 0xb7b7b7b7, 0x62626262, 0x0e0e0e0e,
 794   0xaaaaaaaa, 0x18181818, 0xbebebebe, 0x1b1b1b1b,
 795   0xfcfcfcfc, 0x56565656, 0x3e3e3e3e, 0x4b4b4b4b,
 796   0xc6c6c6c6, 0xd2d2d2d2, 0x79797979, 0x20202020,
 797   0x9a9a9a9a, 0xdbdbdbdb, 0xc0c0c0c0, 0xfefefefe,
 798   0x78787878, 0xcdcdcdcd, 0x5a5a5a5a, 0xf4f4f4f4,
 799   0x1f1f1f1f, 0xdddddddd, 0xa8a8a8a8, 0x33333333,
 800   0x88888888, 0x07070707, 0xc7c7c7c7, 0x31313131,
 801   0xb1b1b1b1, 0x12121212, 0x10101010, 0x59595959,
 802   0x27272727, 0x80808080, 0xecececec, 0x5f5f5f5f,
 803   0x60606060, 0x51515151, 0x7f7f7f7f, 0xa9a9a9a9,
 804   0x19191919, 0xb5b5b5b5, 0x4a4a4a4a, 0x0d0d0d0d,
 805   0x2d2d2d2d, 0xe5e5e5e5, 0x7a7a7a7a, 0x9f9f9f9f,
 806   0x93939393, 0xc9c9c9c9, 0x9c9c9c9c, 0xefefefef,
 807   0xa0a0a0a0, 0xe0e0e0e0, 0x3b3b3b3b, 0x4d4d4d4d,
 808   0xaeaeaeae, 0x2a2a2a2a, 0xf5f5f5f5, 0xb0b0b0b0,
 809   0xc8c8c8c8, 0xebebebeb, 0xbbbbbbbb, 0x3c3c3c3c,
 810   0x83838383, 0x53535353, 0x99999999, 0x61616161,
 811   0x17171717, 0x2b2b2b2b, 0x04040404, 0x7e7e7e7e,
 812   0xbabababa, 0x77777777, 0xd6d6d6d6, 0x26262626,
 813   0xe1e1e1e1, 0x69696969, 0x14141414, 0x63636363,
 814   0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
 815 };
 816
 817 __device__ __constant__ u32 rcon[] =
 818 {
 819   0x01000000, 0x02000000, 0x04000000, 0x08000000,
 820   0x10000000, 0x20000000, 0x40000000, 0x80000000,
 821   0x1b000000, 0x36000000,
 822 };
 823
 824 __device__ static void AES256_ExpandKey (u32 *userkey, u32 *rek, u32 s_te0[256], u32 s_te1[256], u32 s_te2[256], u32 s_te3[256], u32 s_te4[256])
 825 {
 826   rek[0] = userkey[0];
 827   rek[1] = userkey[1];
 828   rek[2] = userkey[2];
 829   rek[3] = userkey[3];
 830   rek[4] = userkey[4];
 831   rek[5] = userkey[5];
 832   rek[6] = userkey[6];
 833   rek[7] = userkey[7];
 834
 835   int i;
 836   int j;
 837
 838   i = 0;
 839   j = 0;
 840
 841   u32 run = 1;
 842
 843   while (run)
 844   {
 845     u32 temp = rek[j +  7];
 846
 847     rek[j +  8] = rek[j +  0]
 848            ^ (s_te2[(temp >> 16) & 0xff] & 0xff000000)
 849            ^ (s_te3[(temp >>  8) & 0xff] & 0x00ff0000)
 850            ^ (s_te0[(temp >>  0) & 0xff] & 0x0000ff00)
 851            ^ (s_te1[(temp >> 24) & 0xff] & 0x000000ff)
 852            ^ rcon[i];
 853
 854     rek[j +  9] = rek[j +  1] ^ rek[j +  8];
 855     rek[j + 10] = rek[j +  2] ^ rek[j +  9];
 856     rek[j + 11] = rek[j +  3] ^ rek[j + 10];
 857
 858     if (++i == 7)
 859     {
 860       run = 0;
 861       continue;
 862     }
 863
 864     temp = rek[j + 11];
 865
 866     rek[j + 12] = rek[j +  4]
 867            ^ (s_te2[(temp >> 24) & 0xff] & 0xff000000)
 868            ^ (s_te3[(temp >> 16) & 0xff] & 0x00ff0000)
 869            ^ (s_te0[(temp >>  8) & 0xff] & 0x0000ff00)
 870            ^ (s_te1[(temp >>  0) & 0xff] & 0x000000ff);
 871
 872     rek[j + 13] = rek[j +  5] ^ rek[j + 12];
 873     rek[j + 14] = rek[j +  6] ^ rek[j + 13];
 874     rek[j + 15] = rek[j +  7] ^ rek[j + 14];
 875
 876     j += 8;
 877   }
 878 }
 879
 880 __device__ static void AES256_InvertKey (u32 *rdk, u32 s_td0[256], u32 s_td1[256], u32 s_td2[256], u32 s_td3[256], u32 s_td4[256], u32 s_te0[256], u32 s_te1[256], u32 s_te2[256], u32 s_te3[256], u32 s_te4[256])
 881 {
 882   for (u32 i = 0, j = 56; i < j; i += 4, j -= 4)
 883   {
 884     u32 temp;
 885
 886     temp = rdk[i + 0]; rdk[i + 0] = rdk[j + 0]; rdk[j + 0] = temp;
 887     temp = rdk[i + 1]; rdk[i + 1] = rdk[j + 1]; rdk[j + 1] = temp;
 888     temp = rdk[i + 2]; rdk[i + 2] = rdk[j + 2]; rdk[j + 2] = temp;
 889     temp = rdk[i + 3]; rdk[i + 3] = rdk[j + 3]; rdk[j + 3] = temp;
 890   }
 891
 892   for (u32 i = 1, j = 4; i < 14; i += 1, j += 4)
 893   {
 894     rdk[j + 0] =
 895       s_td0[s_te1[(rdk[j + 0] >> 24) & 0xff] & 0xff] ^
 896       s_td1[s_te1[(rdk[j + 0] >> 16) & 0xff] & 0xff] ^
 897       s_td2[s_te1[(rdk[j + 0] >>  8) & 0xff] & 0xff] ^
 898       s_td3[s_te1[(rdk[j + 0] >>  0) & 0xff] & 0xff];
 899
 900     rdk[j + 1] =
 901       s_td0[s_te1[(rdk[j + 1] >> 24) & 0xff] & 0xff] ^
 902       s_td1[s_te1[(rdk[j + 1] >> 16) & 0xff] & 0xff] ^
 903       s_td2[s_te1[(rdk[j + 1] >>  8) & 0xff] & 0xff] ^
 904       s_td3[s_te1[(rdk[j + 1] >>  0) & 0xff] & 0xff];
 905
 906     rdk[j + 2] =
 907       s_td0[s_te1[(rdk[j + 2] >> 24) & 0xff] & 0xff] ^
 908       s_td1[s_te1[(rdk[j + 2] >> 16) & 0xff] & 0xff] ^
 909       s_td2[s_te1[(rdk[j + 2] >>  8) & 0xff] & 0xff] ^
 910       s_td3[s_te1[(rdk[j + 2] >>  0) & 0xff] & 0xff];
 911
 912     rdk[j + 3] =
 913       s_td0[s_te1[(rdk[j + 3] >> 24) & 0xff] & 0xff] ^
 914       s_td1[s_te1[(rdk[j + 3] >> 16) & 0xff] & 0xff] ^
 915       s_td2[s_te1[(rdk[j + 3] >>  8) & 0xff] & 0xff] ^
 916       s_td3[s_te1[(rdk[j + 3] >>  0) & 0xff] & 0xff];
 917   }
 918 }
 919
 920 __device__ static void AES256_decrypt (const u32 *in, u32 *out, const u32 *rdk, u32 s_td0[256], u32 s_td1[256], u32 s_td2[256], u32 s_td3[256], u32 s_td4[256])
 921 {
 922   u32 s0 = in[0] ^ rdk[0];
 923   u32 s1 = in[1] ^ rdk[1];
 924   u32 s2 = in[2] ^ rdk[2];
 925   u32 s3 = in[3] ^ rdk[3];
 926
 927   u32 t0;
 928   u32 t1;
 929   u32 t2;
 930   u32 t3;
 931
 932   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[ 4];
 933   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[ 5];
 934   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[ 6];
 935   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[ 7];
 936   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[ 8];
 937   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[ 9];
 938   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[10];
 939   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[11];
 940   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[12];
 941   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[13];
 942   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[14];
 943   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[15];
 944   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[16];
 945   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[17];
 946   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[18];
 947   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[19];
 948   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[20];
 949   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[21];
 950   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[22];
 951   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[23];
 952   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[24];
 953   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[25];
 954   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[26];
 955   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[27];
 956   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[28];
 957   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[29];
 958   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[30];
 959   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[31];
 960   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[32];
 961   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[33];
 962   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[34];
 963   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[35];
 964   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[36];
 965   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[37];
 966   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[38];
 967   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[39];
 968   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[40];
 969   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[41];
 970   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[42];
 971   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[43];
 972   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[44];
 973   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[45];
 974   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[46];
 975   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[47];
 976   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[48];
 977   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[49];
 978   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[50];
 979   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[51];
 980   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[52];
 981   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[53];
 982   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[54];
 983   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[55];
 984
 985   out[0] = (s_td4[(t0 >> 24) & 0xff] & 0xff000000)
 986          ^ (s_td4[(t3 >> 16) & 0xff] & 0x00ff0000)
 987          ^ (s_td4[(t2 >>  8) & 0xff] & 0x0000ff00)
 988          ^ (s_td4[(t1 >>  0) & 0xff] & 0x000000ff)
 989          ^ rdk[56];
 990
 991   out[1] = (s_td4[(t1 >> 24) & 0xff] & 0xff000000)
 992          ^ (s_td4[(t0 >> 16) & 0xff] & 0x00ff0000)
 993          ^ (s_td4[(t3 >>  8) & 0xff] & 0x0000ff00)
 994          ^ (s_td4[(t2 >>  0) & 0xff] & 0x000000ff)
 995          ^ rdk[57];
 996
 997   out[2] = (s_td4[(t2 >> 24) & 0xff] & 0xff000000)
 998          ^ (s_td4[(t1 >> 16) & 0xff] & 0x00ff0000)
 999          ^ (s_td4[(t0 >>  8) & 0xff] & 0x0000ff00)
1000          ^ (s_td4[(t3 >>  0) & 0xff] & 0x000000ff)
1001          ^ rdk[58];
1002
1003   out[3] = (s_td4[(t3 >> 24) & 0xff] & 0xff000000)
1004          ^ (s_td4[(t2 >> 16) & 0xff] & 0x00ff0000)
1005          ^ (s_td4[(t1 >>  8) & 0xff] & 0x0000ff00)
1006          ^ (s_td4[(t0 >>  0) & 0xff] & 0x000000ff)
1007          ^ rdk[59];
1008 }
1009
1010 extern "C" __global__ void __launch_bounds__ (256, 1) m11300_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, bitcoin_wallet_tmp_t *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 bitcoin_wallet_t *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)
1011 {
1012   /**
1013    * base
1014    */
1015
1016   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1017
1018   if (gid >= gid_max) return;
1019
1020   u32x w0[4];
1021
1022   w0[0] = pws[gid].i[ 0];
1023   w0[1] = pws[gid].i[ 1];
1024   w0[2] = pws[gid].i[ 2];
1025   w0[3] = pws[gid].i[ 3];
1026
1027   u32x w1[4];
1028
1029   w1[0] = pws[gid].i[ 4];
1030   w1[1] = pws[gid].i[ 5];
1031   w1[2] = pws[gid].i[ 6];
1032   w1[3] = pws[gid].i[ 7];
1033
1034   u32x w2[4];
1035
1036   w2[0] = pws[gid].i[ 8];
1037   w2[1] = pws[gid].i[ 9];
1038   w2[2] = 0;
1039   w2[3] = 0;
1040
1041   u32x w3[4];
1042
1043   w3[0] = 0;
1044   w3[1] = 0;
1045   w3[2] = 0;
1046   w3[3] = 0;
1047
1048   const u32 pw_len = pws[gid].pw_len;
1049
1050   /**
1051    * salt
1052    */
1053
1054   u32 salt_buf0[4];
1055
1056   salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
1057   salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
1058   salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
1059   salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
1060
1061   u32 salt_buf1[4];
1062
1063   salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
1064   salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
1065   salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
1066   salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
1067
1068   u32 salt_buf2[4];
1069
1070   salt_buf2[0] = 0;
1071   salt_buf2[1] = 0;
1072   salt_buf2[2] = 0;
1073   salt_buf2[3] = 0;
1074
1075   u32 salt_buf3[4];
1076
1077   salt_buf3[0] = 0;
1078   salt_buf3[1] = 0;
1079   salt_buf3[2] = 0;
1080   salt_buf3[3] = 0;
1081
1082   u32 salt_len = salt_bufs[salt_pos].salt_len;
1083
1084   switch_buffer_by_offset (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
1085
1086   w0[0] |= salt_buf0[0];
1087   w0[1] |= salt_buf0[1];
1088   w0[2] |= salt_buf0[2];
1089   w0[3] |= salt_buf0[3];
1090
1091   w1[0] |= salt_buf1[0];
1092   w1[1] |= salt_buf1[1];
1093   w1[2] |= salt_buf1[2];
1094   w1[3] |= salt_buf1[3];
1095
1096   w2[0] |= salt_buf2[0];
1097   w2[1] |= salt_buf2[1];
1098   w2[2] |= salt_buf2[2];
1099   w2[3] |= salt_buf2[3];
1100
1101   w3[0] |= salt_buf3[0];
1102   w3[1] |= salt_buf3[1];
1103   w3[2] |= salt_buf3[2];
1104   w3[3] |= salt_buf3[3];
1105
1106   const u32 block_len = pw_len + salt_len;
1107
1108   /**
1109    * init
1110    */
1111
1112   w0[0] = swap_workaround (w0[0]);
1113   w0[1] = swap_workaround (w0[1]);
1114   w0[2] = swap_workaround (w0[2]);
1115   w0[3] = swap_workaround (w0[3]);
1116   w1[0] = swap_workaround (w1[0]);
1117   w1[1] = swap_workaround (w1[1]);
1118   w1[2] = swap_workaround (w1[2]);
1119   w1[3] = swap_workaround (w1[3]);
1120   w2[0] = swap_workaround (w2[0]);
1121   w2[1] = swap_workaround (w2[1]);
1122   w2[2] = swap_workaround (w2[2]);
1123   w2[3] = swap_workaround (w2[3]);
1124   w3[0] = swap_workaround (w3[0]);
1125   w3[1] = swap_workaround (w3[1]);
1126   w3[2] = swap_workaround (w3[2]);
1127   w3[3] = swap_workaround (w3[3]);
1128
1129   u64x w[16];
1130
1131   w[ 0] = hl32_to_64 (w0[0], w0[1]);
1132   w[ 1] = hl32_to_64 (w0[2], w0[3]);
1133   w[ 2] = hl32_to_64 (w1[0], w1[1]);
1134   w[ 3] = hl32_to_64 (w1[2], w1[3]);
1135   w[ 4] = hl32_to_64 (w2[0], w2[1]);
1136   w[ 5] = hl32_to_64 (w2[2], w2[3]);
1137   w[ 6] = hl32_to_64 (w3[0], w3[1]);
1138   w[ 7] = hl32_to_64 (w3[2], w3[3]);
1139   w[ 8] = 0;
1140   w[ 9] = 0;
1141   w[10] = 0;
1142   w[11] = 0;
1143   w[12] = 0;
1144   w[13] = 0;
1145   w[14] = 0;
1146   w[15] = block_len * 8;
1147
1148   u64x dgst[8];
1149
1150   dgst[0] = SHA512M_A;
1151   dgst[1] = SHA512M_B;
1152   dgst[2] = SHA512M_C;
1153   dgst[3] = SHA512M_D;
1154   dgst[4] = SHA512M_E;
1155   dgst[5] = SHA512M_F;
1156   dgst[6] = SHA512M_G;
1157   dgst[7] = SHA512M_H;
1158
1159   sha512_transform (w, dgst);
1160
1161   tmps[gid].dgst[0] = dgst[0];
1162   tmps[gid].dgst[1] = dgst[1];
1163   tmps[gid].dgst[2] = dgst[2];
1164   tmps[gid].dgst[3] = dgst[3];
1165   tmps[gid].dgst[4] = dgst[4];
1166   tmps[gid].dgst[5] = dgst[5];
1167   tmps[gid].dgst[6] = dgst[6];
1168   tmps[gid].dgst[7] = dgst[7];
1169 }
1170
1171 extern "C" __global__ void __launch_bounds__ (256, 1) m11300_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, bitcoin_wallet_tmp_t *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 bitcoin_wallet_t *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)
1172 {
1173   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1174
1175
1176   if (gid >= gid_max) return;
1177
1178   u64x dgst[8];
1179
1180   dgst[0] = tmps[gid].dgst[0];
1181   dgst[1] = tmps[gid].dgst[1];
1182   dgst[2] = tmps[gid].dgst[2];
1183   dgst[3] = tmps[gid].dgst[3];
1184   dgst[4] = tmps[gid].dgst[4];
1185   dgst[5] = tmps[gid].dgst[5];
1186   dgst[6] = tmps[gid].dgst[6];
1187   dgst[7] = tmps[gid].dgst[7];
1188
1189   u64x w[16];
1190
1191   w[ 0] = dgst[0];
1192   w[ 1] = dgst[1];
1193   w[ 2] = dgst[2];
1194   w[ 3] = dgst[3];
1195   w[ 4] = dgst[4];
1196   w[ 5] = dgst[5];
1197   w[ 6] = dgst[6];
1198   w[ 7] = dgst[7];
1199   w[ 8] = 0x8000000000000000;
1200   w[ 9] = 0;
1201   w[10] = 0;
1202   w[11] = 0;
1203   w[12] = 0;
1204   w[13] = 0;
1205   w[14] = 0;
1206   w[15] = 64 * 8;
1207
1208   for (u32 i = 0; i < loop_cnt; i++)
1209   {
1210     w[0] = dgst[0];
1211     w[1] = dgst[1];
1212     w[2] = dgst[2];
1213     w[3] = dgst[3];
1214     w[4] = dgst[4];
1215     w[5] = dgst[5];
1216     w[6] = dgst[6];
1217     w[7] = dgst[7];
1218
1219     dgst[0] = SHA512M_A;
1220     dgst[1] = SHA512M_B;
1221     dgst[2] = SHA512M_C;
1222     dgst[3] = SHA512M_D;
1223     dgst[4] = SHA512M_E;
1224     dgst[5] = SHA512M_F;
1225     dgst[6] = SHA512M_G;
1226     dgst[7] = SHA512M_H;
1227
1228     sha512_transform (w, dgst);
1229   }
1230
1231   tmps[gid].dgst[0] = dgst[0];
1232   tmps[gid].dgst[1] = dgst[1];
1233   tmps[gid].dgst[2] = dgst[2];
1234   tmps[gid].dgst[3] = dgst[3];
1235   tmps[gid].dgst[4] = dgst[4];
1236   tmps[gid].dgst[5] = dgst[5];
1237   tmps[gid].dgst[6] = dgst[6];
1238   tmps[gid].dgst[7] = dgst[7];
1239 }
1240
1241 extern "C" __global__ void __launch_bounds__ (256, 1) m11300_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, bitcoin_wallet_tmp_t *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 bitcoin_wallet_t *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)
1242 {
1243   /**
1244    * base
1245    */
1246
1247   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1248   const u32 lid = threadIdx.x;
1249
1250   /**
1251    * aes shared
1252    */
1253
1254   __shared__ u32 s_td0[256];
1255   __shared__ u32 s_td1[256];
1256   __shared__ u32 s_td2[256];
1257   __shared__ u32 s_td3[256];
1258   __shared__ u32 s_td4[256];
1259
1260   __shared__ u32 s_te0[256];
1261   __shared__ u32 s_te1[256];
1262   __shared__ u32 s_te2[256];
1263   __shared__ u32 s_te3[256];
1264   __shared__ u32 s_te4[256];
1265
1266   s_td0[lid] = td0[lid];
1267   s_td1[lid] = td1[lid];
1268   s_td2[lid] = td2[lid];
1269   s_td3[lid] = td3[lid];
1270   s_td4[lid] = td4[lid];
1271
1272   s_te0[lid] = te0[lid];
1273   s_te1[lid] = te1[lid];
1274   s_te2[lid] = te2[lid];
1275   s_te3[lid] = te3[lid];
1276   s_te4[lid] = te4[lid];
1277
1278   __syncthreads ();
1279
1280   if (gid >= gid_max) return;
1281
1282   /**
1283    * real code
1284    */
1285
1286   u64 dgst[8];
1287
1288   dgst[0] = tmps[gid].dgst[0];
1289   dgst[1] = tmps[gid].dgst[1];
1290   dgst[2] = tmps[gid].dgst[2];
1291   dgst[3] = tmps[gid].dgst[3];
1292   dgst[4] = tmps[gid].dgst[4];
1293   dgst[5] = tmps[gid].dgst[5];
1294   dgst[6] = tmps[gid].dgst[6];
1295   dgst[7] = tmps[gid].dgst[7];
1296
1297   u32x key[8];
1298
1299   key[0] = h32_from_64 (dgst[0]);
1300   key[1] = l32_from_64 (dgst[0]);
1301   key[2] = h32_from_64 (dgst[1]);
1302   key[3] = l32_from_64 (dgst[1]);
1303   key[4] = h32_from_64 (dgst[2]);
1304   key[5] = l32_from_64 (dgst[2]);
1305   key[6] = h32_from_64 (dgst[3]);
1306   key[7] = l32_from_64 (dgst[3]);
1307
1308   u32x iv[4];
1309
1310   iv[0] = h32_from_64 (dgst[4]);
1311   iv[1] = l32_from_64 (dgst[4]);
1312   iv[2] = h32_from_64 (dgst[5]);
1313   iv[3] = l32_from_64 (dgst[5]);
1314
1315   #define KEYLEN 60
1316
1317   u32 rk[KEYLEN];
1318
1319   AES256_ExpandKey (key, rk, s_te0, s_te1, s_te2, s_te3, s_te4);
1320
1321   AES256_InvertKey (rk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
1322
1323   u32x out[4];
1324
1325   for (u32 i = 0; i < esalt_bufs[salt_pos].cry_master_len; i += 16)
1326   {
1327     u32x data[4];
1328
1329     data[0] = swap_workaround (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 0]);
1330     data[1] = swap_workaround (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 1]);
1331     data[2] = swap_workaround (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 2]);
1332     data[3] = swap_workaround (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 3]);
1333
1334     AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1335
1336     out[0] ^= iv[0];
1337     out[1] ^= iv[1];
1338     out[2] ^= iv[2];
1339     out[3] ^= iv[3];
1340
1341     iv[0] = data[0];
1342     iv[1] = data[1];
1343     iv[2] = data[2];
1344     iv[3] = data[3];
1345   }
1346
1347   if ((out[0] == 0x10101010)
1348    && (out[1] == 0x10101010)
1349    && (out[2] == 0x10101010)
1350    && (out[3] == 0x10101010))
1351   {
1352     mark_hash_s0 (plains_buf, hashes_shown, digests_offset + 0, gid, 0);
1353
1354     d_return_buf[lid] = 1;
1355   }
1356 }