nv/m11600.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _SEVEN_ZIP_
   7
   8 #include "include/constants.h"
   9 #include "include/kernel_vendor.h"
  10
  11 #ifdef  VLIW1
  12 #define VECT_SIZE1
  13 #endif
  14
  15 #ifdef  VLIW2
  16 #define VECT_SIZE1
  17 #endif
  18
  19 #define DGST_R0 0
  20 #define DGST_R1 1
  21 #define DGST_R2 2
  22 #define DGST_R3 3
  23
  24 #include "include/kernel_functions.c"
  25 #include "types_nv.c"
  26 #include "common_nv.c"
  27
  28 #ifdef  VECT_SIZE1
  29 #define VECT_COMPARE_M "check_multi_vect1_comp4.c"
  30 #endif
  31
  32 #ifdef  VECT_SIZE2
  33 #define VECT_COMPARE_M "check_multi_vect2_comp4.c"
  34 #endif
  35
  36 __device__ __constant__ u32 te0[256] =
  37 {
  38   0xc66363a5, 0xf87c7c84, 0xee777799, 0xf67b7b8d,
  39   0xfff2f20d, 0xd66b6bbd, 0xde6f6fb1, 0x91c5c554,
  40   0x60303050, 0x02010103, 0xce6767a9, 0x562b2b7d,
  41   0xe7fefe19, 0xb5d7d762, 0x4dababe6, 0xec76769a,
  42   0x8fcaca45, 0x1f82829d, 0x89c9c940, 0xfa7d7d87,
  43   0xeffafa15, 0xb25959eb, 0x8e4747c9, 0xfbf0f00b,
  44   0x41adadec, 0xb3d4d467, 0x5fa2a2fd, 0x45afafea,
  45   0x239c9cbf, 0x53a4a4f7, 0xe4727296, 0x9bc0c05b,
  46   0x75b7b7c2, 0xe1fdfd1c, 0x3d9393ae, 0x4c26266a,
  47   0x6c36365a, 0x7e3f3f41, 0xf5f7f702, 0x83cccc4f,
  48   0x6834345c, 0x51a5a5f4, 0xd1e5e534, 0xf9f1f108,
  49   0xe2717193, 0xabd8d873, 0x62313153, 0x2a15153f,
  50   0x0804040c, 0x95c7c752, 0x46232365, 0x9dc3c35e,
  51   0x30181828, 0x379696a1, 0x0a05050f, 0x2f9a9ab5,
  52   0x0e070709, 0x24121236, 0x1b80809b, 0xdfe2e23d,
  53   0xcdebeb26, 0x4e272769, 0x7fb2b2cd, 0xea75759f,
  54   0x1209091b, 0x1d83839e, 0x582c2c74, 0x341a1a2e,
  55   0x361b1b2d, 0xdc6e6eb2, 0xb45a5aee, 0x5ba0a0fb,
  56   0xa45252f6, 0x763b3b4d, 0xb7d6d661, 0x7db3b3ce,
  57   0x5229297b, 0xdde3e33e, 0x5e2f2f71, 0x13848497,
  58   0xa65353f5, 0xb9d1d168, 0x00000000, 0xc1eded2c,
  59   0x40202060, 0xe3fcfc1f, 0x79b1b1c8, 0xb65b5bed,
  60   0xd46a6abe, 0x8dcbcb46, 0x67bebed9, 0x7239394b,
  61   0x944a4ade, 0x984c4cd4, 0xb05858e8, 0x85cfcf4a,
  62   0xbbd0d06b, 0xc5efef2a, 0x4faaaae5, 0xedfbfb16,
  63   0x864343c5, 0x9a4d4dd7, 0x66333355, 0x11858594,
  64   0x8a4545cf, 0xe9f9f910, 0x04020206, 0xfe7f7f81,
  65   0xa05050f0, 0x783c3c44, 0x259f9fba, 0x4ba8a8e3,
  66   0xa25151f3, 0x5da3a3fe, 0x804040c0, 0x058f8f8a,
  67   0x3f9292ad, 0x219d9dbc, 0x70383848, 0xf1f5f504,
  68   0x63bcbcdf, 0x77b6b6c1, 0xafdada75, 0x42212163,
  69   0x20101030, 0xe5ffff1a, 0xfdf3f30e, 0xbfd2d26d,
  70   0x81cdcd4c, 0x180c0c14, 0x26131335, 0xc3ecec2f,
  71   0xbe5f5fe1, 0x359797a2, 0x884444cc, 0x2e171739,
  72   0x93c4c457, 0x55a7a7f2, 0xfc7e7e82, 0x7a3d3d47,
  73   0xc86464ac, 0xba5d5de7, 0x3219192b, 0xe6737395,
  74   0xc06060a0, 0x19818198, 0x9e4f4fd1, 0xa3dcdc7f,
  75   0x44222266, 0x542a2a7e, 0x3b9090ab, 0x0b888883,
  76   0x8c4646ca, 0xc7eeee29, 0x6bb8b8d3, 0x2814143c,
  77   0xa7dede79, 0xbc5e5ee2, 0x160b0b1d, 0xaddbdb76,
  78   0xdbe0e03b, 0x64323256, 0x743a3a4e, 0x140a0a1e,
  79   0x924949db, 0x0c06060a, 0x4824246c, 0xb85c5ce4,
  80   0x9fc2c25d, 0xbdd3d36e, 0x43acacef, 0xc46262a6,
  81   0x399191a8, 0x319595a4, 0xd3e4e437, 0xf279798b,
  82   0xd5e7e732, 0x8bc8c843, 0x6e373759, 0xda6d6db7,
  83   0x018d8d8c, 0xb1d5d564, 0x9c4e4ed2, 0x49a9a9e0,
  84   0xd86c6cb4, 0xac5656fa, 0xf3f4f407, 0xcfeaea25,
  85   0xca6565af, 0xf47a7a8e, 0x47aeaee9, 0x10080818,
  86   0x6fbabad5, 0xf0787888, 0x4a25256f, 0x5c2e2e72,
  87   0x381c1c24, 0x57a6a6f1, 0x73b4b4c7, 0x97c6c651,
  88   0xcbe8e823, 0xa1dddd7c, 0xe874749c, 0x3e1f1f21,
  89   0x964b4bdd, 0x61bdbddc, 0x0d8b8b86, 0x0f8a8a85,
  90   0xe0707090, 0x7c3e3e42, 0x71b5b5c4, 0xcc6666aa,
  91   0x904848d8, 0x06030305, 0xf7f6f601, 0x1c0e0e12,
  92   0xc26161a3, 0x6a35355f, 0xae5757f9, 0x69b9b9d0,
  93   0x17868691, 0x99c1c158, 0x3a1d1d27, 0x279e9eb9,
  94   0xd9e1e138, 0xebf8f813, 0x2b9898b3, 0x22111133,
  95   0xd26969bb, 0xa9d9d970, 0x078e8e89, 0x339494a7,
  96   0x2d9b9bb6, 0x3c1e1e22, 0x15878792, 0xc9e9e920,
  97   0x87cece49, 0xaa5555ff, 0x50282878, 0xa5dfdf7a,
  98   0x038c8c8f, 0x59a1a1f8, 0x09898980, 0x1a0d0d17,
  99   0x65bfbfda, 0xd7e6e631, 0x844242c6, 0xd06868b8,
 100   0x824141c3, 0x299999b0, 0x5a2d2d77, 0x1e0f0f11,
 101   0x7bb0b0cb, 0xa85454fc, 0x6dbbbbd6, 0x2c16163a,
 102 };
 103
 104 __device__ __constant__ u32 te1[256] =
 105 {
 106   0xa5c66363, 0x84f87c7c, 0x99ee7777, 0x8df67b7b,
 107   0x0dfff2f2, 0xbdd66b6b, 0xb1de6f6f, 0x5491c5c5,
 108   0x50603030, 0x03020101, 0xa9ce6767, 0x7d562b2b,
 109   0x19e7fefe, 0x62b5d7d7, 0xe64dabab, 0x9aec7676,
 110   0x458fcaca, 0x9d1f8282, 0x4089c9c9, 0x87fa7d7d,
 111   0x15effafa, 0xebb25959, 0xc98e4747, 0x0bfbf0f0,
 112   0xec41adad, 0x67b3d4d4, 0xfd5fa2a2, 0xea45afaf,
 113   0xbf239c9c, 0xf753a4a4, 0x96e47272, 0x5b9bc0c0,
 114   0xc275b7b7, 0x1ce1fdfd, 0xae3d9393, 0x6a4c2626,
 115   0x5a6c3636, 0x417e3f3f, 0x02f5f7f7, 0x4f83cccc,
 116   0x5c683434, 0xf451a5a5, 0x34d1e5e5, 0x08f9f1f1,
 117   0x93e27171, 0x73abd8d8, 0x53623131, 0x3f2a1515,
 118   0x0c080404, 0x5295c7c7, 0x65462323, 0x5e9dc3c3,
 119   0x28301818, 0xa1379696, 0x0f0a0505, 0xb52f9a9a,
 120   0x090e0707, 0x36241212, 0x9b1b8080, 0x3ddfe2e2,
 121   0x26cdebeb, 0x694e2727, 0xcd7fb2b2, 0x9fea7575,
 122   0x1b120909, 0x9e1d8383, 0x74582c2c, 0x2e341a1a,
 123   0x2d361b1b, 0xb2dc6e6e, 0xeeb45a5a, 0xfb5ba0a0,
 124   0xf6a45252, 0x4d763b3b, 0x61b7d6d6, 0xce7db3b3,
 125   0x7b522929, 0x3edde3e3, 0x715e2f2f, 0x97138484,
 126   0xf5a65353, 0x68b9d1d1, 0x00000000, 0x2cc1eded,
 127   0x60402020, 0x1fe3fcfc, 0xc879b1b1, 0xedb65b5b,
 128   0xbed46a6a, 0x468dcbcb, 0xd967bebe, 0x4b723939,
 129   0xde944a4a, 0xd4984c4c, 0xe8b05858, 0x4a85cfcf,
 130   0x6bbbd0d0, 0x2ac5efef, 0xe54faaaa, 0x16edfbfb,
 131   0xc5864343, 0xd79a4d4d, 0x55663333, 0x94118585,
 132   0xcf8a4545, 0x10e9f9f9, 0x06040202, 0x81fe7f7f,
 133   0xf0a05050, 0x44783c3c, 0xba259f9f, 0xe34ba8a8,
 134   0xf3a25151, 0xfe5da3a3, 0xc0804040, 0x8a058f8f,
 135   0xad3f9292, 0xbc219d9d, 0x48703838, 0x04f1f5f5,
 136   0xdf63bcbc, 0xc177b6b6, 0x75afdada, 0x63422121,
 137   0x30201010, 0x1ae5ffff, 0x0efdf3f3, 0x6dbfd2d2,
 138   0x4c81cdcd, 0x14180c0c, 0x35261313, 0x2fc3ecec,
 139   0xe1be5f5f, 0xa2359797, 0xcc884444, 0x392e1717,
 140   0x5793c4c4, 0xf255a7a7, 0x82fc7e7e, 0x477a3d3d,
 141   0xacc86464, 0xe7ba5d5d, 0x2b321919, 0x95e67373,
 142   0xa0c06060, 0x98198181, 0xd19e4f4f, 0x7fa3dcdc,
 143   0x66442222, 0x7e542a2a, 0xab3b9090, 0x830b8888,
 144   0xca8c4646, 0x29c7eeee, 0xd36bb8b8, 0x3c281414,
 145   0x79a7dede, 0xe2bc5e5e, 0x1d160b0b, 0x76addbdb,
 146   0x3bdbe0e0, 0x56643232, 0x4e743a3a, 0x1e140a0a,
 147   0xdb924949, 0x0a0c0606, 0x6c482424, 0xe4b85c5c,
 148   0x5d9fc2c2, 0x6ebdd3d3, 0xef43acac, 0xa6c46262,
 149   0xa8399191, 0xa4319595, 0x37d3e4e4, 0x8bf27979,
 150   0x32d5e7e7, 0x438bc8c8, 0x596e3737, 0xb7da6d6d,
 151   0x8c018d8d, 0x64b1d5d5, 0xd29c4e4e, 0xe049a9a9,
 152   0xb4d86c6c, 0xfaac5656, 0x07f3f4f4, 0x25cfeaea,
 153   0xafca6565, 0x8ef47a7a, 0xe947aeae, 0x18100808,
 154   0xd56fbaba, 0x88f07878, 0x6f4a2525, 0x725c2e2e,
 155   0x24381c1c, 0xf157a6a6, 0xc773b4b4, 0x5197c6c6,
 156   0x23cbe8e8, 0x7ca1dddd, 0x9ce87474, 0x213e1f1f,
 157   0xdd964b4b, 0xdc61bdbd, 0x860d8b8b, 0x850f8a8a,
 158   0x90e07070, 0x427c3e3e, 0xc471b5b5, 0xaacc6666,
 159   0xd8904848, 0x05060303, 0x01f7f6f6, 0x121c0e0e,
 160   0xa3c26161, 0x5f6a3535, 0xf9ae5757, 0xd069b9b9,
 161   0x91178686, 0x5899c1c1, 0x273a1d1d, 0xb9279e9e,
 162   0x38d9e1e1, 0x13ebf8f8, 0xb32b9898, 0x33221111,
 163   0xbbd26969, 0x70a9d9d9, 0x89078e8e, 0xa7339494,
 164   0xb62d9b9b, 0x223c1e1e, 0x92158787, 0x20c9e9e9,
 165   0x4987cece, 0xffaa5555, 0x78502828, 0x7aa5dfdf,
 166   0x8f038c8c, 0xf859a1a1, 0x80098989, 0x171a0d0d,
 167   0xda65bfbf, 0x31d7e6e6, 0xc6844242, 0xb8d06868,
 168   0xc3824141, 0xb0299999, 0x775a2d2d, 0x111e0f0f,
 169   0xcb7bb0b0, 0xfca85454, 0xd66dbbbb, 0x3a2c1616,
 170 };
 171
 172 __device__ __constant__ u32 te2[256] =
 173 {
 174   0x63a5c663, 0x7c84f87c, 0x7799ee77, 0x7b8df67b,
 175   0xf20dfff2, 0x6bbdd66b, 0x6fb1de6f, 0xc55491c5,
 176   0x30506030, 0x01030201, 0x67a9ce67, 0x2b7d562b,
 177   0xfe19e7fe, 0xd762b5d7, 0xabe64dab, 0x769aec76,
 178   0xca458fca, 0x829d1f82, 0xc94089c9, 0x7d87fa7d,
 179   0xfa15effa, 0x59ebb259, 0x47c98e47, 0xf00bfbf0,
 180   0xadec41ad, 0xd467b3d4, 0xa2fd5fa2, 0xafea45af,
 181   0x9cbf239c, 0xa4f753a4, 0x7296e472, 0xc05b9bc0,
 182   0xb7c275b7, 0xfd1ce1fd, 0x93ae3d93, 0x266a4c26,
 183   0x365a6c36, 0x3f417e3f, 0xf702f5f7, 0xcc4f83cc,
 184   0x345c6834, 0xa5f451a5, 0xe534d1e5, 0xf108f9f1,
 185   0x7193e271, 0xd873abd8, 0x31536231, 0x153f2a15,
 186   0x040c0804, 0xc75295c7, 0x23654623, 0xc35e9dc3,
 187   0x18283018, 0x96a13796, 0x050f0a05, 0x9ab52f9a,
 188   0x07090e07, 0x12362412, 0x809b1b80, 0xe23ddfe2,
 189   0xeb26cdeb, 0x27694e27, 0xb2cd7fb2, 0x759fea75,
 190   0x091b1209, 0x839e1d83, 0x2c74582c, 0x1a2e341a,
 191   0x1b2d361b, 0x6eb2dc6e, 0x5aeeb45a, 0xa0fb5ba0,
 192   0x52f6a452, 0x3b4d763b, 0xd661b7d6, 0xb3ce7db3,
 193   0x297b5229, 0xe33edde3, 0x2f715e2f, 0x84971384,
 194   0x53f5a653, 0xd168b9d1, 0x00000000, 0xed2cc1ed,
 195   0x20604020, 0xfc1fe3fc, 0xb1c879b1, 0x5bedb65b,
 196   0x6abed46a, 0xcb468dcb, 0xbed967be, 0x394b7239,
 197   0x4ade944a, 0x4cd4984c, 0x58e8b058, 0xcf4a85cf,
 198   0xd06bbbd0, 0xef2ac5ef, 0xaae54faa, 0xfb16edfb,
 199   0x43c58643, 0x4dd79a4d, 0x33556633, 0x85941185,
 200   0x45cf8a45, 0xf910e9f9, 0x02060402, 0x7f81fe7f,
 201   0x50f0a050, 0x3c44783c, 0x9fba259f, 0xa8e34ba8,
 202   0x51f3a251, 0xa3fe5da3, 0x40c08040, 0x8f8a058f,
 203   0x92ad3f92, 0x9dbc219d, 0x38487038, 0xf504f1f5,
 204   0xbcdf63bc, 0xb6c177b6, 0xda75afda, 0x21634221,
 205   0x10302010, 0xff1ae5ff, 0xf30efdf3, 0xd26dbfd2,
 206   0xcd4c81cd, 0x0c14180c, 0x13352613, 0xec2fc3ec,
 207   0x5fe1be5f, 0x97a23597, 0x44cc8844, 0x17392e17,
 208   0xc45793c4, 0xa7f255a7, 0x7e82fc7e, 0x3d477a3d,
 209   0x64acc864, 0x5de7ba5d, 0x192b3219, 0x7395e673,
 210   0x60a0c060, 0x81981981, 0x4fd19e4f, 0xdc7fa3dc,
 211   0x22664422, 0x2a7e542a, 0x90ab3b90, 0x88830b88,
 212   0x46ca8c46, 0xee29c7ee, 0xb8d36bb8, 0x143c2814,
 213   0xde79a7de, 0x5ee2bc5e, 0x0b1d160b, 0xdb76addb,
 214   0xe03bdbe0, 0x32566432, 0x3a4e743a, 0x0a1e140a,
 215   0x49db9249, 0x060a0c06, 0x246c4824, 0x5ce4b85c,
 216   0xc25d9fc2, 0xd36ebdd3, 0xacef43ac, 0x62a6c462,
 217   0x91a83991, 0x95a43195, 0xe437d3e4, 0x798bf279,
 218   0xe732d5e7, 0xc8438bc8, 0x37596e37, 0x6db7da6d,
 219   0x8d8c018d, 0xd564b1d5, 0x4ed29c4e, 0xa9e049a9,
 220   0x6cb4d86c, 0x56faac56, 0xf407f3f4, 0xea25cfea,
 221   0x65afca65, 0x7a8ef47a, 0xaee947ae, 0x08181008,
 222   0xbad56fba, 0x7888f078, 0x256f4a25, 0x2e725c2e,
 223   0x1c24381c, 0xa6f157a6, 0xb4c773b4, 0xc65197c6,
 224   0xe823cbe8, 0xdd7ca1dd, 0x749ce874, 0x1f213e1f,
 225   0x4bdd964b, 0xbddc61bd, 0x8b860d8b, 0x8a850f8a,
 226   0x7090e070, 0x3e427c3e, 0xb5c471b5, 0x66aacc66,
 227   0x48d89048, 0x03050603, 0xf601f7f6, 0x0e121c0e,
 228   0x61a3c261, 0x355f6a35, 0x57f9ae57, 0xb9d069b9,
 229   0x86911786, 0xc15899c1, 0x1d273a1d, 0x9eb9279e,
 230   0xe138d9e1, 0xf813ebf8, 0x98b32b98, 0x11332211,
 231   0x69bbd269, 0xd970a9d9, 0x8e89078e, 0x94a73394,
 232   0x9bb62d9b, 0x1e223c1e, 0x87921587, 0xe920c9e9,
 233   0xce4987ce, 0x55ffaa55, 0x28785028, 0xdf7aa5df,
 234   0x8c8f038c, 0xa1f859a1, 0x89800989, 0x0d171a0d,
 235   0xbfda65bf, 0xe631d7e6, 0x42c68442, 0x68b8d068,
 236   0x41c38241, 0x99b02999, 0x2d775a2d, 0x0f111e0f,
 237   0xb0cb7bb0, 0x54fca854, 0xbbd66dbb, 0x163a2c16,
 238 };
 239
 240 __device__ __constant__ u32 te3[256] =
 241 {
 242   0x6363a5c6, 0x7c7c84f8, 0x777799ee, 0x7b7b8df6,
 243   0xf2f20dff, 0x6b6bbdd6, 0x6f6fb1de, 0xc5c55491,
 244   0x30305060, 0x01010302, 0x6767a9ce, 0x2b2b7d56,
 245   0xfefe19e7, 0xd7d762b5, 0xababe64d, 0x76769aec,
 246   0xcaca458f, 0x82829d1f, 0xc9c94089, 0x7d7d87fa,
 247   0xfafa15ef, 0x5959ebb2, 0x4747c98e, 0xf0f00bfb,
 248   0xadadec41, 0xd4d467b3, 0xa2a2fd5f, 0xafafea45,
 249   0x9c9cbf23, 0xa4a4f753, 0x727296e4, 0xc0c05b9b,
 250   0xb7b7c275, 0xfdfd1ce1, 0x9393ae3d, 0x26266a4c,
 251   0x36365a6c, 0x3f3f417e, 0xf7f702f5, 0xcccc4f83,
 252   0x34345c68, 0xa5a5f451, 0xe5e534d1, 0xf1f108f9,
 253   0x717193e2, 0xd8d873ab, 0x31315362, 0x15153f2a,
 254   0x04040c08, 0xc7c75295, 0x23236546, 0xc3c35e9d,
 255   0x18182830, 0x9696a137, 0x05050f0a, 0x9a9ab52f,
 256   0x0707090e, 0x12123624, 0x80809b1b, 0xe2e23ddf,
 257   0xebeb26cd, 0x2727694e, 0xb2b2cd7f, 0x75759fea,
 258   0x09091b12, 0x83839e1d, 0x2c2c7458, 0x1a1a2e34,
 259   0x1b1b2d36, 0x6e6eb2dc, 0x5a5aeeb4, 0xa0a0fb5b,
 260   0x5252f6a4, 0x3b3b4d76, 0xd6d661b7, 0xb3b3ce7d,
 261   0x29297b52, 0xe3e33edd, 0x2f2f715e, 0x84849713,
 262   0x5353f5a6, 0xd1d168b9, 0x00000000, 0xeded2cc1,
 263   0x20206040, 0xfcfc1fe3, 0xb1b1c879, 0x5b5bedb6,
 264   0x6a6abed4, 0xcbcb468d, 0xbebed967, 0x39394b72,
 265   0x4a4ade94, 0x4c4cd498, 0x5858e8b0, 0xcfcf4a85,
 266   0xd0d06bbb, 0xefef2ac5, 0xaaaae54f, 0xfbfb16ed,
 267   0x4343c586, 0x4d4dd79a, 0x33335566, 0x85859411,
 268   0x4545cf8a, 0xf9f910e9, 0x02020604, 0x7f7f81fe,
 269   0x5050f0a0, 0x3c3c4478, 0x9f9fba25, 0xa8a8e34b,
 270   0x5151f3a2, 0xa3a3fe5d, 0x4040c080, 0x8f8f8a05,
 271   0x9292ad3f, 0x9d9dbc21, 0x38384870, 0xf5f504f1,
 272   0xbcbcdf63, 0xb6b6c177, 0xdada75af, 0x21216342,
 273   0x10103020, 0xffff1ae5, 0xf3f30efd, 0xd2d26dbf,
 274   0xcdcd4c81, 0x0c0c1418, 0x13133526, 0xecec2fc3,
 275   0x5f5fe1be, 0x9797a235, 0x4444cc88, 0x1717392e,
 276   0xc4c45793, 0xa7a7f255, 0x7e7e82fc, 0x3d3d477a,
 277   0x6464acc8, 0x5d5de7ba, 0x19192b32, 0x737395e6,
 278   0x6060a0c0, 0x81819819, 0x4f4fd19e, 0xdcdc7fa3,
 279   0x22226644, 0x2a2a7e54, 0x9090ab3b, 0x8888830b,
 280   0x4646ca8c, 0xeeee29c7, 0xb8b8d36b, 0x14143c28,
 281   0xdede79a7, 0x5e5ee2bc, 0x0b0b1d16, 0xdbdb76ad,
 282   0xe0e03bdb, 0x32325664, 0x3a3a4e74, 0x0a0a1e14,
 283   0x4949db92, 0x06060a0c, 0x24246c48, 0x5c5ce4b8,
 284   0xc2c25d9f, 0xd3d36ebd, 0xacacef43, 0x6262a6c4,
 285   0x9191a839, 0x9595a431, 0xe4e437d3, 0x79798bf2,
 286   0xe7e732d5, 0xc8c8438b, 0x3737596e, 0x6d6db7da,
 287   0x8d8d8c01, 0xd5d564b1, 0x4e4ed29c, 0xa9a9e049,
 288   0x6c6cb4d8, 0x5656faac, 0xf4f407f3, 0xeaea25cf,
 289   0x6565afca, 0x7a7a8ef4, 0xaeaee947, 0x08081810,
 290   0xbabad56f, 0x787888f0, 0x25256f4a, 0x2e2e725c,
 291   0x1c1c2438, 0xa6a6f157, 0xb4b4c773, 0xc6c65197,
 292   0xe8e823cb, 0xdddd7ca1, 0x74749ce8, 0x1f1f213e,
 293   0x4b4bdd96, 0xbdbddc61, 0x8b8b860d, 0x8a8a850f,
 294   0x707090e0, 0x3e3e427c, 0xb5b5c471, 0x6666aacc,
 295   0x4848d890, 0x03030506, 0xf6f601f7, 0x0e0e121c,
 296   0x6161a3c2, 0x35355f6a, 0x5757f9ae, 0xb9b9d069,
 297   0x86869117, 0xc1c15899, 0x1d1d273a, 0x9e9eb927,
 298   0xe1e138d9, 0xf8f813eb, 0x9898b32b, 0x11113322,
 299   0x6969bbd2, 0xd9d970a9, 0x8e8e8907, 0x9494a733,
 300   0x9b9bb62d, 0x1e1e223c, 0x87879215, 0xe9e920c9,
 301   0xcece4987, 0x5555ffaa, 0x28287850, 0xdfdf7aa5,
 302   0x8c8c8f03, 0xa1a1f859, 0x89898009, 0x0d0d171a,
 303   0xbfbfda65, 0xe6e631d7, 0x4242c684, 0x6868b8d0,
 304   0x4141c382, 0x9999b029, 0x2d2d775a, 0x0f0f111e,
 305   0xb0b0cb7b, 0x5454fca8, 0xbbbbd66d, 0x16163a2c,
 306 };
 307
 308 __device__ __constant__ u32 te4[256] =
 309 {
 310   0x63636363, 0x7c7c7c7c, 0x77777777, 0x7b7b7b7b,
 311   0xf2f2f2f2, 0x6b6b6b6b, 0x6f6f6f6f, 0xc5c5c5c5,
 312   0x30303030, 0x01010101, 0x67676767, 0x2b2b2b2b,
 313   0xfefefefe, 0xd7d7d7d7, 0xabababab, 0x76767676,
 314   0xcacacaca, 0x82828282, 0xc9c9c9c9, 0x7d7d7d7d,
 315   0xfafafafa, 0x59595959, 0x47474747, 0xf0f0f0f0,
 316   0xadadadad, 0xd4d4d4d4, 0xa2a2a2a2, 0xafafafaf,
 317   0x9c9c9c9c, 0xa4a4a4a4, 0x72727272, 0xc0c0c0c0,
 318   0xb7b7b7b7, 0xfdfdfdfd, 0x93939393, 0x26262626,
 319   0x36363636, 0x3f3f3f3f, 0xf7f7f7f7, 0xcccccccc,
 320   0x34343434, 0xa5a5a5a5, 0xe5e5e5e5, 0xf1f1f1f1,
 321   0x71717171, 0xd8d8d8d8, 0x31313131, 0x15151515,
 322   0x04040404, 0xc7c7c7c7, 0x23232323, 0xc3c3c3c3,
 323   0x18181818, 0x96969696, 0x05050505, 0x9a9a9a9a,
 324   0x07070707, 0x12121212, 0x80808080, 0xe2e2e2e2,
 325   0xebebebeb, 0x27272727, 0xb2b2b2b2, 0x75757575,
 326   0x09090909, 0x83838383, 0x2c2c2c2c, 0x1a1a1a1a,
 327   0x1b1b1b1b, 0x6e6e6e6e, 0x5a5a5a5a, 0xa0a0a0a0,
 328   0x52525252, 0x3b3b3b3b, 0xd6d6d6d6, 0xb3b3b3b3,
 329   0x29292929, 0xe3e3e3e3, 0x2f2f2f2f, 0x84848484,
 330   0x53535353, 0xd1d1d1d1, 0x00000000, 0xedededed,
 331   0x20202020, 0xfcfcfcfc, 0xb1b1b1b1, 0x5b5b5b5b,
 332   0x6a6a6a6a, 0xcbcbcbcb, 0xbebebebe, 0x39393939,
 333   0x4a4a4a4a, 0x4c4c4c4c, 0x58585858, 0xcfcfcfcf,
 334   0xd0d0d0d0, 0xefefefef, 0xaaaaaaaa, 0xfbfbfbfb,
 335   0x43434343, 0x4d4d4d4d, 0x33333333, 0x85858585,
 336   0x45454545, 0xf9f9f9f9, 0x02020202, 0x7f7f7f7f,
 337   0x50505050, 0x3c3c3c3c, 0x9f9f9f9f, 0xa8a8a8a8,
 338   0x51515151, 0xa3a3a3a3, 0x40404040, 0x8f8f8f8f,
 339   0x92929292, 0x9d9d9d9d, 0x38383838, 0xf5f5f5f5,
 340   0xbcbcbcbc, 0xb6b6b6b6, 0xdadadada, 0x21212121,
 341   0x10101010, 0xffffffff, 0xf3f3f3f3, 0xd2d2d2d2,
 342   0xcdcdcdcd, 0x0c0c0c0c, 0x13131313, 0xecececec,
 343   0x5f5f5f5f, 0x97979797, 0x44444444, 0x17171717,
 344   0xc4c4c4c4, 0xa7a7a7a7, 0x7e7e7e7e, 0x3d3d3d3d,
 345   0x64646464, 0x5d5d5d5d, 0x19191919, 0x73737373,
 346   0x60606060, 0x81818181, 0x4f4f4f4f, 0xdcdcdcdc,
 347   0x22222222, 0x2a2a2a2a, 0x90909090, 0x88888888,
 348   0x46464646, 0xeeeeeeee, 0xb8b8b8b8, 0x14141414,
 349   0xdededede, 0x5e5e5e5e, 0x0b0b0b0b, 0xdbdbdbdb,
 350   0xe0e0e0e0, 0x32323232, 0x3a3a3a3a, 0x0a0a0a0a,
 351   0x49494949, 0x06060606, 0x24242424, 0x5c5c5c5c,
 352   0xc2c2c2c2, 0xd3d3d3d3, 0xacacacac, 0x62626262,
 353   0x91919191, 0x95959595, 0xe4e4e4e4, 0x79797979,
 354   0xe7e7e7e7, 0xc8c8c8c8, 0x37373737, 0x6d6d6d6d,
 355   0x8d8d8d8d, 0xd5d5d5d5, 0x4e4e4e4e, 0xa9a9a9a9,
 356   0x6c6c6c6c, 0x56565656, 0xf4f4f4f4, 0xeaeaeaea,
 357   0x65656565, 0x7a7a7a7a, 0xaeaeaeae, 0x08080808,
 358   0xbabababa, 0x78787878, 0x25252525, 0x2e2e2e2e,
 359   0x1c1c1c1c, 0xa6a6a6a6, 0xb4b4b4b4, 0xc6c6c6c6,
 360   0xe8e8e8e8, 0xdddddddd, 0x74747474, 0x1f1f1f1f,
 361   0x4b4b4b4b, 0xbdbdbdbd, 0x8b8b8b8b, 0x8a8a8a8a,
 362   0x70707070, 0x3e3e3e3e, 0xb5b5b5b5, 0x66666666,
 363   0x48484848, 0x03030303, 0xf6f6f6f6, 0x0e0e0e0e,
 364   0x61616161, 0x35353535, 0x57575757, 0xb9b9b9b9,
 365   0x86868686, 0xc1c1c1c1, 0x1d1d1d1d, 0x9e9e9e9e,
 366   0xe1e1e1e1, 0xf8f8f8f8, 0x98989898, 0x11111111,
 367   0x69696969, 0xd9d9d9d9, 0x8e8e8e8e, 0x94949494,
 368   0x9b9b9b9b, 0x1e1e1e1e, 0x87878787, 0xe9e9e9e9,
 369   0xcececece, 0x55555555, 0x28282828, 0xdfdfdfdf,
 370   0x8c8c8c8c, 0xa1a1a1a1, 0x89898989, 0x0d0d0d0d,
 371   0xbfbfbfbf, 0xe6e6e6e6, 0x42424242, 0x68686868,
 372   0x41414141, 0x99999999, 0x2d2d2d2d, 0x0f0f0f0f,
 373   0xb0b0b0b0, 0x54545454, 0xbbbbbbbb, 0x16161616,
 374 };
 375
 376 __device__ __constant__ u32 td0[256] =
 377 {
 378   0x51f4a750, 0x7e416553, 0x1a17a4c3, 0x3a275e96,
 379   0x3bab6bcb, 0x1f9d45f1, 0xacfa58ab, 0x4be30393,
 380   0x2030fa55, 0xad766df6, 0x88cc7691, 0xf5024c25,
 381   0x4fe5d7fc, 0xc52acbd7, 0x26354480, 0xb562a38f,
 382   0xdeb15a49, 0x25ba1b67, 0x45ea0e98, 0x5dfec0e1,
 383   0xc32f7502, 0x814cf012, 0x8d4697a3, 0x6bd3f9c6,
 384   0x038f5fe7, 0x15929c95, 0xbf6d7aeb, 0x955259da,
 385   0xd4be832d, 0x587421d3, 0x49e06929, 0x8ec9c844,
 386   0x75c2896a, 0xf48e7978, 0x99583e6b, 0x27b971dd,
 387   0xbee14fb6, 0xf088ad17, 0xc920ac66, 0x7dce3ab4,
 388   0x63df4a18, 0xe51a3182, 0x97513360, 0x62537f45,
 389   0xb16477e0, 0xbb6bae84, 0xfe81a01c, 0xf9082b94,
 390   0x70486858, 0x8f45fd19, 0x94de6c87, 0x527bf8b7,
 391   0xab73d323, 0x724b02e2, 0xe31f8f57, 0x6655ab2a,
 392   0xb2eb2807, 0x2fb5c203, 0x86c57b9a, 0xd33708a5,
 393   0x302887f2, 0x23bfa5b2, 0x02036aba, 0xed16825c,
 394   0x8acf1c2b, 0xa779b492, 0xf307f2f0, 0x4e69e2a1,
 395   0x65daf4cd, 0x0605bed5, 0xd134621f, 0xc4a6fe8a,
 396   0x342e539d, 0xa2f355a0, 0x058ae132, 0xa4f6eb75,
 397   0x0b83ec39, 0x4060efaa, 0x5e719f06, 0xbd6e1051,
 398   0x3e218af9, 0x96dd063d, 0xdd3e05ae, 0x4de6bd46,
 399   0x91548db5, 0x71c45d05, 0x0406d46f, 0x605015ff,
 400   0x1998fb24, 0xd6bde997, 0x894043cc, 0x67d99e77,
 401   0xb0e842bd, 0x07898b88, 0xe7195b38, 0x79c8eedb,
 402   0xa17c0a47, 0x7c420fe9, 0xf8841ec9, 0x00000000,
 403   0x09808683, 0x322bed48, 0x1e1170ac, 0x6c5a724e,
 404   0xfd0efffb, 0x0f853856, 0x3daed51e, 0x362d3927,
 405   0x0a0fd964, 0x685ca621, 0x9b5b54d1, 0x24362e3a,
 406   0x0c0a67b1, 0x9357e70f, 0xb4ee96d2, 0x1b9b919e,
 407   0x80c0c54f, 0x61dc20a2, 0x5a774b69, 0x1c121a16,
 408   0xe293ba0a, 0xc0a02ae5, 0x3c22e043, 0x121b171d,
 409   0x0e090d0b, 0xf28bc7ad, 0x2db6a8b9, 0x141ea9c8,
 410   0x57f11985, 0xaf75074c, 0xee99ddbb, 0xa37f60fd,
 411   0xf701269f, 0x5c72f5bc, 0x44663bc5, 0x5bfb7e34,
 412   0x8b432976, 0xcb23c6dc, 0xb6edfc68, 0xb8e4f163,
 413   0xd731dcca, 0x42638510, 0x13972240, 0x84c61120,
 414   0x854a247d, 0xd2bb3df8, 0xaef93211, 0xc729a16d,
 415   0x1d9e2f4b, 0xdcb230f3, 0x0d8652ec, 0x77c1e3d0,
 416   0x2bb3166c, 0xa970b999, 0x119448fa, 0x47e96422,
 417   0xa8fc8cc4, 0xa0f03f1a, 0x567d2cd8, 0x223390ef,
 418   0x87494ec7, 0xd938d1c1, 0x8ccaa2fe, 0x98d40b36,
 419   0xa6f581cf, 0xa57ade28, 0xdab78e26, 0x3fadbfa4,
 420   0x2c3a9de4, 0x5078920d, 0x6a5fcc9b, 0x547e4662,
 421   0xf68d13c2, 0x90d8b8e8, 0x2e39f75e, 0x82c3aff5,
 422   0x9f5d80be, 0x69d0937c, 0x6fd52da9, 0xcf2512b3,
 423   0xc8ac993b, 0x10187da7, 0xe89c636e, 0xdb3bbb7b,
 424   0xcd267809, 0x6e5918f4, 0xec9ab701, 0x834f9aa8,
 425   0xe6956e65, 0xaaffe67e, 0x21bccf08, 0xef15e8e6,
 426   0xbae79bd9, 0x4a6f36ce, 0xea9f09d4, 0x29b07cd6,
 427   0x31a4b2af, 0x2a3f2331, 0xc6a59430, 0x35a266c0,
 428   0x744ebc37, 0xfc82caa6, 0xe090d0b0, 0x33a7d815,
 429   0xf104984a, 0x41ecdaf7, 0x7fcd500e, 0x1791f62f,
 430   0x764dd68d, 0x43efb04d, 0xccaa4d54, 0xe49604df,
 431   0x9ed1b5e3, 0x4c6a881b, 0xc12c1fb8, 0x4665517f,
 432   0x9d5eea04, 0x018c355d, 0xfa877473, 0xfb0b412e,
 433   0xb3671d5a, 0x92dbd252, 0xe9105633, 0x6dd64713,
 434   0x9ad7618c, 0x37a10c7a, 0x59f8148e, 0xeb133c89,
 435   0xcea927ee, 0xb761c935, 0xe11ce5ed, 0x7a47b13c,
 436   0x9cd2df59, 0x55f2733f, 0x1814ce79, 0x73c737bf,
 437   0x53f7cdea, 0x5ffdaa5b, 0xdf3d6f14, 0x7844db86,
 438   0xcaaff381, 0xb968c43e, 0x3824342c, 0xc2a3405f,
 439   0x161dc372, 0xbce2250c, 0x283c498b, 0xff0d9541,
 440   0x39a80171, 0x080cb3de, 0xd8b4e49c, 0x6456c190,
 441   0x7bcb8461, 0xd532b670, 0x486c5c74, 0xd0b85742,
 442 };
 443
 444 __device__ __constant__ u32 td1[256] =
 445 {
 446   0x5051f4a7, 0x537e4165, 0xc31a17a4, 0x963a275e,
 447   0xcb3bab6b, 0xf11f9d45, 0xabacfa58, 0x934be303,
 448   0x552030fa, 0xf6ad766d, 0x9188cc76, 0x25f5024c,
 449   0xfc4fe5d7, 0xd7c52acb, 0x80263544, 0x8fb562a3,
 450   0x49deb15a, 0x6725ba1b, 0x9845ea0e, 0xe15dfec0,
 451   0x02c32f75, 0x12814cf0, 0xa38d4697, 0xc66bd3f9,
 452   0xe7038f5f, 0x9515929c, 0xebbf6d7a, 0xda955259,
 453   0x2dd4be83, 0xd3587421, 0x2949e069, 0x448ec9c8,
 454   0x6a75c289, 0x78f48e79, 0x6b99583e, 0xdd27b971,
 455   0xb6bee14f, 0x17f088ad, 0x66c920ac, 0xb47dce3a,
 456   0x1863df4a, 0x82e51a31, 0x60975133, 0x4562537f,
 457   0xe0b16477, 0x84bb6bae, 0x1cfe81a0, 0x94f9082b,
 458   0x58704868, 0x198f45fd, 0x8794de6c, 0xb7527bf8,
 459   0x23ab73d3, 0xe2724b02, 0x57e31f8f, 0x2a6655ab,
 460   0x07b2eb28, 0x032fb5c2, 0x9a86c57b, 0xa5d33708,
 461   0xf2302887, 0xb223bfa5, 0xba02036a, 0x5ced1682,
 462   0x2b8acf1c, 0x92a779b4, 0xf0f307f2, 0xa14e69e2,
 463   0xcd65daf4, 0xd50605be, 0x1fd13462, 0x8ac4a6fe,
 464   0x9d342e53, 0xa0a2f355, 0x32058ae1, 0x75a4f6eb,
 465   0x390b83ec, 0xaa4060ef, 0x065e719f, 0x51bd6e10,
 466   0xf93e218a, 0x3d96dd06, 0xaedd3e05, 0x464de6bd,
 467   0xb591548d, 0x0571c45d, 0x6f0406d4, 0xff605015,
 468   0x241998fb, 0x97d6bde9, 0xcc894043, 0x7767d99e,
 469   0xbdb0e842, 0x8807898b, 0x38e7195b, 0xdb79c8ee,
 470   0x47a17c0a, 0xe97c420f, 0xc9f8841e, 0x00000000,
 471   0x83098086, 0x48322bed, 0xac1e1170, 0x4e6c5a72,
 472   0xfbfd0eff, 0x560f8538, 0x1e3daed5, 0x27362d39,
 473   0x640a0fd9, 0x21685ca6, 0xd19b5b54, 0x3a24362e,
 474   0xb10c0a67, 0x0f9357e7, 0xd2b4ee96, 0x9e1b9b91,
 475   0x4f80c0c5, 0xa261dc20, 0x695a774b, 0x161c121a,
 476   0x0ae293ba, 0xe5c0a02a, 0x433c22e0, 0x1d121b17,
 477   0x0b0e090d, 0xadf28bc7, 0xb92db6a8, 0xc8141ea9,
 478   0x8557f119, 0x4caf7507, 0xbbee99dd, 0xfda37f60,
 479   0x9ff70126, 0xbc5c72f5, 0xc544663b, 0x345bfb7e,
 480   0x768b4329, 0xdccb23c6, 0x68b6edfc, 0x63b8e4f1,
 481   0xcad731dc, 0x10426385, 0x40139722, 0x2084c611,
 482   0x7d854a24, 0xf8d2bb3d, 0x11aef932, 0x6dc729a1,
 483   0x4b1d9e2f, 0xf3dcb230, 0xec0d8652, 0xd077c1e3,
 484   0x6c2bb316, 0x99a970b9, 0xfa119448, 0x2247e964,
 485   0xc4a8fc8c, 0x1aa0f03f, 0xd8567d2c, 0xef223390,
 486   0xc787494e, 0xc1d938d1, 0xfe8ccaa2, 0x3698d40b,
 487   0xcfa6f581, 0x28a57ade, 0x26dab78e, 0xa43fadbf,
 488   0xe42c3a9d, 0x0d507892, 0x9b6a5fcc, 0x62547e46,
 489   0xc2f68d13, 0xe890d8b8, 0x5e2e39f7, 0xf582c3af,
 490   0xbe9f5d80, 0x7c69d093, 0xa96fd52d, 0xb3cf2512,
 491   0x3bc8ac99, 0xa710187d, 0x6ee89c63, 0x7bdb3bbb,
 492   0x09cd2678, 0xf46e5918, 0x01ec9ab7, 0xa8834f9a,
 493   0x65e6956e, 0x7eaaffe6, 0x0821bccf, 0xe6ef15e8,
 494   0xd9bae79b, 0xce4a6f36, 0xd4ea9f09, 0xd629b07c,
 495   0xaf31a4b2, 0x312a3f23, 0x30c6a594, 0xc035a266,
 496   0x37744ebc, 0xa6fc82ca, 0xb0e090d0, 0x1533a7d8,
 497   0x4af10498, 0xf741ecda, 0x0e7fcd50, 0x2f1791f6,
 498   0x8d764dd6, 0x4d43efb0, 0x54ccaa4d, 0xdfe49604,
 499   0xe39ed1b5, 0x1b4c6a88, 0xb8c12c1f, 0x7f466551,
 500   0x049d5eea, 0x5d018c35, 0x73fa8774, 0x2efb0b41,
 501   0x5ab3671d, 0x5292dbd2, 0x33e91056, 0x136dd647,
 502   0x8c9ad761, 0x7a37a10c, 0x8e59f814, 0x89eb133c,
 503   0xeecea927, 0x35b761c9, 0xede11ce5, 0x3c7a47b1,
 504   0x599cd2df, 0x3f55f273, 0x791814ce, 0xbf73c737,
 505   0xea53f7cd, 0x5b5ffdaa, 0x14df3d6f, 0x867844db,
 506   0x81caaff3, 0x3eb968c4, 0x2c382434, 0x5fc2a340,
 507   0x72161dc3, 0x0cbce225, 0x8b283c49, 0x41ff0d95,
 508   0x7139a801, 0xde080cb3, 0x9cd8b4e4, 0x906456c1,
 509   0x617bcb84, 0x70d532b6, 0x74486c5c, 0x42d0b857,
 510 };
 511
 512 __device__ __constant__ u32 td2[256] =
 513 {
 514   0xa75051f4, 0x65537e41, 0xa4c31a17, 0x5e963a27,
 515   0x6bcb3bab, 0x45f11f9d, 0x58abacfa, 0x03934be3,
 516   0xfa552030, 0x6df6ad76, 0x769188cc, 0x4c25f502,
 517   0xd7fc4fe5, 0xcbd7c52a, 0x44802635, 0xa38fb562,
 518   0x5a49deb1, 0x1b6725ba, 0x0e9845ea, 0xc0e15dfe,
 519   0x7502c32f, 0xf012814c, 0x97a38d46, 0xf9c66bd3,
 520   0x5fe7038f, 0x9c951592, 0x7aebbf6d, 0x59da9552,
 521   0x832dd4be, 0x21d35874, 0x692949e0, 0xc8448ec9,
 522   0x896a75c2, 0x7978f48e, 0x3e6b9958, 0x71dd27b9,
 523   0x4fb6bee1, 0xad17f088, 0xac66c920, 0x3ab47dce,
 524   0x4a1863df, 0x3182e51a, 0x33609751, 0x7f456253,
 525   0x77e0b164, 0xae84bb6b, 0xa01cfe81, 0x2b94f908,
 526   0x68587048, 0xfd198f45, 0x6c8794de, 0xf8b7527b,
 527   0xd323ab73, 0x02e2724b, 0x8f57e31f, 0xab2a6655,
 528   0x2807b2eb, 0xc2032fb5, 0x7b9a86c5, 0x08a5d337,
 529   0x87f23028, 0xa5b223bf, 0x6aba0203, 0x825ced16,
 530   0x1c2b8acf, 0xb492a779, 0xf2f0f307, 0xe2a14e69,
 531   0xf4cd65da, 0xbed50605, 0x621fd134, 0xfe8ac4a6,
 532   0x539d342e, 0x55a0a2f3, 0xe132058a, 0xeb75a4f6,
 533   0xec390b83, 0xefaa4060, 0x9f065e71, 0x1051bd6e,
 534   0x8af93e21, 0x063d96dd, 0x05aedd3e, 0xbd464de6,
 535   0x8db59154, 0x5d0571c4, 0xd46f0406, 0x15ff6050,
 536   0xfb241998, 0xe997d6bd, 0x43cc8940, 0x9e7767d9,
 537   0x42bdb0e8, 0x8b880789, 0x5b38e719, 0xeedb79c8,
 538   0x0a47a17c, 0x0fe97c42, 0x1ec9f884, 0x00000000,
 539   0x86830980, 0xed48322b, 0x70ac1e11, 0x724e6c5a,
 540   0xfffbfd0e, 0x38560f85, 0xd51e3dae, 0x3927362d,
 541   0xd9640a0f, 0xa621685c, 0x54d19b5b, 0x2e3a2436,
 542   0x67b10c0a, 0xe70f9357, 0x96d2b4ee, 0x919e1b9b,
 543   0xc54f80c0, 0x20a261dc, 0x4b695a77, 0x1a161c12,
 544   0xba0ae293, 0x2ae5c0a0, 0xe0433c22, 0x171d121b,
 545   0x0d0b0e09, 0xc7adf28b, 0xa8b92db6, 0xa9c8141e,
 546   0x198557f1, 0x074caf75, 0xddbbee99, 0x60fda37f,
 547   0x269ff701, 0xf5bc5c72, 0x3bc54466, 0x7e345bfb,
 548   0x29768b43, 0xc6dccb23, 0xfc68b6ed, 0xf163b8e4,
 549   0xdccad731, 0x85104263, 0x22401397, 0x112084c6,
 550   0x247d854a, 0x3df8d2bb, 0x3211aef9, 0xa16dc729,
 551   0x2f4b1d9e, 0x30f3dcb2, 0x52ec0d86, 0xe3d077c1,
 552   0x166c2bb3, 0xb999a970, 0x48fa1194, 0x642247e9,
 553   0x8cc4a8fc, 0x3f1aa0f0, 0x2cd8567d, 0x90ef2233,
 554   0x4ec78749, 0xd1c1d938, 0xa2fe8cca, 0x0b3698d4,
 555   0x81cfa6f5, 0xde28a57a, 0x8e26dab7, 0xbfa43fad,
 556   0x9de42c3a, 0x920d5078, 0xcc9b6a5f, 0x4662547e,
 557   0x13c2f68d, 0xb8e890d8, 0xf75e2e39, 0xaff582c3,
 558   0x80be9f5d, 0x937c69d0, 0x2da96fd5, 0x12b3cf25,
 559   0x993bc8ac, 0x7da71018, 0x636ee89c, 0xbb7bdb3b,
 560   0x7809cd26, 0x18f46e59, 0xb701ec9a, 0x9aa8834f,
 561   0x6e65e695, 0xe67eaaff, 0xcf0821bc, 0xe8e6ef15,
 562   0x9bd9bae7, 0x36ce4a6f, 0x09d4ea9f, 0x7cd629b0,
 563   0xb2af31a4, 0x23312a3f, 0x9430c6a5, 0x66c035a2,
 564   0xbc37744e, 0xcaa6fc82, 0xd0b0e090, 0xd81533a7,
 565   0x984af104, 0xdaf741ec, 0x500e7fcd, 0xf62f1791,
 566   0xd68d764d, 0xb04d43ef, 0x4d54ccaa, 0x04dfe496,
 567   0xb5e39ed1, 0x881b4c6a, 0x1fb8c12c, 0x517f4665,
 568   0xea049d5e, 0x355d018c, 0x7473fa87, 0x412efb0b,
 569   0x1d5ab367, 0xd25292db, 0x5633e910, 0x47136dd6,
 570   0x618c9ad7, 0x0c7a37a1, 0x148e59f8, 0x3c89eb13,
 571   0x27eecea9, 0xc935b761, 0xe5ede11c, 0xb13c7a47,
 572   0xdf599cd2, 0x733f55f2, 0xce791814, 0x37bf73c7,
 573   0xcdea53f7, 0xaa5b5ffd, 0x6f14df3d, 0xdb867844,
 574   0xf381caaf, 0xc43eb968, 0x342c3824, 0x405fc2a3,
 575   0xc372161d, 0x250cbce2, 0x498b283c, 0x9541ff0d,
 576   0x017139a8, 0xb3de080c, 0xe49cd8b4, 0xc1906456,
 577   0x84617bcb, 0xb670d532, 0x5c74486c, 0x5742d0b8,
 578 };
 579
 580 __device__ __constant__ u32 td3[256] =
 581 {
 582   0xf4a75051, 0x4165537e, 0x17a4c31a, 0x275e963a,
 583   0xab6bcb3b, 0x9d45f11f, 0xfa58abac, 0xe303934b,
 584   0x30fa5520, 0x766df6ad, 0xcc769188, 0x024c25f5,
 585   0xe5d7fc4f, 0x2acbd7c5, 0x35448026, 0x62a38fb5,
 586   0xb15a49de, 0xba1b6725, 0xea0e9845, 0xfec0e15d,
 587   0x2f7502c3, 0x4cf01281, 0x4697a38d, 0xd3f9c66b,
 588   0x8f5fe703, 0x929c9515, 0x6d7aebbf, 0x5259da95,
 589   0xbe832dd4, 0x7421d358, 0xe0692949, 0xc9c8448e,
 590   0xc2896a75, 0x8e7978f4, 0x583e6b99, 0xb971dd27,
 591   0xe14fb6be, 0x88ad17f0, 0x20ac66c9, 0xce3ab47d,
 592   0xdf4a1863, 0x1a3182e5, 0x51336097, 0x537f4562,
 593   0x6477e0b1, 0x6bae84bb, 0x81a01cfe, 0x082b94f9,
 594   0x48685870, 0x45fd198f, 0xde6c8794, 0x7bf8b752,
 595   0x73d323ab, 0x4b02e272, 0x1f8f57e3, 0x55ab2a66,
 596   0xeb2807b2, 0xb5c2032f, 0xc57b9a86, 0x3708a5d3,
 597   0x2887f230, 0xbfa5b223, 0x036aba02, 0x16825ced,
 598   0xcf1c2b8a, 0x79b492a7, 0x07f2f0f3, 0x69e2a14e,
 599   0xdaf4cd65, 0x05bed506, 0x34621fd1, 0xa6fe8ac4,
 600   0x2e539d34, 0xf355a0a2, 0x8ae13205, 0xf6eb75a4,
 601   0x83ec390b, 0x60efaa40, 0x719f065e, 0x6e1051bd,
 602   0x218af93e, 0xdd063d96, 0x3e05aedd, 0xe6bd464d,
 603   0x548db591, 0xc45d0571, 0x06d46f04, 0x5015ff60,
 604   0x98fb2419, 0xbde997d6, 0x4043cc89, 0xd99e7767,
 605   0xe842bdb0, 0x898b8807, 0x195b38e7, 0xc8eedb79,
 606   0x7c0a47a1, 0x420fe97c, 0x841ec9f8, 0x00000000,
 607   0x80868309, 0x2bed4832, 0x1170ac1e, 0x5a724e6c,
 608   0x0efffbfd, 0x8538560f, 0xaed51e3d, 0x2d392736,
 609   0x0fd9640a, 0x5ca62168, 0x5b54d19b, 0x362e3a24,
 610   0x0a67b10c, 0x57e70f93, 0xee96d2b4, 0x9b919e1b,
 611   0xc0c54f80, 0xdc20a261, 0x774b695a, 0x121a161c,
 612   0x93ba0ae2, 0xa02ae5c0, 0x22e0433c, 0x1b171d12,
 613   0x090d0b0e, 0x8bc7adf2, 0xb6a8b92d, 0x1ea9c814,
 614   0xf1198557, 0x75074caf, 0x99ddbbee, 0x7f60fda3,
 615   0x01269ff7, 0x72f5bc5c, 0x663bc544, 0xfb7e345b,
 616   0x4329768b, 0x23c6dccb, 0xedfc68b6, 0xe4f163b8,
 617   0x31dccad7, 0x63851042, 0x97224013, 0xc6112084,
 618   0x4a247d85, 0xbb3df8d2, 0xf93211ae, 0x29a16dc7,
 619   0x9e2f4b1d, 0xb230f3dc, 0x8652ec0d, 0xc1e3d077,
 620   0xb3166c2b, 0x70b999a9, 0x9448fa11, 0xe9642247,
 621   0xfc8cc4a8, 0xf03f1aa0, 0x7d2cd856, 0x3390ef22,
 622   0x494ec787, 0x38d1c1d9, 0xcaa2fe8c, 0xd40b3698,
 623   0xf581cfa6, 0x7ade28a5, 0xb78e26da, 0xadbfa43f,
 624   0x3a9de42c, 0x78920d50, 0x5fcc9b6a, 0x7e466254,
 625   0x8d13c2f6, 0xd8b8e890, 0x39f75e2e, 0xc3aff582,
 626   0x5d80be9f, 0xd0937c69, 0xd52da96f, 0x2512b3cf,
 627   0xac993bc8, 0x187da710, 0x9c636ee8, 0x3bbb7bdb,
 628   0x267809cd, 0x5918f46e, 0x9ab701ec, 0x4f9aa883,
 629   0x956e65e6, 0xffe67eaa, 0xbccf0821, 0x15e8e6ef,
 630   0xe79bd9ba, 0x6f36ce4a, 0x9f09d4ea, 0xb07cd629,
 631   0xa4b2af31, 0x3f23312a, 0xa59430c6, 0xa266c035,
 632   0x4ebc3774, 0x82caa6fc, 0x90d0b0e0, 0xa7d81533,
 633   0x04984af1, 0xecdaf741, 0xcd500e7f, 0x91f62f17,
 634   0x4dd68d76, 0xefb04d43, 0xaa4d54cc, 0x9604dfe4,
 635   0xd1b5e39e, 0x6a881b4c, 0x2c1fb8c1, 0x65517f46,
 636   0x5eea049d, 0x8c355d01, 0x877473fa, 0x0b412efb,
 637   0x671d5ab3, 0xdbd25292, 0x105633e9, 0xd647136d,
 638   0xd7618c9a, 0xa10c7a37, 0xf8148e59, 0x133c89eb,
 639   0xa927eece, 0x61c935b7, 0x1ce5ede1, 0x47b13c7a,
 640   0xd2df599c, 0xf2733f55, 0x14ce7918, 0xc737bf73,
 641   0xf7cdea53, 0xfdaa5b5f, 0x3d6f14df, 0x44db8678,
 642   0xaff381ca, 0x68c43eb9, 0x24342c38, 0xa3405fc2,
 643   0x1dc37216, 0xe2250cbc, 0x3c498b28, 0x0d9541ff,
 644   0xa8017139, 0x0cb3de08, 0xb4e49cd8, 0x56c19064,
 645   0xcb84617b, 0x32b670d5, 0x6c5c7448, 0xb85742d0,
 646 };
 647
 648 __device__ __constant__ u32 td4[256] =
 649 {
 650   0x52525252, 0x09090909, 0x6a6a6a6a, 0xd5d5d5d5,
 651   0x30303030, 0x36363636, 0xa5a5a5a5, 0x38383838,
 652   0xbfbfbfbf, 0x40404040, 0xa3a3a3a3, 0x9e9e9e9e,
 653   0x81818181, 0xf3f3f3f3, 0xd7d7d7d7, 0xfbfbfbfb,
 654   0x7c7c7c7c, 0xe3e3e3e3, 0x39393939, 0x82828282,
 655   0x9b9b9b9b, 0x2f2f2f2f, 0xffffffff, 0x87878787,
 656   0x34343434, 0x8e8e8e8e, 0x43434343, 0x44444444,
 657   0xc4c4c4c4, 0xdededede, 0xe9e9e9e9, 0xcbcbcbcb,
 658   0x54545454, 0x7b7b7b7b, 0x94949494, 0x32323232,
 659   0xa6a6a6a6, 0xc2c2c2c2, 0x23232323, 0x3d3d3d3d,
 660   0xeeeeeeee, 0x4c4c4c4c, 0x95959595, 0x0b0b0b0b,
 661   0x42424242, 0xfafafafa, 0xc3c3c3c3, 0x4e4e4e4e,
 662   0x08080808, 0x2e2e2e2e, 0xa1a1a1a1, 0x66666666,
 663   0x28282828, 0xd9d9d9d9, 0x24242424, 0xb2b2b2b2,
 664   0x76767676, 0x5b5b5b5b, 0xa2a2a2a2, 0x49494949,
 665   0x6d6d6d6d, 0x8b8b8b8b, 0xd1d1d1d1, 0x25252525,
 666   0x72727272, 0xf8f8f8f8, 0xf6f6f6f6, 0x64646464,
 667   0x86868686, 0x68686868, 0x98989898, 0x16161616,
 668   0xd4d4d4d4, 0xa4a4a4a4, 0x5c5c5c5c, 0xcccccccc,
 669   0x5d5d5d5d, 0x65656565, 0xb6b6b6b6, 0x92929292,
 670   0x6c6c6c6c, 0x70707070, 0x48484848, 0x50505050,
 671   0xfdfdfdfd, 0xedededed, 0xb9b9b9b9, 0xdadadada,
 672   0x5e5e5e5e, 0x15151515, 0x46464646, 0x57575757,
 673   0xa7a7a7a7, 0x8d8d8d8d, 0x9d9d9d9d, 0x84848484,
 674   0x90909090, 0xd8d8d8d8, 0xabababab, 0x00000000,
 675   0x8c8c8c8c, 0xbcbcbcbc, 0xd3d3d3d3, 0x0a0a0a0a,
 676   0xf7f7f7f7, 0xe4e4e4e4, 0x58585858, 0x05050505,
 677   0xb8b8b8b8, 0xb3b3b3b3, 0x45454545, 0x06060606,
 678   0xd0d0d0d0, 0x2c2c2c2c, 0x1e1e1e1e, 0x8f8f8f8f,
 679   0xcacacaca, 0x3f3f3f3f, 0x0f0f0f0f, 0x02020202,
 680   0xc1c1c1c1, 0xafafafaf, 0xbdbdbdbd, 0x03030303,
 681   0x01010101, 0x13131313, 0x8a8a8a8a, 0x6b6b6b6b,
 682   0x3a3a3a3a, 0x91919191, 0x11111111, 0x41414141,
 683   0x4f4f4f4f, 0x67676767, 0xdcdcdcdc, 0xeaeaeaea,
 684   0x97979797, 0xf2f2f2f2, 0xcfcfcfcf, 0xcececece,
 685   0xf0f0f0f0, 0xb4b4b4b4, 0xe6e6e6e6, 0x73737373,
 686   0x96969696, 0xacacacac, 0x74747474, 0x22222222,
 687   0xe7e7e7e7, 0xadadadad, 0x35353535, 0x85858585,
 688   0xe2e2e2e2, 0xf9f9f9f9, 0x37373737, 0xe8e8e8e8,
 689   0x1c1c1c1c, 0x75757575, 0xdfdfdfdf, 0x6e6e6e6e,
 690   0x47474747, 0xf1f1f1f1, 0x1a1a1a1a, 0x71717171,
 691   0x1d1d1d1d, 0x29292929, 0xc5c5c5c5, 0x89898989,
 692   0x6f6f6f6f, 0xb7b7b7b7, 0x62626262, 0x0e0e0e0e,
 693   0xaaaaaaaa, 0x18181818, 0xbebebebe, 0x1b1b1b1b,
 694   0xfcfcfcfc, 0x56565656, 0x3e3e3e3e, 0x4b4b4b4b,
 695   0xc6c6c6c6, 0xd2d2d2d2, 0x79797979, 0x20202020,
 696   0x9a9a9a9a, 0xdbdbdbdb, 0xc0c0c0c0, 0xfefefefe,
 697   0x78787878, 0xcdcdcdcd, 0x5a5a5a5a, 0xf4f4f4f4,
 698   0x1f1f1f1f, 0xdddddddd, 0xa8a8a8a8, 0x33333333,
 699   0x88888888, 0x07070707, 0xc7c7c7c7, 0x31313131,
 700   0xb1b1b1b1, 0x12121212, 0x10101010, 0x59595959,
 701   0x27272727, 0x80808080, 0xecececec, 0x5f5f5f5f,
 702   0x60606060, 0x51515151, 0x7f7f7f7f, 0xa9a9a9a9,
 703   0x19191919, 0xb5b5b5b5, 0x4a4a4a4a, 0x0d0d0d0d,
 704   0x2d2d2d2d, 0xe5e5e5e5, 0x7a7a7a7a, 0x9f9f9f9f,
 705   0x93939393, 0xc9c9c9c9, 0x9c9c9c9c, 0xefefefef,
 706   0xa0a0a0a0, 0xe0e0e0e0, 0x3b3b3b3b, 0x4d4d4d4d,
 707   0xaeaeaeae, 0x2a2a2a2a, 0xf5f5f5f5, 0xb0b0b0b0,
 708   0xc8c8c8c8, 0xebebebeb, 0xbbbbbbbb, 0x3c3c3c3c,
 709   0x83838383, 0x53535353, 0x99999999, 0x61616161,
 710   0x17171717, 0x2b2b2b2b, 0x04040404, 0x7e7e7e7e,
 711   0xbabababa, 0x77777777, 0xd6d6d6d6, 0x26262626,
 712   0xe1e1e1e1, 0x69696969, 0x14141414, 0x63636363,
 713   0x55555555, 0x21212121, 0x0c0c0c0c, 0x7d7d7d7d,
 714 };
 715
 716 __device__ __constant__ u32 rcon[] =
 717 {
 718   0x01000000, 0x02000000, 0x04000000, 0x08000000,
 719   0x10000000, 0x20000000, 0x40000000, 0x80000000,
 720   0x1b000000, 0x36000000,
 721 };
 722
 723 __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])
 724 {
 725   rek[0] = userkey[0];
 726   rek[1] = userkey[1];
 727   rek[2] = userkey[2];
 728   rek[3] = userkey[3];
 729   rek[4] = userkey[4];
 730   rek[5] = userkey[5];
 731   rek[6] = userkey[6];
 732   rek[7] = userkey[7];
 733
 734   int i;
 735   int j;
 736
 737   i = 0;
 738   j = 0;
 739
 740   u32 run = 1;
 741
 742   while (run)
 743   {
 744     u32 temp = rek[j +  7];
 745
 746     rek[j +  8] = rek[j +  0]
 747            ^ (s_te2[(temp >> 16) & 0xff] & 0xff000000)
 748            ^ (s_te3[(temp >>  8) & 0xff] & 0x00ff0000)
 749            ^ (s_te0[(temp >>  0) & 0xff] & 0x0000ff00)
 750            ^ (s_te1[(temp >> 24) & 0xff] & 0x000000ff)
 751            ^ rcon[i];
 752
 753     rek[j +  9] = rek[j +  1] ^ rek[j +  8];
 754     rek[j + 10] = rek[j +  2] ^ rek[j +  9];
 755     rek[j + 11] = rek[j +  3] ^ rek[j + 10];
 756
 757     if (++i == 7)
 758     {
 759       run = 0;
 760       continue;
 761     }
 762
 763     temp = rek[j + 11];
 764
 765     rek[j + 12] = rek[j +  4]
 766            ^ (s_te2[(temp >> 24) & 0xff] & 0xff000000)
 767            ^ (s_te3[(temp >> 16) & 0xff] & 0x00ff0000)
 768            ^ (s_te0[(temp >>  8) & 0xff] & 0x0000ff00)
 769            ^ (s_te1[(temp >>  0) & 0xff] & 0x000000ff);
 770
 771     rek[j + 13] = rek[j +  5] ^ rek[j + 12];
 772     rek[j + 14] = rek[j +  6] ^ rek[j + 13];
 773     rek[j + 15] = rek[j +  7] ^ rek[j + 14];
 774
 775     j += 8;
 776   }
 777 }
 778
 779 __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])
 780 {
 781   for (u32 i = 0, j = 56; i < j; i += 4, j -= 4)
 782   {
 783     u32 temp;
 784
 785     temp = rdk[i + 0]; rdk[i + 0] = rdk[j + 0]; rdk[j + 0] = temp;
 786     temp = rdk[i + 1]; rdk[i + 1] = rdk[j + 1]; rdk[j + 1] = temp;
 787     temp = rdk[i + 2]; rdk[i + 2] = rdk[j + 2]; rdk[j + 2] = temp;
 788     temp = rdk[i + 3]; rdk[i + 3] = rdk[j + 3]; rdk[j + 3] = temp;
 789   }
 790
 791   for (u32 i = 1, j = 4; i < 14; i += 1, j += 4)
 792   {
 793     rdk[j + 0] =
 794       s_td0[s_te1[(rdk[j + 0] >> 24) & 0xff] & 0xff] ^
 795       s_td1[s_te1[(rdk[j + 0] >> 16) & 0xff] & 0xff] ^
 796       s_td2[s_te1[(rdk[j + 0] >>  8) & 0xff] & 0xff] ^
 797       s_td3[s_te1[(rdk[j + 0] >>  0) & 0xff] & 0xff];
 798
 799     rdk[j + 1] =
 800       s_td0[s_te1[(rdk[j + 1] >> 24) & 0xff] & 0xff] ^
 801       s_td1[s_te1[(rdk[j + 1] >> 16) & 0xff] & 0xff] ^
 802       s_td2[s_te1[(rdk[j + 1] >>  8) & 0xff] & 0xff] ^
 803       s_td3[s_te1[(rdk[j + 1] >>  0) & 0xff] & 0xff];
 804
 805     rdk[j + 2] =
 806       s_td0[s_te1[(rdk[j + 2] >> 24) & 0xff] & 0xff] ^
 807       s_td1[s_te1[(rdk[j + 2] >> 16) & 0xff] & 0xff] ^
 808       s_td2[s_te1[(rdk[j + 2] >>  8) & 0xff] & 0xff] ^
 809       s_td3[s_te1[(rdk[j + 2] >>  0) & 0xff] & 0xff];
 810
 811     rdk[j + 3] =
 812       s_td0[s_te1[(rdk[j + 3] >> 24) & 0xff] & 0xff] ^
 813       s_td1[s_te1[(rdk[j + 3] >> 16) & 0xff] & 0xff] ^
 814       s_td2[s_te1[(rdk[j + 3] >>  8) & 0xff] & 0xff] ^
 815       s_td3[s_te1[(rdk[j + 3] >>  0) & 0xff] & 0xff];
 816   }
 817 }
 818
 819 __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])
 820 {
 821   u32 s0 = in[0] ^ rdk[0];
 822   u32 s1 = in[1] ^ rdk[1];
 823   u32 s2 = in[2] ^ rdk[2];
 824   u32 s3 = in[3] ^ rdk[3];
 825
 826   u32 t0;
 827   u32 t1;
 828   u32 t2;
 829   u32 t3;
 830
 831   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[ 4];
 832   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[ 5];
 833   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[ 6];
 834   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[ 7];
 835   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[ 8];
 836   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[ 9];
 837   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[10];
 838   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[11];
 839   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[12];
 840   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[13];
 841   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[14];
 842   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[15];
 843   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[16];
 844   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[17];
 845   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[18];
 846   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[19];
 847   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[20];
 848   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[21];
 849   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[22];
 850   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[23];
 851   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[24];
 852   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[25];
 853   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[26];
 854   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[27];
 855   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[28];
 856   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[29];
 857   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[30];
 858   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[31];
 859   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[32];
 860   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[33];
 861   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[34];
 862   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[35];
 863   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[36];
 864   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[37];
 865   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[38];
 866   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[39];
 867   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[40];
 868   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[41];
 869   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[42];
 870   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[43];
 871   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[44];
 872   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[45];
 873   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[46];
 874   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[47];
 875   s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >>  8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[48];
 876   s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >>  8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[49];
 877   s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >>  8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[50];
 878   s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >>  8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[51];
 879   t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >>  8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[52];
 880   t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >>  8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[53];
 881   t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >>  8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[54];
 882   t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >>  8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[55];
 883
 884   out[0] = (s_td4[(t0 >> 24) & 0xff] & 0xff000000)
 885          ^ (s_td4[(t3 >> 16) & 0xff] & 0x00ff0000)
 886          ^ (s_td4[(t2 >>  8) & 0xff] & 0x0000ff00)
 887          ^ (s_td4[(t1 >>  0) & 0xff] & 0x000000ff)
 888          ^ rdk[56];
 889
 890   out[1] = (s_td4[(t1 >> 24) & 0xff] & 0xff000000)
 891          ^ (s_td4[(t0 >> 16) & 0xff] & 0x00ff0000)
 892          ^ (s_td4[(t3 >>  8) & 0xff] & 0x0000ff00)
 893          ^ (s_td4[(t2 >>  0) & 0xff] & 0x000000ff)
 894          ^ rdk[57];
 895
 896   out[2] = (s_td4[(t2 >> 24) & 0xff] & 0xff000000)
 897          ^ (s_td4[(t1 >> 16) & 0xff] & 0x00ff0000)
 898          ^ (s_td4[(t0 >>  8) & 0xff] & 0x0000ff00)
 899          ^ (s_td4[(t3 >>  0) & 0xff] & 0x000000ff)
 900          ^ rdk[58];
 901
 902   out[3] = (s_td4[(t3 >> 24) & 0xff] & 0xff000000)
 903          ^ (s_td4[(t2 >> 16) & 0xff] & 0x00ff0000)
 904          ^ (s_td4[(t1 >>  8) & 0xff] & 0x0000ff00)
 905          ^ (s_td4[(t0 >>  0) & 0xff] & 0x000000ff)
 906          ^ rdk[59];
 907 }
 908
 909 __device__ __constant__ u32 k_sha256[64] =
 910 {
 911   SHA256C00, SHA256C01, SHA256C02, SHA256C03,
 912   SHA256C04, SHA256C05, SHA256C06, SHA256C07,
 913   SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
 914   SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
 915   SHA256C10, SHA256C11, SHA256C12, SHA256C13,
 916   SHA256C14, SHA256C15, SHA256C16, SHA256C17,
 917   SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
 918   SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
 919   SHA256C20, SHA256C21, SHA256C22, SHA256C23,
 920   SHA256C24, SHA256C25, SHA256C26, SHA256C27,
 921   SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
 922   SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
 923   SHA256C30, SHA256C31, SHA256C32, SHA256C33,
 924   SHA256C34, SHA256C35, SHA256C36, SHA256C37,
 925   SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
 926   SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
 927 };
 928
 929 __device__ static void sha256_transform (const u32x w[16], u32x digest[8])
 930 {
 931   u32x a = digest[0];
 932   u32x b = digest[1];
 933   u32x c = digest[2];
 934   u32x d = digest[3];
 935   u32x e = digest[4];
 936   u32x f = digest[5];
 937   u32x g = digest[6];
 938   u32x h = digest[7];
 939
 940   u32x w0_t = swap_workaround (w[ 0]);
 941   u32x w1_t = swap_workaround (w[ 1]);
 942   u32x w2_t = swap_workaround (w[ 2]);
 943   u32x w3_t = swap_workaround (w[ 3]);
 944   u32x w4_t = swap_workaround (w[ 4]);
 945   u32x w5_t = swap_workaround (w[ 5]);
 946   u32x w6_t = swap_workaround (w[ 6]);
 947   u32x w7_t = swap_workaround (w[ 7]);
 948   u32x w8_t = swap_workaround (w[ 8]);
 949   u32x w9_t = swap_workaround (w[ 9]);
 950   u32x wa_t = swap_workaround (w[10]);
 951   u32x wb_t = swap_workaround (w[11]);
 952   u32x wc_t = swap_workaround (w[12]);
 953   u32x wd_t = swap_workaround (w[13]);
 954   u32x we_t = swap_workaround (w[14]);
 955   u32x wf_t = swap_workaround (w[15]);
 956
 957   #define ROUND_EXPAND()                            \
 958   {                                                 \
 959     w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t);  \
 960     w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t);  \
 961     w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t);  \
 962     w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t);  \
 963     w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t);  \
 964     w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t);  \
 965     w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t);  \
 966     w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t);  \
 967     w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t);  \
 968     w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t);  \
 969     wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t);  \
 970     wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t);  \
 971     wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t);  \
 972     wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t);  \
 973     we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t);  \
 974     wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t);  \
 975   }
 976
 977   #define ROUND_STEP(i)                                                                   \
 978   {                                                                                       \
 979     SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i +  0]); \
 980     SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i +  1]); \
 981     SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i +  2]); \
 982     SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i +  3]); \
 983     SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i +  4]); \
 984     SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i +  5]); \
 985     SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i +  6]); \
 986     SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i +  7]); \
 987     SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i +  8]); \
 988     SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i +  9]); \
 989     SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
 990     SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
 991     SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
 992     SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
 993     SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
 994     SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
 995   }
 996
 997   ROUND_STEP (0);
 998
 999   for (int i = 16; i < 64; i += 16)
1000   {
1001     ROUND_EXPAND (); ROUND_STEP (i);
1002   }
1003
1004   digest[0] += a;
1005   digest[1] += b;
1006   digest[2] += c;
1007   digest[3] += d;
1008   digest[4] += e;
1009   digest[5] += f;
1010   digest[6] += g;
1011   digest[7] += h;
1012 }
1013
1014 __device__ const u32 crc32tab[0x100] =
1015 {
1016   0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
1017   0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
1018   0x0edb8832, 0x79dcb8a4, 0xe0d5e91e, 0x97d2d988,
1019   0x09b64c2b, 0x7eb17cbd, 0xe7b82d07, 0x90bf1d91,
1020   0x1db71064, 0x6ab020f2, 0xf3b97148, 0x84be41de,
1021   0x1adad47d, 0x6ddde4eb, 0xf4d4b551, 0x83d385c7,
1022   0x136c9856, 0x646ba8c0, 0xfd62f97a, 0x8a65c9ec,
1023   0x14015c4f, 0x63066cd9, 0xfa0f3d63, 0x8d080df5,
1024   0x3b6e20c8, 0x4c69105e, 0xd56041e4, 0xa2677172,
1025   0x3c03e4d1, 0x4b04d447, 0xd20d85fd, 0xa50ab56b,
1026   0x35b5a8fa, 0x42b2986c, 0xdbbbc9d6, 0xacbcf940,
1027   0x32d86ce3, 0x45df5c75, 0xdcd60dcf, 0xabd13d59,
1028   0x26d930ac, 0x51de003a, 0xc8d75180, 0xbfd06116,
1029   0x21b4f4b5, 0x56b3c423, 0xcfba9599, 0xb8bda50f,
1030   0x2802b89e, 0x5f058808, 0xc60cd9b2, 0xb10be924,
1031   0x2f6f7c87, 0x58684c11, 0xc1611dab, 0xb6662d3d,
1032   0x76dc4190, 0x01db7106, 0x98d220bc, 0xefd5102a,
1033   0x71b18589, 0x06b6b51f, 0x9fbfe4a5, 0xe8b8d433,
1034   0x7807c9a2, 0x0f00f934, 0x9609a88e, 0xe10e9818,
1035   0x7f6a0dbb, 0x086d3d2d, 0x91646c97, 0xe6635c01,
1036   0x6b6b51f4, 0x1c6c6162, 0x856530d8, 0xf262004e,
1037   0x6c0695ed, 0x1b01a57b, 0x8208f4c1, 0xf50fc457,
1038   0x65b0d9c6, 0x12b7e950, 0x8bbeb8ea, 0xfcb9887c,
1039   0x62dd1ddf, 0x15da2d49, 0x8cd37cf3, 0xfbd44c65,
1040   0x4db26158, 0x3ab551ce, 0xa3bc0074, 0xd4bb30e2,
1041   0x4adfa541, 0x3dd895d7, 0xa4d1c46d, 0xd3d6f4fb,
1042   0x4369e96a, 0x346ed9fc, 0xad678846, 0xda60b8d0,
1043   0x44042d73, 0x33031de5, 0xaa0a4c5f, 0xdd0d7cc9,
1044   0x5005713c, 0x270241aa, 0xbe0b1010, 0xc90c2086,
1045   0x5768b525, 0x206f85b3, 0xb966d409, 0xce61e49f,
1046   0x5edef90e, 0x29d9c998, 0xb0d09822, 0xc7d7a8b4,
1047   0x59b33d17, 0x2eb40d81, 0xb7bd5c3b, 0xc0ba6cad,
1048   0xedb88320, 0x9abfb3b6, 0x03b6e20c, 0x74b1d29a,
1049   0xead54739, 0x9dd277af, 0x04db2615, 0x73dc1683,
1050   0xe3630b12, 0x94643b84, 0x0d6d6a3e, 0x7a6a5aa8,
1051   0xe40ecf0b, 0x9309ff9d, 0x0a00ae27, 0x7d079eb1,
1052   0xf00f9344, 0x8708a3d2, 0x1e01f268, 0x6906c2fe,
1053   0xf762575d, 0x806567cb, 0x196c3671, 0x6e6b06e7,
1054   0xfed41b76, 0x89d32be0, 0x10da7a5a, 0x67dd4acc,
1055   0xf9b9df6f, 0x8ebeeff9, 0x17b7be43, 0x60b08ed5,
1056   0xd6d6a3e8, 0xa1d1937e, 0x38d8c2c4, 0x4fdff252,
1057   0xd1bb67f1, 0xa6bc5767, 0x3fb506dd, 0x48b2364b,
1058   0xd80d2bda, 0xaf0a1b4c, 0x36034af6, 0x41047a60,
1059   0xdf60efc3, 0xa867df55, 0x316e8eef, 0x4669be79,
1060   0xcb61b38c, 0xbc66831a, 0x256fd2a0, 0x5268e236,
1061   0xcc0c7795, 0xbb0b4703, 0x220216b9, 0x5505262f,
1062   0xc5ba3bbe, 0xb2bd0b28, 0x2bb45a92, 0x5cb36a04,
1063   0xc2d7ffa7, 0xb5d0cf31, 0x2cd99e8b, 0x5bdeae1d,
1064   0x9b64c2b0, 0xec63f226, 0x756aa39c, 0x026d930a,
1065   0x9c0906a9, 0xeb0e363f, 0x72076785, 0x05005713,
1066   0x95bf4a82, 0xe2b87a14, 0x7bb12bae, 0x0cb61b38,
1067   0x92d28e9b, 0xe5d5be0d, 0x7cdcefb7, 0x0bdbdf21,
1068   0x86d3d2d4, 0xf1d4e242, 0x68ddb3f8, 0x1fda836e,
1069   0x81be16cd, 0xf6b9265b, 0x6fb077e1, 0x18b74777,
1070   0x88085ae6, 0xff0f6a70, 0x66063bca, 0x11010b5c,
1071   0x8f659eff, 0xf862ae69, 0x616bffd3, 0x166ccf45,
1072   0xa00ae278, 0xd70dd2ee, 0x4e048354, 0x3903b3c2,
1073   0xa7672661, 0xd06016f7, 0x4969474d, 0x3e6e77db,
1074   0xaed16a4a, 0xd9d65adc, 0x40df0b66, 0x37d83bf0,
1075   0xa9bcae53, 0xdebb9ec5, 0x47b2cf7f, 0x30b5ffe9,
1076   0xbdbdf21c, 0xcabac28a, 0x53b39330, 0x24b4a3a6,
1077   0xbad03605, 0xcdd70693, 0x54de5729, 0x23d967bf,
1078   0xb3667a2e, 0xc4614ab8, 0x5d681b02, 0x2a6f2b94,
1079   0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
1080 };
1081
1082 __device__ static u32x round_crc32 (u32x a, const u32x v)
1083 {
1084   const u32x k = (a ^ v) & 0xff;
1085
1086   const u32x s = a >> 8;
1087
1088   #ifdef VECT_SIZE1
1089   a    = crc32tab[k];
1090   #endif
1091
1092   #ifdef VECT_SIZE2
1093   a.x = crc32tab[k.x];
1094   a.y = crc32tab[k.y];
1095   #endif
1096
1097   #ifdef VECT_SIZE4
1098   a.x = crc32tab[k.x];
1099   a.y = crc32tab[k.y];
1100   a.z = crc32tab[k.z];
1101   a.w = crc32tab[k.w];
1102   #endif
1103
1104   a ^= s;
1105
1106   return a;
1107 }
1108
1109 __device__ static u32x crc32 (const u32x w[16], const u32 pw_len, const u32 iv)
1110 {
1111   u32x a = iv ^ ~0;
1112
1113   if (pw_len >=  1) a = round_crc32 (a, w[0] >>  0);
1114   if (pw_len >=  2) a = round_crc32 (a, w[0] >>  8);
1115   if (pw_len >=  3) a = round_crc32 (a, w[0] >> 16);
1116   if (pw_len >=  4) a = round_crc32 (a, w[0] >> 24);
1117
1118   for (u32 i = 4, j = 1; i < pw_len; i += 4, j += 1)
1119   {
1120     if (pw_len >= (i + 1)) a = round_crc32 (a, w[j] >>  0);
1121     if (pw_len >= (i + 2)) a = round_crc32 (a, w[j] >>  8);
1122     if (pw_len >= (i + 3)) a = round_crc32 (a, w[j] >> 16);
1123     if (pw_len >= (i + 4)) a = round_crc32 (a, w[j] >> 24);
1124   }
1125
1126   return ~a;
1127 }
1128
1129 __device__ static void bzero16 (u32x block[16])
1130 {
1131   block[ 0] = 0;
1132   block[ 1] = 0;
1133   block[ 2] = 0;
1134   block[ 3] = 0;
1135   block[ 4] = 0;
1136   block[ 5] = 0;
1137   block[ 6] = 0;
1138   block[ 7] = 0;
1139   block[ 8] = 0;
1140   block[ 9] = 0;
1141   block[10] = 0;
1142   block[11] = 0;
1143   block[12] = 0;
1144   block[13] = 0;
1145   block[14] = 0;
1146   block[15] = 0;
1147 }
1148
1149 __device__ static u32 memcat8c (u32x block[16], const u32 block_len, const u32x append[4], const u32 append_len, u32x digest[8])
1150 {
1151   const u32 div = block_len / 4;
1152
1153   u32x tmp0;
1154   u32x tmp1;
1155   u32x tmp2;
1156
1157   #if __CUDA_ARCH__ >= 200
1158
1159   const int offset_minus_4 = 4 - (block_len & 3);
1160
1161   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
1162
1163   tmp0 = __byte_perm (        0, append[0], selector);
1164   tmp1 = __byte_perm (append[0], append[1], selector);
1165   tmp2 = __byte_perm (append[1],         0, selector);
1166
1167   #endif
1168
1169   u32x carry[2] = { 0, 0 };
1170
1171   switch (div)
1172   {
1173     case  0:  block[ 0] |= tmp0;
1174               block[ 1]  = tmp1;
1175               block[ 2]  = tmp2;
1176               break;
1177     case  1:  block[ 1] |= tmp0;
1178               block[ 2]  = tmp1;
1179               block[ 3]  = tmp2;
1180               break;
1181     case  2:  block[ 2] |= tmp0;
1182               block[ 3]  = tmp1;
1183               block[ 4]  = tmp2;
1184               break;
1185     case  3:  block[ 3] |= tmp0;
1186               block[ 4]  = tmp1;
1187               block[ 5]  = tmp2;
1188               break;
1189     case  4:  block[ 4] |= tmp0;
1190               block[ 5]  = tmp1;
1191               block[ 6]  = tmp2;
1192               break;
1193     case  5:  block[ 5] |= tmp0;
1194               block[ 6]  = tmp1;
1195               block[ 7]  = tmp2;
1196               break;
1197     case  6:  block[ 6] |= tmp0;
1198               block[ 7]  = tmp1;
1199               block[ 8]  = tmp2;
1200               break;
1201     case  7:  block[ 7] |= tmp0;
1202               block[ 8]  = tmp1;
1203               block[ 9]  = tmp2;
1204               break;
1205     case  8:  block[ 8] |= tmp0;
1206               block[ 9]  = tmp1;
1207               block[10]  = tmp2;
1208               break;
1209     case  9:  block[ 9] |= tmp0;
1210               block[10]  = tmp1;
1211               block[11]  = tmp2;
1212               break;
1213     case 10:  block[10] |= tmp0;
1214               block[11]  = tmp1;
1215               block[12]  = tmp2;
1216               break;
1217     case 11:  block[11] |= tmp0;
1218               block[12]  = tmp1;
1219               block[13]  = tmp2;
1220               break;
1221     case 12:  block[12] |= tmp0;
1222               block[13]  = tmp1;
1223               block[14]  = tmp2;
1224               break;
1225     case 13:  block[13] |= tmp0;
1226               block[14]  = tmp1;
1227               block[15]  = tmp2;
1228               break;
1229     case 14:  block[14] |= tmp0;
1230               block[15]  = tmp1;
1231               carry[ 0]  = tmp2;
1232               break;
1233     case 15:  block[15] |= tmp0;
1234               carry[ 0]  = tmp1;
1235               carry[ 1]  = tmp2;
1236               break;
1237   }
1238
1239   u32 new_len = block_len + append_len;
1240
1241   if (new_len >= 64)
1242   {
1243     new_len -= 64;
1244
1245     sha256_transform (block, digest);
1246
1247     bzero16 (block);
1248
1249     block[0] = carry[0];
1250     block[1] = carry[1];
1251   }
1252
1253   return new_len;
1254 }
1255
1256 __device__ static u32 memcat32c (u32x block[16], const u32 block_len, const u32x append[4], const u32 append_len, u32x digest[8])
1257 {
1258   const u32 div = block_len / 4;
1259
1260   u32x tmp0;
1261   u32x tmp1;
1262   u32x tmp2;
1263   u32x tmp3;
1264   u32x tmp4;
1265   u32x tmp5;
1266   u32x tmp6;
1267   u32x tmp7;
1268   u32x tmp8;
1269
1270   #if __CUDA_ARCH__ >= 200
1271
1272   const int offset_minus_4 = 4 - (block_len & 3);
1273
1274   const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
1275
1276   tmp0 = __byte_perm (        0, append[0], selector);
1277   tmp1 = __byte_perm (append[0], append[1], selector);
1278   tmp2 = __byte_perm (append[1], append[2], selector);
1279   tmp3 = __byte_perm (append[2], append[3], selector);
1280   tmp4 = __byte_perm (append[3], append[4], selector);
1281   tmp5 = __byte_perm (append[4], append[5], selector);
1282   tmp6 = __byte_perm (append[5], append[6], selector);
1283   tmp7 = __byte_perm (append[6], append[7], selector);
1284   tmp8 = __byte_perm (append[7],         0, selector);
1285
1286   #endif
1287
1288   u32x carry[8] = { 0, 0, 0, 0, 0, 0, 0, 0 };
1289
1290   switch (div)
1291   {
1292     case  0:  block[ 0] |= tmp0;
1293               block[ 1]  = tmp1;
1294               block[ 2]  = tmp2;
1295               block[ 3]  = tmp3;
1296               block[ 4]  = tmp4;
1297               block[ 5]  = tmp5;
1298               block[ 6]  = tmp6;
1299               block[ 7]  = tmp7;
1300               block[ 8]  = tmp8;
1301               break;
1302     case  1:  block[ 1] |= tmp0;
1303               block[ 2]  = tmp1;
1304               block[ 3]  = tmp2;
1305               block[ 4]  = tmp3;
1306               block[ 5]  = tmp4;
1307               block[ 6]  = tmp5;
1308               block[ 7]  = tmp6;
1309               block[ 8]  = tmp7;
1310               block[ 9]  = tmp8;
1311               break;
1312     case  2:  block[ 2] |= tmp0;
1313               block[ 3]  = tmp1;
1314               block[ 4]  = tmp2;
1315               block[ 5]  = tmp3;
1316               block[ 6]  = tmp4;
1317               block[ 7]  = tmp5;
1318               block[ 8]  = tmp6;
1319               block[ 9]  = tmp7;
1320               block[10]  = tmp8;
1321               break;
1322     case  3:  block[ 3] |= tmp0;
1323               block[ 4]  = tmp1;
1324               block[ 5]  = tmp2;
1325               block[ 6]  = tmp3;
1326               block[ 7]  = tmp4;
1327               block[ 8]  = tmp5;
1328               block[ 9]  = tmp6;
1329               block[10]  = tmp7;
1330               block[11]  = tmp8;
1331               break;
1332     case  4:  block[ 4] |= tmp0;
1333               block[ 5]  = tmp1;
1334               block[ 6]  = tmp2;
1335               block[ 7]  = tmp3;
1336               block[ 8]  = tmp4;
1337               block[ 9]  = tmp5;
1338               block[10]  = tmp6;
1339               block[11]  = tmp7;
1340               block[12]  = tmp8;
1341               break;
1342     case  5:  block[ 5] |= tmp0;
1343               block[ 6]  = tmp1;
1344               block[ 7]  = tmp2;
1345               block[ 8]  = tmp3;
1346               block[ 9]  = tmp4;
1347               block[10]  = tmp5;
1348               block[11]  = tmp6;
1349               block[12]  = tmp7;
1350               block[13]  = tmp8;
1351               break;
1352     case  6:  block[ 6] |= tmp0;
1353               block[ 7]  = tmp1;
1354               block[ 8]  = tmp2;
1355               block[ 9]  = tmp3;
1356               block[10]  = tmp4;
1357               block[11]  = tmp5;
1358               block[12]  = tmp6;
1359               block[13]  = tmp7;
1360               block[14]  = tmp8;
1361               break;
1362     case  7:  block[ 7] |= tmp0;
1363               block[ 8]  = tmp1;
1364               block[ 9]  = tmp2;
1365               block[10]  = tmp3;
1366               block[11]  = tmp4;
1367               block[12]  = tmp5;
1368               block[13]  = tmp6;
1369               block[14]  = tmp7;
1370               block[15]  = tmp8;
1371               break;
1372     case  8:  block[ 8] |= tmp0;
1373               block[ 9]  = tmp1;
1374               block[10]  = tmp2;
1375               block[11]  = tmp3;
1376               block[12]  = tmp4;
1377               block[13]  = tmp5;
1378               block[14]  = tmp6;
1379               block[15]  = tmp7;
1380               carry[ 0]  = tmp8;
1381               break;
1382     case  9:  block[ 9] |= tmp0;
1383               block[10]  = tmp1;
1384               block[11]  = tmp2;
1385               block[12]  = tmp3;
1386               block[13]  = tmp4;
1387               block[14]  = tmp5;
1388               block[15]  = tmp6;
1389               carry[ 0]  = tmp7;
1390               carry[ 1]  = tmp8;
1391               break;
1392     case 10:  block[10] |= tmp0;
1393               block[11]  = tmp1;
1394               block[12]  = tmp2;
1395               block[13]  = tmp3;
1396               block[14]  = tmp4;
1397               block[15]  = tmp5;
1398               carry[ 0]  = tmp6;
1399               carry[ 1]  = tmp7;
1400               carry[ 2]  = tmp8;
1401               break;
1402     case 11:  block[11] |= tmp0;
1403               block[12]  = tmp1;
1404               block[13]  = tmp2;
1405               block[14]  = tmp3;
1406               block[15]  = tmp4;
1407               carry[ 0]  = tmp5;
1408               carry[ 1]  = tmp6;
1409               carry[ 2]  = tmp7;
1410               carry[ 3]  = tmp8;
1411               break;
1412     case 12:  block[12] |= tmp0;
1413               block[13]  = tmp1;
1414               block[14]  = tmp2;
1415               block[15]  = tmp3;
1416               carry[ 0]  = tmp4;
1417               carry[ 1]  = tmp5;
1418               carry[ 2]  = tmp6;
1419               carry[ 3]  = tmp7;
1420               carry[ 4]  = tmp8;
1421               break;
1422     case 13:  block[13] |= tmp0;
1423               block[14]  = tmp1;
1424               block[15]  = tmp2;
1425               carry[ 0]  = tmp3;
1426               carry[ 1]  = tmp4;
1427               carry[ 2]  = tmp5;
1428               carry[ 3]  = tmp6;
1429               carry[ 4]  = tmp7;
1430               carry[ 5]  = tmp8;
1431               break;
1432     case 14:  block[14] |= tmp0;
1433               block[15]  = tmp1;
1434               carry[ 0]  = tmp2;
1435               carry[ 1]  = tmp3;
1436               carry[ 2]  = tmp4;
1437               carry[ 3]  = tmp5;
1438               carry[ 4]  = tmp6;
1439               carry[ 5]  = tmp7;
1440               carry[ 6]  = tmp8;
1441               break;
1442     case 15:  block[15] |= tmp0;
1443               carry[ 0]  = tmp1;
1444               carry[ 1]  = tmp2;
1445               carry[ 2]  = tmp3;
1446               carry[ 3]  = tmp4;
1447               carry[ 4]  = tmp5;
1448               carry[ 5]  = tmp6;
1449               carry[ 6]  = tmp7;
1450               carry[ 7]  = tmp8;
1451               break;
1452   }
1453
1454   u32 new_len = block_len + append_len;
1455
1456   if (new_len >= 64)
1457   {
1458     new_len -= 64;
1459
1460     sha256_transform (block, digest);
1461
1462     bzero16 (block);
1463
1464     block[0] = carry[0];
1465     block[1] = carry[1];
1466     block[2] = carry[2];
1467     block[3] = carry[3];
1468     block[4] = carry[4];
1469     block[5] = carry[5];
1470     block[6] = carry[6];
1471     block[7] = carry[7];
1472   }
1473
1474   return new_len;
1475 }
1476
1477 extern "C" __global__ void __launch_bounds__ (256, 1) m11600_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, seven_zip_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 seven_zip_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)
1478 {
1479   /**
1480    * base
1481    */
1482
1483   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1484
1485   if (gid >= gid_max) return;
1486
1487   /**
1488    * algo starts here already
1489    */
1490
1491   u32x dgst[8];
1492
1493   dgst[0] = SHA256M_A;
1494   dgst[1] = SHA256M_B;
1495   dgst[2] = SHA256M_C;
1496   dgst[3] = SHA256M_D;
1497   dgst[4] = SHA256M_E;
1498   dgst[5] = SHA256M_F;
1499   dgst[6] = SHA256M_G;
1500   dgst[7] = SHA256M_H;
1501
1502   u32x block[16];
1503
1504   bzero16 (block);
1505
1506   u32 block_len = 0;
1507   u32 final_len = 0;
1508
1509   /**
1510    * context save
1511    */
1512
1513   tmps[gid].dgst[0] = dgst[0];
1514   tmps[gid].dgst[1] = dgst[1];
1515   tmps[gid].dgst[2] = dgst[2];
1516   tmps[gid].dgst[3] = dgst[3];
1517   tmps[gid].dgst[4] = dgst[4];
1518   tmps[gid].dgst[5] = dgst[5];
1519   tmps[gid].dgst[6] = dgst[6];
1520   tmps[gid].dgst[7] = dgst[7];
1521
1522   tmps[gid].block[ 0] = block[ 0];
1523   tmps[gid].block[ 1] = block[ 1];
1524   tmps[gid].block[ 2] = block[ 2];
1525   tmps[gid].block[ 3] = block[ 3];
1526   tmps[gid].block[ 4] = block[ 4];
1527   tmps[gid].block[ 5] = block[ 5];
1528   tmps[gid].block[ 6] = block[ 6];
1529   tmps[gid].block[ 7] = block[ 7];
1530   tmps[gid].block[ 8] = block[ 8];
1531   tmps[gid].block[ 9] = block[ 9];
1532   tmps[gid].block[10] = block[10];
1533   tmps[gid].block[11] = block[11];
1534   tmps[gid].block[12] = block[12];
1535   tmps[gid].block[13] = block[13];
1536   tmps[gid].block[14] = block[14];
1537   tmps[gid].block[15] = block[15];
1538
1539   tmps[gid].block_len = block_len;
1540   tmps[gid].final_len = final_len;
1541 }
1542
1543 extern "C" __global__ void __launch_bounds__ (256, 1) m11600_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, seven_zip_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 seven_zip_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)
1544 {
1545   /**
1546    * base
1547    */
1548
1549   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1550
1551   if (gid >= gid_max) return;
1552
1553   u32x pw[8];
1554
1555   pw[0] = pws[gid].i[ 0];
1556   pw[1] = pws[gid].i[ 1];
1557   pw[2] = pws[gid].i[ 2];
1558   pw[3] = pws[gid].i[ 3];
1559   pw[4] = 0;
1560   pw[5] = 0;
1561   pw[6] = 0;
1562   pw[7] = 0;
1563
1564   u32 pw_len = pws[gid].pw_len;
1565
1566   make_unicode (&pw[0], &pw[0], &pw[4]);
1567
1568   pw_len *= 2;
1569
1570   /**
1571    * context load
1572    */
1573
1574   u32x dgst[8];
1575
1576   dgst[0] = tmps[gid].dgst[0];
1577   dgst[1] = tmps[gid].dgst[1];
1578   dgst[2] = tmps[gid].dgst[2];
1579   dgst[3] = tmps[gid].dgst[3];
1580   dgst[4] = tmps[gid].dgst[4];
1581   dgst[5] = tmps[gid].dgst[5];
1582   dgst[6] = tmps[gid].dgst[6];
1583   dgst[7] = tmps[gid].dgst[7];
1584
1585   u32x block[16];
1586
1587   block[ 0] = tmps[gid].block[ 0];
1588   block[ 1] = tmps[gid].block[ 1];
1589   block[ 2] = tmps[gid].block[ 2];
1590   block[ 3] = tmps[gid].block[ 3];
1591   block[ 4] = tmps[gid].block[ 4];
1592   block[ 5] = tmps[gid].block[ 5];
1593   block[ 6] = tmps[gid].block[ 6];
1594   block[ 7] = tmps[gid].block[ 7];
1595   block[ 8] = tmps[gid].block[ 8];
1596   block[ 9] = tmps[gid].block[ 9];
1597   block[10] = tmps[gid].block[10];
1598   block[11] = tmps[gid].block[11];
1599   block[12] = tmps[gid].block[12];
1600   block[13] = tmps[gid].block[13];
1601   block[14] = tmps[gid].block[14];
1602   block[15] = tmps[gid].block[15];
1603
1604   u32 block_len = tmps[gid].block_len;
1605   u32 final_len = tmps[gid].final_len;
1606
1607   /**
1608    * base
1609    */
1610
1611   for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1612   {
1613     u32 it[2];
1614
1615     it[0] = j;
1616     it[1] = 0;
1617
1618     block_len = memcat32c (block, block_len, pw, pw_len, dgst); final_len += pw_len;
1619     block_len = memcat8c  (block, block_len, it,      8, dgst); final_len += 8;
1620   }
1621
1622   /**
1623    * context save
1624    */
1625
1626   tmps[gid].dgst[0] = dgst[0];
1627   tmps[gid].dgst[1] = dgst[1];
1628   tmps[gid].dgst[2] = dgst[2];
1629   tmps[gid].dgst[3] = dgst[3];
1630   tmps[gid].dgst[4] = dgst[4];
1631   tmps[gid].dgst[5] = dgst[5];
1632   tmps[gid].dgst[6] = dgst[6];
1633   tmps[gid].dgst[7] = dgst[7];
1634
1635   tmps[gid].block[ 0] = block[ 0];
1636   tmps[gid].block[ 1] = block[ 1];
1637   tmps[gid].block[ 2] = block[ 2];
1638   tmps[gid].block[ 3] = block[ 3];
1639   tmps[gid].block[ 4] = block[ 4];
1640   tmps[gid].block[ 5] = block[ 5];
1641   tmps[gid].block[ 6] = block[ 6];
1642   tmps[gid].block[ 7] = block[ 7];
1643   tmps[gid].block[ 8] = block[ 8];
1644   tmps[gid].block[ 9] = block[ 9];
1645   tmps[gid].block[10] = block[10];
1646   tmps[gid].block[11] = block[11];
1647   tmps[gid].block[12] = block[12];
1648   tmps[gid].block[13] = block[13];
1649   tmps[gid].block[14] = block[14];
1650   tmps[gid].block[15] = block[15];
1651
1652   tmps[gid].block_len = block_len;
1653   tmps[gid].final_len = final_len;
1654 }
1655
1656 extern "C" __global__ void __launch_bounds__ (256, 1) m11600_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, seven_zip_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 seven_zip_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)
1657 {
1658   /**
1659    * base
1660    */
1661
1662   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1663   const u32 lid = threadIdx.x;
1664
1665   /**
1666    * aes shared
1667    */
1668
1669   __shared__ u32 s_td0[256];
1670   __shared__ u32 s_td1[256];
1671   __shared__ u32 s_td2[256];
1672   __shared__ u32 s_td3[256];
1673   __shared__ u32 s_td4[256];
1674
1675   __shared__ u32 s_te0[256];
1676   __shared__ u32 s_te1[256];
1677   __shared__ u32 s_te2[256];
1678   __shared__ u32 s_te3[256];
1679   __shared__ u32 s_te4[256];
1680
1681   s_td0[lid] = td0[lid];
1682   s_td1[lid] = td1[lid];
1683   s_td2[lid] = td2[lid];
1684   s_td3[lid] = td3[lid];
1685   s_td4[lid] = td4[lid];
1686
1687   s_te0[lid] = te0[lid];
1688   s_te1[lid] = te1[lid];
1689   s_te2[lid] = te2[lid];
1690   s_te3[lid] = te3[lid];
1691   s_te4[lid] = te4[lid];
1692
1693   __syncthreads ();
1694
1695   if (gid >= gid_max) return;
1696
1697   /**
1698    * context load
1699    */
1700
1701   u32x dgst[8];
1702
1703   dgst[0] = tmps[gid].dgst[0];
1704   dgst[1] = tmps[gid].dgst[1];
1705   dgst[2] = tmps[gid].dgst[2];
1706   dgst[3] = tmps[gid].dgst[3];
1707   dgst[4] = tmps[gid].dgst[4];
1708   dgst[5] = tmps[gid].dgst[5];
1709   dgst[6] = tmps[gid].dgst[6];
1710   dgst[7] = tmps[gid].dgst[7];
1711
1712   u32x block[16];
1713
1714   block[ 0] = tmps[gid].block[ 0];
1715   block[ 1] = tmps[gid].block[ 1];
1716   block[ 2] = tmps[gid].block[ 2];
1717   block[ 3] = tmps[gid].block[ 3];
1718   block[ 4] = tmps[gid].block[ 4];
1719   block[ 5] = tmps[gid].block[ 5];
1720   block[ 6] = tmps[gid].block[ 6];
1721   block[ 7] = tmps[gid].block[ 7];
1722   block[ 8] = tmps[gid].block[ 8];
1723   block[ 9] = tmps[gid].block[ 9];
1724   block[10] = tmps[gid].block[10];
1725   block[11] = tmps[gid].block[11];
1726   block[12] = tmps[gid].block[12];
1727   block[13] = tmps[gid].block[13];
1728   block[14] = tmps[gid].block[14];
1729   block[15] = tmps[gid].block[15];
1730
1731   u32 block_len = tmps[gid].block_len;
1732   u32 final_len = tmps[gid].final_len;
1733
1734   append_0x80_4 (block, block_len);
1735
1736   if (block_len >= 56)
1737   {
1738     sha256_transform (block, dgst);
1739
1740     bzero16 (block);
1741   }
1742
1743   block[15] = swap_workaround (final_len * 8);
1744
1745   sha256_transform (block, dgst);
1746
1747   /**
1748    * final key operations
1749    */
1750
1751   u32 iv[4];
1752
1753   iv[0] = esalt_bufs[salt_pos].iv_buf[0];
1754   iv[1] = esalt_bufs[salt_pos].iv_buf[1];
1755   iv[2] = esalt_bufs[salt_pos].iv_buf[2];
1756   iv[3] = esalt_bufs[salt_pos].iv_buf[3];
1757
1758   u32x ukey[8];
1759
1760   ukey[0] = dgst[0];
1761   ukey[1] = dgst[1];
1762   ukey[2] = dgst[2];
1763   ukey[3] = dgst[3];
1764   ukey[4] = dgst[4];
1765   ukey[5] = dgst[5];
1766   ukey[6] = dgst[6];
1767   ukey[7] = dgst[7];
1768
1769   #define KEYLEN 60
1770
1771   u32 rk[KEYLEN];
1772
1773   AES256_ExpandKey (ukey, rk, s_te0, s_te1, s_te2, s_te3, s_te4);
1774
1775   AES256_InvertKey (rk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
1776
1777   u32 crc = 0;
1778
1779   int data_len    = esalt_bufs[salt_pos].data_len;
1780   int unpack_size = esalt_bufs[salt_pos].unpack_size;
1781
1782   int i;
1783   int j;
1784
1785   for (i = 0, j = 0; i < data_len - 16; i += 16, j += 4)
1786   {
1787     u32 data[4];
1788
1789     data[0] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 0]);
1790     data[1] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 1]);
1791     data[2] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 2]);
1792     data[3] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 3]);
1793
1794     u32 out[4];
1795
1796     AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1797
1798     out[0] ^= iv[0];
1799     out[1] ^= iv[1];
1800     out[2] ^= iv[2];
1801     out[3] ^= iv[3];
1802
1803     iv[0] = data[0];
1804     iv[1] = data[1];
1805     iv[2] = data[2];
1806     iv[3] = data[3];
1807
1808     out[0] = swap_workaround (out[0]);
1809     out[1] = swap_workaround (out[1]);
1810     out[2] = swap_workaround (out[2]);
1811     out[3] = swap_workaround (out[3]);
1812
1813     crc = crc32 (out, 16, crc);
1814   }
1815
1816   u32 data[4];
1817
1818   data[0] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 0]);
1819   data[1] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 1]);
1820   data[2] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 2]);
1821   data[3] = swap_workaround (esalt_bufs[salt_pos].data_buf[j + 3]);
1822
1823   u32 out[4];
1824
1825   AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1826
1827   out[0] ^= iv[0];
1828   out[1] ^= iv[1];
1829   out[2] ^= iv[2];
1830   out[3] ^= iv[3];
1831
1832   iv[0] = data[0];
1833   iv[1] = data[1];
1834   iv[2] = data[2];
1835   iv[3] = data[3];
1836
1837   out[0] = swap_workaround (out[0]);
1838   out[1] = swap_workaround (out[1]);
1839   out[2] = swap_workaround (out[2]);
1840   out[3] = swap_workaround (out[3]);
1841
1842   const u32 margin = data_len - unpack_size;
1843
1844   const u32 left = 16 - margin;
1845
1846   crc = crc32 (out, left, crc);
1847
1848   // use padding attack in that case
1849
1850   if (margin >= 4)
1851   {
1852     switch (margin)
1853     {
1854       case 15:  out[0] &= 0xffffff00;
1855                 break;
1856       case 14:  out[0] &= 0xffff0000;
1857                 break;
1858       case 13:  out[0] &= 0xff000000;
1859                 break;
1860       case 12:  out[0]  = 0;
1861                 break;
1862       case 11:  out[0]  = 0;
1863                 out[1] &= 0xffffff00;
1864                 break;
1865       case 10:  out[0]  = 0;
1866                 out[1] &= 0xffff0000;
1867                 break;
1868       case  9:  out[0]  = 0;
1869                 out[1] &= 0xff000000;
1870                 break;
1871       case  8:  out[0]  = 0;
1872                 out[1]  = 0;
1873                 break;
1874       case  7:  out[0]  = 0;
1875                 out[1]  = 0;
1876                 out[2] &= 0xffffff00;
1877                 break;
1878       case  6:  out[0]  = 0;
1879                 out[1]  = 0;
1880                 out[2] &= 0xffff0000;
1881                 break;
1882       case  5:  out[0]  = 0;
1883                 out[1]  = 0;
1884                 out[2] &= 0xff000000;
1885                 break;
1886       case  4:  out[0]  = 0;
1887                 out[1]  = 0;
1888                 out[2]  = 0;
1889                 break;
1890       case  3:  out[0]  = 0;
1891                 out[1]  = 0;
1892                 out[2]  = 0;
1893                 out[3] &= 0xffffff00;
1894                 break;
1895       case  2:  out[0]  = 0;
1896                 out[1]  = 0;
1897                 out[2]  = 0;
1898                 out[3] &= 0xffff0000;
1899                 break;
1900       case  1:  out[0]  = 0;
1901                 out[1]  = 0;
1902                 out[2]  = 0;
1903                 out[3] &= 0xff000000;
1904                 break;
1905     }
1906
1907     if ((out[0] == 0) && (out[1] == 0) && (out[2] == 0) && (out[3] == 0))
1908     {
1909       mark_hash_s0 (plains_buf, hashes_shown, digests_offset + 0, gid, 0);
1910
1911       d_return_buf[lid] = 1;
1912     }
1913   }
1914
1915   const u32x r0 = crc;
1916   const u32x r1 = 0;
1917   const u32x r2 = 0;
1918   const u32x r3 = 0;
1919
1920   #define il_pos 0
1921
1922   #include VECT_COMPARE_M
1923 }