nv/m06232.cu

   1 /**
   2  * Author......: Jens Steube <jens.steube@gmail.com>
   3  * License.....: MIT
   4  */
   5
   6 #define _WHIRLPOOL_
   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 #include "gpu_aes256_nv.c"
  29 #include "gpu_twofish256_nv.c"
  30 #include "gpu_serpent256_nv.c"
  31
  32 #define R 10
  33
  34 #define BOX(S,n,i) (u32) ((S)[(n)][(i)])
  35
  36 __device__ __constant__ u32 Ch[8][256] =
  37 {
  38   {
  39     0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
  40     0x87872687, 0xb8b8dab8, 0x01010401, 0x4f4f214f,
  41     0x3636d836, 0xa6a6a2a6, 0xd2d26fd2, 0xf5f5f3f5,
  42     0x7979f979, 0x6f6fa16f, 0x91917e91, 0x52525552,
  43     0x60609d60, 0xbcbccabc, 0x9b9b569b, 0x8e8e028e,
  44     0xa3a3b6a3, 0x0c0c300c, 0x7b7bf17b, 0x3535d435,
  45     0x1d1d741d, 0xe0e0a7e0, 0xd7d77bd7, 0xc2c22fc2,
  46     0x2e2eb82e, 0x4b4b314b, 0xfefedffe, 0x57574157,
  47     0x15155415, 0x7777c177, 0x3737dc37, 0xe5e5b3e5,
  48     0x9f9f469f, 0xf0f0e7f0, 0x4a4a354a, 0xdada4fda,
  49     0x58587d58, 0xc9c903c9, 0x2929a429, 0x0a0a280a,
  50     0xb1b1feb1, 0xa0a0baa0, 0x6b6bb16b, 0x85852e85,
  51     0xbdbdcebd, 0x5d5d695d, 0x10104010, 0xf4f4f7f4,
  52     0xcbcb0bcb, 0x3e3ef83e, 0x05051405, 0x67678167,
  53     0xe4e4b7e4, 0x27279c27, 0x41411941, 0x8b8b168b,
  54     0xa7a7a6a7, 0x7d7de97d, 0x95956e95, 0xd8d847d8,
  55     0xfbfbcbfb, 0xeeee9fee, 0x7c7ced7c, 0x66668566,
  56     0xdddd53dd, 0x17175c17, 0x47470147, 0x9e9e429e,
  57     0xcaca0fca, 0x2d2db42d, 0xbfbfc6bf, 0x07071c07,
  58     0xadad8ead, 0x5a5a755a, 0x83833683, 0x3333cc33,
  59     0x63639163, 0x02020802, 0xaaaa92aa, 0x7171d971,
  60     0xc8c807c8, 0x19196419, 0x49493949, 0xd9d943d9,
  61     0xf2f2eff2, 0xe3e3abe3, 0x5b5b715b, 0x88881a88,
  62     0x9a9a529a, 0x26269826, 0x3232c832, 0xb0b0fab0,
  63     0xe9e983e9, 0x0f0f3c0f, 0xd5d573d5, 0x80803a80,
  64     0xbebec2be, 0xcdcd13cd, 0x3434d034, 0x48483d48,
  65     0xffffdbff, 0x7a7af57a, 0x90907a90, 0x5f5f615f,
  66     0x20208020, 0x6868bd68, 0x1a1a681a, 0xaeae82ae,
  67     0xb4b4eab4, 0x54544d54, 0x93937693, 0x22228822,
  68     0x64648d64, 0xf1f1e3f1, 0x7373d173, 0x12124812,
  69     0x40401d40, 0x08082008, 0xc3c32bc3, 0xecec97ec,
  70     0xdbdb4bdb, 0xa1a1bea1, 0x8d8d0e8d, 0x3d3df43d,
  71     0x97976697, 0x00000000, 0xcfcf1bcf, 0x2b2bac2b,
  72     0x7676c576, 0x82823282, 0xd6d67fd6, 0x1b1b6c1b,
  73     0xb5b5eeb5, 0xafaf86af, 0x6a6ab56a, 0x50505d50,
  74     0x45450945, 0xf3f3ebf3, 0x3030c030, 0xefef9bef,
  75     0x3f3ffc3f, 0x55554955, 0xa2a2b2a2, 0xeaea8fea,
  76     0x65658965, 0xbabad2ba, 0x2f2fbc2f, 0xc0c027c0,
  77     0xdede5fde, 0x1c1c701c, 0xfdfdd3fd, 0x4d4d294d,
  78     0x92927292, 0x7575c975, 0x06061806, 0x8a8a128a,
  79     0xb2b2f2b2, 0xe6e6bfe6, 0x0e0e380e, 0x1f1f7c1f,
  80     0x62629562, 0xd4d477d4, 0xa8a89aa8, 0x96966296,
  81     0xf9f9c3f9, 0xc5c533c5, 0x25259425, 0x59597959,
  82     0x84842a84, 0x7272d572, 0x3939e439, 0x4c4c2d4c,
  83     0x5e5e655e, 0x7878fd78, 0x3838e038, 0x8c8c0a8c,
  84     0xd1d163d1, 0xa5a5aea5, 0xe2e2afe2, 0x61619961,
  85     0xb3b3f6b3, 0x21218421, 0x9c9c4a9c, 0x1e1e781e,
  86     0x43431143, 0xc7c73bc7, 0xfcfcd7fc, 0x04041004,
  87     0x51515951, 0x99995e99, 0x6d6da96d, 0x0d0d340d,
  88     0xfafacffa, 0xdfdf5bdf, 0x7e7ee57e, 0x24249024,
  89     0x3b3bec3b, 0xabab96ab, 0xcece1fce, 0x11114411,
  90     0x8f8f068f, 0x4e4e254e, 0xb7b7e6b7, 0xebeb8beb,
  91     0x3c3cf03c, 0x81813e81, 0x94946a94, 0xf7f7fbf7,
  92     0xb9b9deb9, 0x13134c13, 0x2c2cb02c, 0xd3d36bd3,
  93     0xe7e7bbe7, 0x6e6ea56e, 0xc4c437c4, 0x03030c03,
  94     0x56564556, 0x44440d44, 0x7f7fe17f, 0xa9a99ea9,
  95     0x2a2aa82a, 0xbbbbd6bb, 0xc1c123c1, 0x53535153,
  96     0xdcdc57dc, 0x0b0b2c0b, 0x9d9d4e9d, 0x6c6cad6c,
  97     0x3131c431, 0x7474cd74, 0xf6f6fff6, 0x46460546,
  98     0xacac8aac, 0x89891e89, 0x14145014, 0xe1e1a3e1,
  99     0x16165816, 0x3a3ae83a, 0x6969b969, 0x09092409,
 100     0x7070dd70, 0xb6b6e2b6, 0xd0d067d0, 0xeded93ed,
 101     0xcccc17cc, 0x42421542, 0x98985a98, 0xa4a4aaa4,
 102     0x2828a028, 0x5c5c6d5c, 0xf8f8c7f8, 0x86862286,
 103   },
 104   {
 105     0xd8181860, 0x2623238c, 0xb8c6c63f, 0xfbe8e887,
 106     0xcb878726, 0x11b8b8da, 0x09010104, 0x0d4f4f21,
 107     0x9b3636d8, 0xffa6a6a2, 0x0cd2d26f, 0x0ef5f5f3,
 108     0x967979f9, 0x306f6fa1, 0x6d91917e, 0xf8525255,
 109     0x4760609d, 0x35bcbcca, 0x379b9b56, 0x8a8e8e02,
 110     0xd2a3a3b6, 0x6c0c0c30, 0x847b7bf1, 0x803535d4,
 111     0xf51d1d74, 0xb3e0e0a7, 0x21d7d77b, 0x9cc2c22f,
 112     0x432e2eb8, 0x294b4b31, 0x5dfefedf, 0xd5575741,
 113     0xbd151554, 0xe87777c1, 0x923737dc, 0x9ee5e5b3,
 114     0x139f9f46, 0x23f0f0e7, 0x204a4a35, 0x44dada4f,
 115     0xa258587d, 0xcfc9c903, 0x7c2929a4, 0x5a0a0a28,
 116     0x50b1b1fe, 0xc9a0a0ba, 0x146b6bb1, 0xd985852e,
 117     0x3cbdbdce, 0x8f5d5d69, 0x90101040, 0x07f4f4f7,
 118     0xddcbcb0b, 0xd33e3ef8, 0x2d050514, 0x78676781,
 119     0x97e4e4b7, 0x0227279c, 0x73414119, 0xa78b8b16,
 120     0xf6a7a7a6, 0xb27d7de9, 0x4995956e, 0x56d8d847,
 121     0x70fbfbcb, 0xcdeeee9f, 0xbb7c7ced, 0x71666685,
 122     0x7bdddd53, 0xaf17175c, 0x45474701, 0x1a9e9e42,
 123     0xd4caca0f, 0x582d2db4, 0x2ebfbfc6, 0x3f07071c,
 124     0xacadad8e, 0xb05a5a75, 0xef838336, 0xb63333cc,
 125     0x5c636391, 0x12020208, 0x93aaaa92, 0xde7171d9,
 126     0xc6c8c807, 0xd1191964, 0x3b494939, 0x5fd9d943,
 127     0x31f2f2ef, 0xa8e3e3ab, 0xb95b5b71, 0xbc88881a,
 128     0x3e9a9a52, 0x0b262698, 0xbf3232c8, 0x59b0b0fa,
 129     0xf2e9e983, 0x770f0f3c, 0x33d5d573, 0xf480803a,
 130     0x27bebec2, 0xebcdcd13, 0x893434d0, 0x3248483d,
 131     0x54ffffdb, 0x8d7a7af5, 0x6490907a, 0x9d5f5f61,
 132     0x3d202080, 0x0f6868bd, 0xca1a1a68, 0xb7aeae82,
 133     0x7db4b4ea, 0xce54544d, 0x7f939376, 0x2f222288,
 134     0x6364648d, 0x2af1f1e3, 0xcc7373d1, 0x82121248,
 135     0x7a40401d, 0x48080820, 0x95c3c32b, 0xdfecec97,
 136     0x4ddbdb4b, 0xc0a1a1be, 0x918d8d0e, 0xc83d3df4,
 137     0x5b979766, 0x00000000, 0xf9cfcf1b, 0x6e2b2bac,
 138     0xe17676c5, 0xe6828232, 0x28d6d67f, 0xc31b1b6c,
 139     0x74b5b5ee, 0xbeafaf86, 0x1d6a6ab5, 0xea50505d,
 140     0x57454509, 0x38f3f3eb, 0xad3030c0, 0xc4efef9b,
 141     0xda3f3ffc, 0xc7555549, 0xdba2a2b2, 0xe9eaea8f,
 142     0x6a656589, 0x03babad2, 0x4a2f2fbc, 0x8ec0c027,
 143     0x60dede5f, 0xfc1c1c70, 0x46fdfdd3, 0x1f4d4d29,
 144     0x76929272, 0xfa7575c9, 0x36060618, 0xae8a8a12,
 145     0x4bb2b2f2, 0x85e6e6bf, 0x7e0e0e38, 0xe71f1f7c,
 146     0x55626295, 0x3ad4d477, 0x81a8a89a, 0x52969662,
 147     0x62f9f9c3, 0xa3c5c533, 0x10252594, 0xab595979,
 148     0xd084842a, 0xc57272d5, 0xec3939e4, 0x164c4c2d,
 149     0x945e5e65, 0x9f7878fd, 0xe53838e0, 0x988c8c0a,
 150     0x17d1d163, 0xe4a5a5ae, 0xa1e2e2af, 0x4e616199,
 151     0x42b3b3f6, 0x34212184, 0x089c9c4a, 0xee1e1e78,
 152     0x61434311, 0xb1c7c73b, 0x4ffcfcd7, 0x24040410,
 153     0xe3515159, 0x2599995e, 0x226d6da9, 0x650d0d34,
 154     0x79fafacf, 0x69dfdf5b, 0xa97e7ee5, 0x19242490,
 155     0xfe3b3bec, 0x9aabab96, 0xf0cece1f, 0x99111144,
 156     0x838f8f06, 0x044e4e25, 0x66b7b7e6, 0xe0ebeb8b,
 157     0xc13c3cf0, 0xfd81813e, 0x4094946a, 0x1cf7f7fb,
 158     0x18b9b9de, 0x8b13134c, 0x512c2cb0, 0x05d3d36b,
 159     0x8ce7e7bb, 0x396e6ea5, 0xaac4c437, 0x1b03030c,
 160     0xdc565645, 0x5e44440d, 0xa07f7fe1, 0x88a9a99e,
 161     0x672a2aa8, 0x0abbbbd6, 0x87c1c123, 0xf1535351,
 162     0x72dcdc57, 0x530b0b2c, 0x019d9d4e, 0x2b6c6cad,
 163     0xa43131c4, 0xf37474cd, 0x15f6f6ff, 0x4c464605,
 164     0xa5acac8a, 0xb589891e, 0xb4141450, 0xbae1e1a3,
 165     0xa6161658, 0xf73a3ae8, 0x066969b9, 0x41090924,
 166     0xd77070dd, 0x6fb6b6e2, 0x1ed0d067, 0xd6eded93,
 167     0xe2cccc17, 0x68424215, 0x2c98985a, 0xeda4a4aa,
 168     0x752828a0, 0x865c5c6d, 0x6bf8f8c7, 0xc2868622,
 169   },
 170   {
 171     0x30d81818, 0x46262323, 0x91b8c6c6, 0xcdfbe8e8,
 172     0x13cb8787, 0x6d11b8b8, 0x02090101, 0x9e0d4f4f,
 173     0x6c9b3636, 0x51ffa6a6, 0xb90cd2d2, 0xf70ef5f5,
 174     0xf2967979, 0xde306f6f, 0x3f6d9191, 0xa4f85252,
 175     0xc0476060, 0x6535bcbc, 0x2b379b9b, 0x018a8e8e,
 176     0x5bd2a3a3, 0x186c0c0c, 0xf6847b7b, 0x6a803535,
 177     0x3af51d1d, 0xddb3e0e0, 0xb321d7d7, 0x999cc2c2,
 178     0x5c432e2e, 0x96294b4b, 0xe15dfefe, 0xaed55757,
 179     0x2abd1515, 0xeee87777, 0x6e923737, 0xd79ee5e5,
 180     0x23139f9f, 0xfd23f0f0, 0x94204a4a, 0xa944dada,
 181     0xb0a25858, 0x8fcfc9c9, 0x527c2929, 0x145a0a0a,
 182     0x7f50b1b1, 0x5dc9a0a0, 0xd6146b6b, 0x17d98585,
 183     0x673cbdbd, 0xba8f5d5d, 0x20901010, 0xf507f4f4,
 184     0x8bddcbcb, 0x7cd33e3e, 0x0a2d0505, 0xce786767,
 185     0xd597e4e4, 0x4e022727, 0x82734141, 0x0ba78b8b,
 186     0x53f6a7a7, 0xfab27d7d, 0x37499595, 0xad56d8d8,
 187     0xeb70fbfb, 0xc1cdeeee, 0xf8bb7c7c, 0xcc716666,
 188     0xa77bdddd, 0x2eaf1717, 0x8e454747, 0x211a9e9e,
 189     0x89d4caca, 0x5a582d2d, 0x632ebfbf, 0x0e3f0707,
 190     0x47acadad, 0xb4b05a5a, 0x1bef8383, 0x66b63333,
 191     0xc65c6363, 0x04120202, 0x4993aaaa, 0xe2de7171,
 192     0x8dc6c8c8, 0x32d11919, 0x923b4949, 0xaf5fd9d9,
 193     0xf931f2f2, 0xdba8e3e3, 0xb6b95b5b, 0x0dbc8888,
 194     0x293e9a9a, 0x4c0b2626, 0x64bf3232, 0x7d59b0b0,
 195     0xcff2e9e9, 0x1e770f0f, 0xb733d5d5, 0x1df48080,
 196     0x6127bebe, 0x87ebcdcd, 0x68893434, 0x90324848,
 197     0xe354ffff, 0xf48d7a7a, 0x3d649090, 0xbe9d5f5f,
 198     0x403d2020, 0xd00f6868, 0x34ca1a1a, 0x41b7aeae,
 199     0x757db4b4, 0xa8ce5454, 0x3b7f9393, 0x442f2222,
 200     0xc8636464, 0xff2af1f1, 0xe6cc7373, 0x24821212,
 201     0x807a4040, 0x10480808, 0x9b95c3c3, 0xc5dfecec,
 202     0xab4ddbdb, 0x5fc0a1a1, 0x07918d8d, 0x7ac83d3d,
 203     0x335b9797, 0x00000000, 0x83f9cfcf, 0x566e2b2b,
 204     0xece17676, 0x19e68282, 0xb128d6d6, 0x36c31b1b,
 205     0x7774b5b5, 0x43beafaf, 0xd41d6a6a, 0xa0ea5050,
 206     0x8a574545, 0xfb38f3f3, 0x60ad3030, 0xc3c4efef,
 207     0x7eda3f3f, 0xaac75555, 0x59dba2a2, 0xc9e9eaea,
 208     0xca6a6565, 0x6903baba, 0x5e4a2f2f, 0x9d8ec0c0,
 209     0xa160dede, 0x38fc1c1c, 0xe746fdfd, 0x9a1f4d4d,
 210     0x39769292, 0xeafa7575, 0x0c360606, 0x09ae8a8a,
 211     0x794bb2b2, 0xd185e6e6, 0x1c7e0e0e, 0x3ee71f1f,
 212     0xc4556262, 0xb53ad4d4, 0x4d81a8a8, 0x31529696,
 213     0xef62f9f9, 0x97a3c5c5, 0x4a102525, 0xb2ab5959,
 214     0x15d08484, 0xe4c57272, 0x72ec3939, 0x98164c4c,
 215     0xbc945e5e, 0xf09f7878, 0x70e53838, 0x05988c8c,
 216     0xbf17d1d1, 0x57e4a5a5, 0xd9a1e2e2, 0xc24e6161,
 217     0x7b42b3b3, 0x42342121, 0x25089c9c, 0x3cee1e1e,
 218     0x86614343, 0x93b1c7c7, 0xe54ffcfc, 0x08240404,
 219     0xa2e35151, 0x2f259999, 0xda226d6d, 0x1a650d0d,
 220     0xe979fafa, 0xa369dfdf, 0xfca97e7e, 0x48192424,
 221     0x76fe3b3b, 0x4b9aabab, 0x81f0cece, 0x22991111,
 222     0x03838f8f, 0x9c044e4e, 0x7366b7b7, 0xcbe0ebeb,
 223     0x78c13c3c, 0x1ffd8181, 0x35409494, 0xf31cf7f7,
 224     0x6f18b9b9, 0x268b1313, 0x58512c2c, 0xbb05d3d3,
 225     0xd38ce7e7, 0xdc396e6e, 0x95aac4c4, 0x061b0303,
 226     0xacdc5656, 0x885e4444, 0xfea07f7f, 0x4f88a9a9,
 227     0x54672a2a, 0x6b0abbbb, 0x9f87c1c1, 0xa6f15353,
 228     0xa572dcdc, 0x16530b0b, 0x27019d9d, 0xd82b6c6c,
 229     0x62a43131, 0xe8f37474, 0xf115f6f6, 0x8c4c4646,
 230     0x45a5acac, 0x0fb58989, 0x28b41414, 0xdfbae1e1,
 231     0x2ca61616, 0x74f73a3a, 0xd2066969, 0x12410909,
 232     0xe0d77070, 0x716fb6b6, 0xbd1ed0d0, 0xc7d6eded,
 233     0x85e2cccc, 0x84684242, 0x2d2c9898, 0x55eda4a4,
 234     0x50752828, 0xb8865c5c, 0xed6bf8f8, 0x11c28686,
 235   },
 236   {
 237     0x7830d818, 0xaf462623, 0xf991b8c6, 0x6fcdfbe8,
 238     0xa113cb87, 0x626d11b8, 0x05020901, 0x6e9e0d4f,
 239     0xee6c9b36, 0x0451ffa6, 0xbdb90cd2, 0x06f70ef5,
 240     0x80f29679, 0xcede306f, 0xef3f6d91, 0x07a4f852,
 241     0xfdc04760, 0x766535bc, 0xcd2b379b, 0x8c018a8e,
 242     0x155bd2a3, 0x3c186c0c, 0x8af6847b, 0xe16a8035,
 243     0x693af51d, 0x47ddb3e0, 0xacb321d7, 0xed999cc2,
 244     0x965c432e, 0x7a96294b, 0x21e15dfe, 0x16aed557,
 245     0x412abd15, 0xb6eee877, 0xeb6e9237, 0x56d79ee5,
 246     0xd923139f, 0x17fd23f0, 0x7f94204a, 0x95a944da,
 247     0x25b0a258, 0xca8fcfc9, 0x8d527c29, 0x22145a0a,
 248     0x4f7f50b1, 0x1a5dc9a0, 0xdad6146b, 0xab17d985,
 249     0x73673cbd, 0x34ba8f5d, 0x50209010, 0x03f507f4,
 250     0xc08bddcb, 0xc67cd33e, 0x110a2d05, 0xe6ce7867,
 251     0x53d597e4, 0xbb4e0227, 0x58827341, 0x9d0ba78b,
 252     0x0153f6a7, 0x94fab27d, 0xfb374995, 0x9fad56d8,
 253     0x30eb70fb, 0x71c1cdee, 0x91f8bb7c, 0xe3cc7166,
 254     0x8ea77bdd, 0x4b2eaf17, 0x468e4547, 0xdc211a9e,
 255     0xc589d4ca, 0x995a582d, 0x79632ebf, 0x1b0e3f07,
 256     0x2347acad, 0x2fb4b05a, 0xb51bef83, 0xff66b633,
 257     0xf2c65c63, 0x0a041202, 0x384993aa, 0xa8e2de71,
 258     0xcf8dc6c8, 0x7d32d119, 0x70923b49, 0x9aaf5fd9,
 259     0x1df931f2, 0x48dba8e3, 0x2ab6b95b, 0x920dbc88,
 260     0xc8293e9a, 0xbe4c0b26, 0xfa64bf32, 0x4a7d59b0,
 261     0x6acff2e9, 0x331e770f, 0xa6b733d5, 0xba1df480,
 262     0x7c6127be, 0xde87ebcd, 0xe4688934, 0x75903248,
 263     0x24e354ff, 0x8ff48d7a, 0xea3d6490, 0x3ebe9d5f,
 264     0xa0403d20, 0xd5d00f68, 0x7234ca1a, 0x2c41b7ae,
 265     0x5e757db4, 0x19a8ce54, 0xe53b7f93, 0xaa442f22,
 266     0xe9c86364, 0x12ff2af1, 0xa2e6cc73, 0x5a248212,
 267     0x5d807a40, 0x28104808, 0xe89b95c3, 0x7bc5dfec,
 268     0x90ab4ddb, 0x1f5fc0a1, 0x8307918d, 0xc97ac83d,
 269     0xf1335b97, 0x00000000, 0xd483f9cf, 0x87566e2b,
 270     0xb3ece176, 0xb019e682, 0xa9b128d6, 0x7736c31b,
 271     0x5b7774b5, 0x2943beaf, 0xdfd41d6a, 0x0da0ea50,
 272     0x4c8a5745, 0x18fb38f3, 0xf060ad30, 0x74c3c4ef,
 273     0xc37eda3f, 0x1caac755, 0x1059dba2, 0x65c9e9ea,
 274     0xecca6a65, 0x686903ba, 0x935e4a2f, 0xe79d8ec0,
 275     0x81a160de, 0x6c38fc1c, 0x2ee746fd, 0x649a1f4d,
 276     0xe0397692, 0xbceafa75, 0x1e0c3606, 0x9809ae8a,
 277     0x40794bb2, 0x59d185e6, 0x361c7e0e, 0x633ee71f,
 278     0xf7c45562, 0xa3b53ad4, 0x324d81a8, 0xf4315296,
 279     0x3aef62f9, 0xf697a3c5, 0xb14a1025, 0x20b2ab59,
 280     0xae15d084, 0xa7e4c572, 0xdd72ec39, 0x6198164c,
 281     0x3bbc945e, 0x85f09f78, 0xd870e538, 0x8605988c,
 282     0xb2bf17d1, 0x0b57e4a5, 0x4dd9a1e2, 0xf8c24e61,
 283     0x457b42b3, 0xa5423421, 0xd625089c, 0x663cee1e,
 284     0x52866143, 0xfc93b1c7, 0x2be54ffc, 0x14082404,
 285     0x08a2e351, 0xc72f2599, 0xc4da226d, 0x391a650d,
 286     0x35e979fa, 0x84a369df, 0x9bfca97e, 0xb4481924,
 287     0xd776fe3b, 0x3d4b9aab, 0xd181f0ce, 0x55229911,
 288     0x8903838f, 0x6b9c044e, 0x517366b7, 0x60cbe0eb,
 289     0xcc78c13c, 0xbf1ffd81, 0xfe354094, 0x0cf31cf7,
 290     0x676f18b9, 0x5f268b13, 0x9c58512c, 0xb8bb05d3,
 291     0x5cd38ce7, 0xcbdc396e, 0xf395aac4, 0x0f061b03,
 292     0x13acdc56, 0x49885e44, 0x9efea07f, 0x374f88a9,
 293     0x8254672a, 0x6d6b0abb, 0xe29f87c1, 0x02a6f153,
 294     0x8ba572dc, 0x2716530b, 0xd327019d, 0xc1d82b6c,
 295     0xf562a431, 0xb9e8f374, 0x09f115f6, 0x438c4c46,
 296     0x2645a5ac, 0x970fb589, 0x4428b414, 0x42dfbae1,
 297     0x4e2ca616, 0xd274f73a, 0xd0d20669, 0x2d124109,
 298     0xade0d770, 0x54716fb6, 0xb7bd1ed0, 0x7ec7d6ed,
 299     0xdb85e2cc, 0x57846842, 0xc22d2c98, 0x0e55eda4,
 300     0x88507528, 0x31b8865c, 0x3fed6bf8, 0xa411c286,
 301   },
 302   {
 303     0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
 304     0x4ca113cb, 0xa9626d11, 0x08050209, 0x426e9e0d,
 305     0xadee6c9b, 0x590451ff, 0xdebdb90c, 0xfb06f70e,
 306     0xef80f296, 0x5fcede30, 0xfcef3f6d, 0xaa07a4f8,
 307     0x27fdc047, 0x89766535, 0xaccd2b37, 0x048c018a,
 308     0x71155bd2, 0x603c186c, 0xff8af684, 0xb5e16a80,
 309     0xe8693af5, 0x5347ddb3, 0xf6acb321, 0x5eed999c,
 310     0x6d965c43, 0x627a9629, 0xa321e15d, 0x8216aed5,
 311     0xa8412abd, 0x9fb6eee8, 0xa5eb6e92, 0x7b56d79e,
 312     0x8cd92313, 0xd317fd23, 0x6a7f9420, 0x9e95a944,
 313     0xfa25b0a2, 0x06ca8fcf, 0x558d527c, 0x5022145a,
 314     0xe14f7f50, 0x691a5dc9, 0x7fdad614, 0x5cab17d9,
 315     0x8173673c, 0xd234ba8f, 0x80502090, 0xf303f507,
 316     0x16c08bdd, 0xedc67cd3, 0x28110a2d, 0x1fe6ce78,
 317     0x7353d597, 0x25bb4e02, 0x32588273, 0x2c9d0ba7,
 318     0x510153f6, 0xcf94fab2, 0xdcfb3749, 0x8e9fad56,
 319     0x8b30eb70, 0x2371c1cd, 0xc791f8bb, 0x17e3cc71,
 320     0xa68ea77b, 0xb84b2eaf, 0x02468e45, 0x84dc211a,
 321     0x1ec589d4, 0x75995a58, 0x9179632e, 0x381b0e3f,
 322     0x012347ac, 0xea2fb4b0, 0x6cb51bef, 0x85ff66b6,
 323     0x3ff2c65c, 0x100a0412, 0x39384993, 0xafa8e2de,
 324     0x0ecf8dc6, 0xc87d32d1, 0x7270923b, 0x869aaf5f,
 325     0xc31df931, 0x4b48dba8, 0xe22ab6b9, 0x34920dbc,
 326     0xa4c8293e, 0x2dbe4c0b, 0x8dfa64bf, 0xe94a7d59,
 327     0x1b6acff2, 0x78331e77, 0xe6a6b733, 0x74ba1df4,
 328     0x997c6127, 0x26de87eb, 0xbde46889, 0x7a759032,
 329     0xab24e354, 0xf78ff48d, 0xf4ea3d64, 0xc23ebe9d,
 330     0x1da0403d, 0x67d5d00f, 0xd07234ca, 0x192c41b7,
 331     0xc95e757d, 0x9a19a8ce, 0xece53b7f, 0x0daa442f,
 332     0x07e9c863, 0xdb12ff2a, 0xbfa2e6cc, 0x905a2482,
 333     0x3a5d807a, 0x40281048, 0x56e89b95, 0x337bc5df,
 334     0x9690ab4d, 0x611f5fc0, 0x1c830791, 0xf5c97ac8,
 335     0xccf1335b, 0x00000000, 0x36d483f9, 0x4587566e,
 336     0x97b3ece1, 0x64b019e6, 0xfea9b128, 0xd87736c3,
 337     0xc15b7774, 0x112943be, 0x77dfd41d, 0xba0da0ea,
 338     0x124c8a57, 0xcb18fb38, 0x9df060ad, 0x2b74c3c4,
 339     0xe5c37eda, 0x921caac7, 0x791059db, 0x0365c9e9,
 340     0x0fecca6a, 0xb9686903, 0x65935e4a, 0x4ee79d8e,
 341     0xbe81a160, 0xe06c38fc, 0xbb2ee746, 0x52649a1f,
 342     0xe4e03976, 0x8fbceafa, 0x301e0c36, 0x249809ae,
 343     0xf940794b, 0x6359d185, 0x70361c7e, 0xf8633ee7,
 344     0x37f7c455, 0xeea3b53a, 0x29324d81, 0xc4f43152,
 345     0x9b3aef62, 0x66f697a3, 0x35b14a10, 0xf220b2ab,
 346     0x54ae15d0, 0xb7a7e4c5, 0xd5dd72ec, 0x5a619816,
 347     0xca3bbc94, 0xe785f09f, 0xddd870e5, 0x14860598,
 348     0xc6b2bf17, 0x410b57e4, 0x434dd9a1, 0x2ff8c24e,
 349     0xf1457b42, 0x15a54234, 0x94d62508, 0xf0663cee,
 350     0x22528661, 0x76fc93b1, 0xb32be54f, 0x20140824,
 351     0xb208a2e3, 0xbcc72f25, 0x4fc4da22, 0x68391a65,
 352     0x8335e979, 0xb684a369, 0xd79bfca9, 0x3db44819,
 353     0xc5d776fe, 0x313d4b9a, 0x3ed181f0, 0x88552299,
 354     0x0c890383, 0x4a6b9c04, 0xd1517366, 0x0b60cbe0,
 355     0xfdcc78c1, 0x7cbf1ffd, 0xd4fe3540, 0xeb0cf31c,
 356     0xa1676f18, 0x985f268b, 0x7d9c5851, 0xd6b8bb05,
 357     0x6b5cd38c, 0x57cbdc39, 0x6ef395aa, 0x180f061b,
 358     0x8a13acdc, 0x1a49885e, 0xdf9efea0, 0x21374f88,
 359     0x4d825467, 0xb16d6b0a, 0x46e29f87, 0xa202a6f1,
 360     0xae8ba572, 0x58271653, 0x9cd32701, 0x47c1d82b,
 361     0x95f562a4, 0x87b9e8f3, 0xe309f115, 0x0a438c4c,
 362     0x092645a5, 0x3c970fb5, 0xa04428b4, 0x5b42dfba,
 363     0xb04e2ca6, 0xcdd274f7, 0x6fd0d206, 0x482d1241,
 364     0xa7ade0d7, 0xd954716f, 0xceb7bd1e, 0x3b7ec7d6,
 365     0x2edb85e2, 0x2a578468, 0xb4c22d2c, 0x490e55ed,
 366     0x5d885075, 0xda31b886, 0x933fed6b, 0x44a411c2,
 367   },
 368   {
 369     0x18c07830, 0x2305af46, 0xc67ef991, 0xe8136fcd,
 370     0x874ca113, 0xb8a9626d, 0x01080502, 0x4f426e9e,
 371     0x36adee6c, 0xa6590451, 0xd2debdb9, 0xf5fb06f7,
 372     0x79ef80f2, 0x6f5fcede, 0x91fcef3f, 0x52aa07a4,
 373     0x6027fdc0, 0xbc897665, 0x9baccd2b, 0x8e048c01,
 374     0xa371155b, 0x0c603c18, 0x7bff8af6, 0x35b5e16a,
 375     0x1de8693a, 0xe05347dd, 0xd7f6acb3, 0xc25eed99,
 376     0x2e6d965c, 0x4b627a96, 0xfea321e1, 0x578216ae,
 377     0x15a8412a, 0x779fb6ee, 0x37a5eb6e, 0xe57b56d7,
 378     0x9f8cd923, 0xf0d317fd, 0x4a6a7f94, 0xda9e95a9,
 379     0x58fa25b0, 0xc906ca8f, 0x29558d52, 0x0a502214,
 380     0xb1e14f7f, 0xa0691a5d, 0x6b7fdad6, 0x855cab17,
 381     0xbd817367, 0x5dd234ba, 0x10805020, 0xf4f303f5,
 382     0xcb16c08b, 0x3eedc67c, 0x0528110a, 0x671fe6ce,
 383     0xe47353d5, 0x2725bb4e, 0x41325882, 0x8b2c9d0b,
 384     0xa7510153, 0x7dcf94fa, 0x95dcfb37, 0xd88e9fad,
 385     0xfb8b30eb, 0xee2371c1, 0x7cc791f8, 0x6617e3cc,
 386     0xdda68ea7, 0x17b84b2e, 0x4702468e, 0x9e84dc21,
 387     0xca1ec589, 0x2d75995a, 0xbf917963, 0x07381b0e,
 388     0xad012347, 0x5aea2fb4, 0x836cb51b, 0x3385ff66,
 389     0x633ff2c6, 0x02100a04, 0xaa393849, 0x71afa8e2,
 390     0xc80ecf8d, 0x19c87d32, 0x49727092, 0xd9869aaf,
 391     0xf2c31df9, 0xe34b48db, 0x5be22ab6, 0x8834920d,
 392     0x9aa4c829, 0x262dbe4c, 0x328dfa64, 0xb0e94a7d,
 393     0xe91b6acf, 0x0f78331e, 0xd5e6a6b7, 0x8074ba1d,
 394     0xbe997c61, 0xcd26de87, 0x34bde468, 0x487a7590,
 395     0xffab24e3, 0x7af78ff4, 0x90f4ea3d, 0x5fc23ebe,
 396     0x201da040, 0x6867d5d0, 0x1ad07234, 0xae192c41,
 397     0xb4c95e75, 0x549a19a8, 0x93ece53b, 0x220daa44,
 398     0x6407e9c8, 0xf1db12ff, 0x73bfa2e6, 0x12905a24,
 399     0x403a5d80, 0x08402810, 0xc356e89b, 0xec337bc5,
 400     0xdb9690ab, 0xa1611f5f, 0x8d1c8307, 0x3df5c97a,
 401     0x97ccf133, 0x00000000, 0xcf36d483, 0x2b458756,
 402     0x7697b3ec, 0x8264b019, 0xd6fea9b1, 0x1bd87736,
 403     0xb5c15b77, 0xaf112943, 0x6a77dfd4, 0x50ba0da0,
 404     0x45124c8a, 0xf3cb18fb, 0x309df060, 0xef2b74c3,
 405     0x3fe5c37e, 0x55921caa, 0xa2791059, 0xea0365c9,
 406     0x650fecca, 0xbab96869, 0x2f65935e, 0xc04ee79d,
 407     0xdebe81a1, 0x1ce06c38, 0xfdbb2ee7, 0x4d52649a,
 408     0x92e4e039, 0x758fbcea, 0x06301e0c, 0x8a249809,
 409     0xb2f94079, 0xe66359d1, 0x0e70361c, 0x1ff8633e,
 410     0x6237f7c4, 0xd4eea3b5, 0xa829324d, 0x96c4f431,
 411     0xf99b3aef, 0xc566f697, 0x2535b14a, 0x59f220b2,
 412     0x8454ae15, 0x72b7a7e4, 0x39d5dd72, 0x4c5a6198,
 413     0x5eca3bbc, 0x78e785f0, 0x38ddd870, 0x8c148605,
 414     0xd1c6b2bf, 0xa5410b57, 0xe2434dd9, 0x612ff8c2,
 415     0xb3f1457b, 0x2115a542, 0x9c94d625, 0x1ef0663c,
 416     0x43225286, 0xc776fc93, 0xfcb32be5, 0x04201408,
 417     0x51b208a2, 0x99bcc72f, 0x6d4fc4da, 0x0d68391a,
 418     0xfa8335e9, 0xdfb684a3, 0x7ed79bfc, 0x243db448,
 419     0x3bc5d776, 0xab313d4b, 0xce3ed181, 0x11885522,
 420     0x8f0c8903, 0x4e4a6b9c, 0xb7d15173, 0xeb0b60cb,
 421     0x3cfdcc78, 0x817cbf1f, 0x94d4fe35, 0xf7eb0cf3,
 422     0xb9a1676f, 0x13985f26, 0x2c7d9c58, 0xd3d6b8bb,
 423     0xe76b5cd3, 0x6e57cbdc, 0xc46ef395, 0x03180f06,
 424     0x568a13ac, 0x441a4988, 0x7fdf9efe, 0xa921374f,
 425     0x2a4d8254, 0xbbb16d6b, 0xc146e29f, 0x53a202a6,
 426     0xdcae8ba5, 0x0b582716, 0x9d9cd327, 0x6c47c1d8,
 427     0x3195f562, 0x7487b9e8, 0xf6e309f1, 0x460a438c,
 428     0xac092645, 0x893c970f, 0x14a04428, 0xe15b42df,
 429     0x16b04e2c, 0x3acdd274, 0x696fd0d2, 0x09482d12,
 430     0x70a7ade0, 0xb6d95471, 0xd0ceb7bd, 0xed3b7ec7,
 431     0xcc2edb85, 0x422a5784, 0x98b4c22d, 0xa4490e55,
 432     0x285d8850, 0x5cda31b8, 0xf8933fed, 0x8644a411,
 433   },
 434   {
 435     0x6018c078, 0x8c2305af, 0x3fc67ef9, 0x87e8136f,
 436     0x26874ca1, 0xdab8a962, 0x04010805, 0x214f426e,
 437     0xd836adee, 0xa2a65904, 0x6fd2debd, 0xf3f5fb06,
 438     0xf979ef80, 0xa16f5fce, 0x7e91fcef, 0x5552aa07,
 439     0x9d6027fd, 0xcabc8976, 0x569baccd, 0x028e048c,
 440     0xb6a37115, 0x300c603c, 0xf17bff8a, 0xd435b5e1,
 441     0x741de869, 0xa7e05347, 0x7bd7f6ac, 0x2fc25eed,
 442     0xb82e6d96, 0x314b627a, 0xdffea321, 0x41578216,
 443     0x5415a841, 0xc1779fb6, 0xdc37a5eb, 0xb3e57b56,
 444     0x469f8cd9, 0xe7f0d317, 0x354a6a7f, 0x4fda9e95,
 445     0x7d58fa25, 0x03c906ca, 0xa429558d, 0x280a5022,
 446     0xfeb1e14f, 0xbaa0691a, 0xb16b7fda, 0x2e855cab,
 447     0xcebd8173, 0x695dd234, 0x40108050, 0xf7f4f303,
 448     0x0bcb16c0, 0xf83eedc6, 0x14052811, 0x81671fe6,
 449     0xb7e47353, 0x9c2725bb, 0x19413258, 0x168b2c9d,
 450     0xa6a75101, 0xe97dcf94, 0x6e95dcfb, 0x47d88e9f,
 451     0xcbfb8b30, 0x9fee2371, 0xed7cc791, 0x856617e3,
 452     0x53dda68e, 0x5c17b84b, 0x01470246, 0x429e84dc,
 453     0x0fca1ec5, 0xb42d7599, 0xc6bf9179, 0x1c07381b,
 454     0x8ead0123, 0x755aea2f, 0x36836cb5, 0xcc3385ff,
 455     0x91633ff2, 0x0802100a, 0x92aa3938, 0xd971afa8,
 456     0x07c80ecf, 0x6419c87d, 0x39497270, 0x43d9869a,
 457     0xeff2c31d, 0xabe34b48, 0x715be22a, 0x1a883492,
 458     0x529aa4c8, 0x98262dbe, 0xc8328dfa, 0xfab0e94a,
 459     0x83e91b6a, 0x3c0f7833, 0x73d5e6a6, 0x3a8074ba,
 460     0xc2be997c, 0x13cd26de, 0xd034bde4, 0x3d487a75,
 461     0xdbffab24, 0xf57af78f, 0x7a90f4ea, 0x615fc23e,
 462     0x80201da0, 0xbd6867d5, 0x681ad072, 0x82ae192c,
 463     0xeab4c95e, 0x4d549a19, 0x7693ece5, 0x88220daa,
 464     0x8d6407e9, 0xe3f1db12, 0xd173bfa2, 0x4812905a,
 465     0x1d403a5d, 0x20084028, 0x2bc356e8, 0x97ec337b,
 466     0x4bdb9690, 0xbea1611f, 0x0e8d1c83, 0xf43df5c9,
 467     0x6697ccf1, 0x00000000, 0x1bcf36d4, 0xac2b4587,
 468     0xc57697b3, 0x328264b0, 0x7fd6fea9, 0x6c1bd877,
 469     0xeeb5c15b, 0x86af1129, 0xb56a77df, 0x5d50ba0d,
 470     0x0945124c, 0xebf3cb18, 0xc0309df0, 0x9bef2b74,
 471     0xfc3fe5c3, 0x4955921c, 0xb2a27910, 0x8fea0365,
 472     0x89650fec, 0xd2bab968, 0xbc2f6593, 0x27c04ee7,
 473     0x5fdebe81, 0x701ce06c, 0xd3fdbb2e, 0x294d5264,
 474     0x7292e4e0, 0xc9758fbc, 0x1806301e, 0x128a2498,
 475     0xf2b2f940, 0xbfe66359, 0x380e7036, 0x7c1ff863,
 476     0x956237f7, 0x77d4eea3, 0x9aa82932, 0x6296c4f4,
 477     0xc3f99b3a, 0x33c566f6, 0x942535b1, 0x7959f220,
 478     0x2a8454ae, 0xd572b7a7, 0xe439d5dd, 0x2d4c5a61,
 479     0x655eca3b, 0xfd78e785, 0xe038ddd8, 0x0a8c1486,
 480     0x63d1c6b2, 0xaea5410b, 0xafe2434d, 0x99612ff8,
 481     0xf6b3f145, 0x842115a5, 0x4a9c94d6, 0x781ef066,
 482     0x11432252, 0x3bc776fc, 0xd7fcb32b, 0x10042014,
 483     0x5951b208, 0x5e99bcc7, 0xa96d4fc4, 0x340d6839,
 484     0xcffa8335, 0x5bdfb684, 0xe57ed79b, 0x90243db4,
 485     0xec3bc5d7, 0x96ab313d, 0x1fce3ed1, 0x44118855,
 486     0x068f0c89, 0x254e4a6b, 0xe6b7d151, 0x8beb0b60,
 487     0xf03cfdcc, 0x3e817cbf, 0x6a94d4fe, 0xfbf7eb0c,
 488     0xdeb9a167, 0x4c13985f, 0xb02c7d9c, 0x6bd3d6b8,
 489     0xbbe76b5c, 0xa56e57cb, 0x37c46ef3, 0x0c03180f,
 490     0x45568a13, 0x0d441a49, 0xe17fdf9e, 0x9ea92137,
 491     0xa82a4d82, 0xd6bbb16d, 0x23c146e2, 0x5153a202,
 492     0x57dcae8b, 0x2c0b5827, 0x4e9d9cd3, 0xad6c47c1,
 493     0xc43195f5, 0xcd7487b9, 0xfff6e309, 0x05460a43,
 494     0x8aac0926, 0x1e893c97, 0x5014a044, 0xa3e15b42,
 495     0x5816b04e, 0xe83acdd2, 0xb9696fd0, 0x2409482d,
 496     0xdd70a7ad, 0xe2b6d954, 0x67d0ceb7, 0x93ed3b7e,
 497     0x17cc2edb, 0x15422a57, 0x5a98b4c2, 0xaaa4490e,
 498     0xa0285d88, 0x6d5cda31, 0xc7f8933f, 0x228644a4,
 499   },
 500   {
 501     0x186018c0, 0x238c2305, 0xc63fc67e, 0xe887e813,
 502     0x8726874c, 0xb8dab8a9, 0x01040108, 0x4f214f42,
 503     0x36d836ad, 0xa6a2a659, 0xd26fd2de, 0xf5f3f5fb,
 504     0x79f979ef, 0x6fa16f5f, 0x917e91fc, 0x525552aa,
 505     0x609d6027, 0xbccabc89, 0x9b569bac, 0x8e028e04,
 506     0xa3b6a371, 0x0c300c60, 0x7bf17bff, 0x35d435b5,
 507     0x1d741de8, 0xe0a7e053, 0xd77bd7f6, 0xc22fc25e,
 508     0x2eb82e6d, 0x4b314b62, 0xfedffea3, 0x57415782,
 509     0x155415a8, 0x77c1779f, 0x37dc37a5, 0xe5b3e57b,
 510     0x9f469f8c, 0xf0e7f0d3, 0x4a354a6a, 0xda4fda9e,
 511     0x587d58fa, 0xc903c906, 0x29a42955, 0x0a280a50,
 512     0xb1feb1e1, 0xa0baa069, 0x6bb16b7f, 0x852e855c,
 513     0xbdcebd81, 0x5d695dd2, 0x10401080, 0xf4f7f4f3,
 514     0xcb0bcb16, 0x3ef83eed, 0x05140528, 0x6781671f,
 515     0xe4b7e473, 0x279c2725, 0x41194132, 0x8b168b2c,
 516     0xa7a6a751, 0x7de97dcf, 0x956e95dc, 0xd847d88e,
 517     0xfbcbfb8b, 0xee9fee23, 0x7ced7cc7, 0x66856617,
 518     0xdd53dda6, 0x175c17b8, 0x47014702, 0x9e429e84,
 519     0xca0fca1e, 0x2db42d75, 0xbfc6bf91, 0x071c0738,
 520     0xad8ead01, 0x5a755aea, 0x8336836c, 0x33cc3385,
 521     0x6391633f, 0x02080210, 0xaa92aa39, 0x71d971af,
 522     0xc807c80e, 0x196419c8, 0x49394972, 0xd943d986,
 523     0xf2eff2c3, 0xe3abe34b, 0x5b715be2, 0x881a8834,
 524     0x9a529aa4, 0x2698262d, 0x32c8328d, 0xb0fab0e9,
 525     0xe983e91b, 0x0f3c0f78, 0xd573d5e6, 0x803a8074,
 526     0xbec2be99, 0xcd13cd26, 0x34d034bd, 0x483d487a,
 527     0xffdbffab, 0x7af57af7, 0x907a90f4, 0x5f615fc2,
 528     0x2080201d, 0x68bd6867, 0x1a681ad0, 0xae82ae19,
 529     0xb4eab4c9, 0x544d549a, 0x937693ec, 0x2288220d,
 530     0x648d6407, 0xf1e3f1db, 0x73d173bf, 0x12481290,
 531     0x401d403a, 0x08200840, 0xc32bc356, 0xec97ec33,
 532     0xdb4bdb96, 0xa1bea161, 0x8d0e8d1c, 0x3df43df5,
 533     0x976697cc, 0x00000000, 0xcf1bcf36, 0x2bac2b45,
 534     0x76c57697, 0x82328264, 0xd67fd6fe, 0x1b6c1bd8,
 535     0xb5eeb5c1, 0xaf86af11, 0x6ab56a77, 0x505d50ba,
 536     0x45094512, 0xf3ebf3cb, 0x30c0309d, 0xef9bef2b,
 537     0x3ffc3fe5, 0x55495592, 0xa2b2a279, 0xea8fea03,
 538     0x6589650f, 0xbad2bab9, 0x2fbc2f65, 0xc027c04e,
 539     0xde5fdebe, 0x1c701ce0, 0xfdd3fdbb, 0x4d294d52,
 540     0x927292e4, 0x75c9758f, 0x06180630, 0x8a128a24,
 541     0xb2f2b2f9, 0xe6bfe663, 0x0e380e70, 0x1f7c1ff8,
 542     0x62956237, 0xd477d4ee, 0xa89aa829, 0x966296c4,
 543     0xf9c3f99b, 0xc533c566, 0x25942535, 0x597959f2,
 544     0x842a8454, 0x72d572b7, 0x39e439d5, 0x4c2d4c5a,
 545     0x5e655eca, 0x78fd78e7, 0x38e038dd, 0x8c0a8c14,
 546     0xd163d1c6, 0xa5aea541, 0xe2afe243, 0x6199612f,
 547     0xb3f6b3f1, 0x21842115, 0x9c4a9c94, 0x1e781ef0,
 548     0x43114322, 0xc73bc776, 0xfcd7fcb3, 0x04100420,
 549     0x515951b2, 0x995e99bc, 0x6da96d4f, 0x0d340d68,
 550     0xfacffa83, 0xdf5bdfb6, 0x7ee57ed7, 0x2490243d,
 551     0x3bec3bc5, 0xab96ab31, 0xce1fce3e, 0x11441188,
 552     0x8f068f0c, 0x4e254e4a, 0xb7e6b7d1, 0xeb8beb0b,
 553     0x3cf03cfd, 0x813e817c, 0x946a94d4, 0xf7fbf7eb,
 554     0xb9deb9a1, 0x134c1398, 0x2cb02c7d, 0xd36bd3d6,
 555     0xe7bbe76b, 0x6ea56e57, 0xc437c46e, 0x030c0318,
 556     0x5645568a, 0x440d441a, 0x7fe17fdf, 0xa99ea921,
 557     0x2aa82a4d, 0xbbd6bbb1, 0xc123c146, 0x535153a2,
 558     0xdc57dcae, 0x0b2c0b58, 0x9d4e9d9c, 0x6cad6c47,
 559     0x31c43195, 0x74cd7487, 0xf6fff6e3, 0x4605460a,
 560     0xac8aac09, 0x891e893c, 0x145014a0, 0xe1a3e15b,
 561     0x165816b0, 0x3ae83acd, 0x69b9696f, 0x09240948,
 562     0x70dd70a7, 0xb6e2b6d9, 0xd067d0ce, 0xed93ed3b,
 563     0xcc17cc2e, 0x4215422a, 0x985a98b4, 0xa4aaa449,
 564     0x28a0285d, 0x5c6d5cda, 0xf8c7f893, 0x86228644,
 565   }
 566 };
 567
 568 __device__ __constant__ u32 Cl[8][256] =
 569 {
 570   {
 571     0xc07830d8, 0x05af4626, 0x7ef991b8, 0x136fcdfb,
 572     0x4ca113cb, 0xa9626d11, 0x08050209, 0x426e9e0d,
 573     0xadee6c9b, 0x590451ff, 0xdebdb90c, 0xfb06f70e,
 574     0xef80f296, 0x5fcede30, 0xfcef3f6d, 0xaa07a4f8,
 575     0x27fdc047, 0x89766535, 0xaccd2b37, 0x048c018a,
 576     0x71155bd2, 0x603c186c, 0xff8af684, 0xb5e16a80,
 577     0xe8693af5, 0x5347ddb3, 0xf6acb321, 0x5eed999c,
 578     0x6d965c43, 0x627a9629, 0xa321e15d, 0x8216aed5,
 579     0xa8412abd, 0x9fb6eee8, 0xa5eb6e92, 0x7b56d79e,
 580     0x8cd92313, 0xd317fd23, 0x6a7f9420, 0x9e95a944,
 581     0xfa25b0a2, 0x06ca8fcf, 0x558d527c, 0x5022145a,
 582     0xe14f7f50, 0x691a5dc9, 0x7fdad614, 0x5cab17d9,
 583     0x8173673c, 0xd234ba8f, 0x80502090, 0xf303f507,
 584     0x16c08bdd, 0xedc67cd3, 0x28110a2d, 0x1fe6ce78,
 585     0x7353d597, 0x25bb4e02, 0x32588273, 0x2c9d0ba7,
 586     0x510153f6, 0xcf94fab2, 0xdcfb3749, 0x8e9fad56,
 587     0x8b30eb70, 0x2371c1cd, 0xc791f8bb, 0x17e3cc71,
 588     0xa68ea77b, 0xb84b2eaf, 0x02468e45, 0x84dc211a,
 589     0x1ec589d4, 0x75995a58, 0x9179632e, 0x381b0e3f,
 590     0x012347ac, 0xea2fb4b0, 0x6cb51bef, 0x85ff66b6,
 591     0x3ff2c65c, 0x100a0412, 0x39384993, 0xafa8e2de,
 592     0x0ecf8dc6, 0xc87d32d1, 0x7270923b, 0x869aaf5f,
 593     0xc31df931, 0x4b48dba8, 0xe22ab6b9, 0x34920dbc,
 594     0xa4c8293e, 0x2dbe4c0b, 0x8dfa64bf, 0xe94a7d59,
 595     0x1b6acff2, 0x78331e77, 0xe6a6b733, 0x74ba1df4,
 596     0x997c6127, 0x26de87eb, 0xbde46889, 0x7a759032,
 597     0xab24e354, 0xf78ff48d, 0xf4ea3d64, 0xc23ebe9d,
 598     0x1da0403d, 0x67d5d00f, 0xd07234ca, 0x192c41b7,
 599     0xc95e757d, 0x9a19a8ce, 0xece53b7f, 0x0daa442f,
 600     0x07e9c863, 0xdb12ff2a, 0xbfa2e6cc, 0x905a2482,
 601     0x3a5d807a, 0x40281048, 0x56e89b95, 0x337bc5df,
 602     0x9690ab4d, 0x611f5fc0, 0x1c830791, 0xf5c97ac8,
 603     0xccf1335b, 0x00000000, 0x36d483f9, 0x4587566e,
 604     0x97b3ece1, 0x64b019e6, 0xfea9b128, 0xd87736c3,
 605     0xc15b7774, 0x112943be, 0x77dfd41d, 0xba0da0ea,
 606     0x124c8a57, 0xcb18fb38, 0x9df060ad, 0x2b74c3c4,
 607     0xe5c37eda, 0x921caac7, 0x791059db, 0x0365c9e9,
 608     0x0fecca6a, 0xb9686903, 0x65935e4a, 0x4ee79d8e,
 609     0xbe81a160, 0xe06c38fc, 0xbb2ee746, 0x52649a1f,
 610     0xe4e03976, 0x8fbceafa, 0x301e0c36, 0x249809ae,
 611     0xf940794b, 0x6359d185, 0x70361c7e, 0xf8633ee7,
 612     0x37f7c455, 0xeea3b53a, 0x29324d81, 0xc4f43152,
 613     0x9b3aef62, 0x66f697a3, 0x35b14a10, 0xf220b2ab,
 614     0x54ae15d0, 0xb7a7e4c5, 0xd5dd72ec, 0x5a619816,
 615     0xca3bbc94, 0xe785f09f, 0xddd870e5, 0x14860598,
 616     0xc6b2bf17, 0x410b57e4, 0x434dd9a1, 0x2ff8c24e,
 617     0xf1457b42, 0x15a54234, 0x94d62508, 0xf0663cee,
 618     0x22528661, 0x76fc93b1, 0xb32be54f, 0x20140824,
 619     0xb208a2e3, 0xbcc72f25, 0x4fc4da22, 0x68391a65,
 620     0x8335e979, 0xb684a369, 0xd79bfca9, 0x3db44819,
 621     0xc5d776fe, 0x313d4b9a, 0x3ed181f0, 0x88552299,
 622     0x0c890383, 0x4a6b9c04, 0xd1517366, 0x0b60cbe0,
 623     0xfdcc78c1, 0x7cbf1ffd, 0xd4fe3540, 0xeb0cf31c,
 624     0xa1676f18, 0x985f268b, 0x7d9c5851, 0xd6b8bb05,
 625     0x6b5cd38c, 0x57cbdc39, 0x6ef395aa, 0x180f061b,
 626     0x8a13acdc, 0x1a49885e, 0xdf9efea0, 0x21374f88,
 627     0x4d825467, 0xb16d6b0a, 0x46e29f87, 0xa202a6f1,
 628     0xae8ba572, 0x58271653, 0x9cd32701, 0x47c1d82b,
 629     0x95f562a4, 0x87b9e8f3, 0xe309f115, 0x0a438c4c,
 630     0x092645a5, 0x3c970fb5, 0xa04428b4, 0x5b42dfba,
 631     0xb04e2ca6, 0xcdd274f7, 0x6fd0d206, 0x482d1241,
 632     0xa7ade0d7, 0xd954716f, 0xceb7bd1e, 0x3b7ec7d6,
 633     0x2edb85e2, 0x2a578468, 0xb4c22d2c, 0x490e55ed,
 634     0x5d885075, 0xda31b886, 0x933fed6b, 0x44a411c2,
 635   },
 636   {
 637     0x18c07830, 0x2305af46, 0xc67ef991, 0xe8136fcd,
 638     0x874ca113, 0xb8a9626d, 0x01080502, 0x4f426e9e,
 639     0x36adee6c, 0xa6590451, 0xd2debdb9, 0xf5fb06f7,
 640     0x79ef80f2, 0x6f5fcede, 0x91fcef3f, 0x52aa07a4,
 641     0x6027fdc0, 0xbc897665, 0x9baccd2b, 0x8e048c01,
 642     0xa371155b, 0x0c603c18, 0x7bff8af6, 0x35b5e16a,
 643     0x1de8693a, 0xe05347dd, 0xd7f6acb3, 0xc25eed99,
 644     0x2e6d965c, 0x4b627a96, 0xfea321e1, 0x578216ae,
 645     0x15a8412a, 0x779fb6ee, 0x37a5eb6e, 0xe57b56d7,
 646     0x9f8cd923, 0xf0d317fd, 0x4a6a7f94, 0xda9e95a9,
 647     0x58fa25b0, 0xc906ca8f, 0x29558d52, 0x0a502214,
 648     0xb1e14f7f, 0xa0691a5d, 0x6b7fdad6, 0x855cab17,
 649     0xbd817367, 0x5dd234ba, 0x10805020, 0xf4f303f5,
 650     0xcb16c08b, 0x3eedc67c, 0x0528110a, 0x671fe6ce,
 651     0xe47353d5, 0x2725bb4e, 0x41325882, 0x8b2c9d0b,
 652     0xa7510153, 0x7dcf94fa, 0x95dcfb37, 0xd88e9fad,
 653     0xfb8b30eb, 0xee2371c1, 0x7cc791f8, 0x6617e3cc,
 654     0xdda68ea7, 0x17b84b2e, 0x4702468e, 0x9e84dc21,
 655     0xca1ec589, 0x2d75995a, 0xbf917963, 0x07381b0e,
 656     0xad012347, 0x5aea2fb4, 0x836cb51b, 0x3385ff66,
 657     0x633ff2c6, 0x02100a04, 0xaa393849, 0x71afa8e2,
 658     0xc80ecf8d, 0x19c87d32, 0x49727092, 0xd9869aaf,
 659     0xf2c31df9, 0xe34b48db, 0x5be22ab6, 0x8834920d,
 660     0x9aa4c829, 0x262dbe4c, 0x328dfa64, 0xb0e94a7d,
 661     0xe91b6acf, 0x0f78331e, 0xd5e6a6b7, 0x8074ba1d,
 662     0xbe997c61, 0xcd26de87, 0x34bde468, 0x487a7590,
 663     0xffab24e3, 0x7af78ff4, 0x90f4ea3d, 0x5fc23ebe,
 664     0x201da040, 0x6867d5d0, 0x1ad07234, 0xae192c41,
 665     0xb4c95e75, 0x549a19a8, 0x93ece53b, 0x220daa44,
 666     0x6407e9c8, 0xf1db12ff, 0x73bfa2e6, 0x12905a24,
 667     0x403a5d80, 0x08402810, 0xc356e89b, 0xec337bc5,
 668     0xdb9690ab, 0xa1611f5f, 0x8d1c8307, 0x3df5c97a,
 669     0x97ccf133, 0x00000000, 0xcf36d483, 0x2b458756,
 670     0x7697b3ec, 0x8264b019, 0xd6fea9b1, 0x1bd87736,
 671     0xb5c15b77, 0xaf112943, 0x6a77dfd4, 0x50ba0da0,
 672     0x45124c8a, 0xf3cb18fb, 0x309df060, 0xef2b74c3,
 673     0x3fe5c37e, 0x55921caa, 0xa2791059, 0xea0365c9,
 674     0x650fecca, 0xbab96869, 0x2f65935e, 0xc04ee79d,
 675     0xdebe81a1, 0x1ce06c38, 0xfdbb2ee7, 0x4d52649a,
 676     0x92e4e039, 0x758fbcea, 0x06301e0c, 0x8a249809,
 677     0xb2f94079, 0xe66359d1, 0x0e70361c, 0x1ff8633e,
 678     0x6237f7c4, 0xd4eea3b5, 0xa829324d, 0x96c4f431,
 679     0xf99b3aef, 0xc566f697, 0x2535b14a, 0x59f220b2,
 680     0x8454ae15, 0x72b7a7e4, 0x39d5dd72, 0x4c5a6198,
 681     0x5eca3bbc, 0x78e785f0, 0x38ddd870, 0x8c148605,
 682     0xd1c6b2bf, 0xa5410b57, 0xe2434dd9, 0x612ff8c2,
 683     0xb3f1457b, 0x2115a542, 0x9c94d625, 0x1ef0663c,
 684     0x43225286, 0xc776fc93, 0xfcb32be5, 0x04201408,
 685     0x51b208a2, 0x99bcc72f, 0x6d4fc4da, 0x0d68391a,
 686     0xfa8335e9, 0xdfb684a3, 0x7ed79bfc, 0x243db448,
 687     0x3bc5d776, 0xab313d4b, 0xce3ed181, 0x11885522,
 688     0x8f0c8903, 0x4e4a6b9c, 0xb7d15173, 0xeb0b60cb,
 689     0x3cfdcc78, 0x817cbf1f, 0x94d4fe35, 0xf7eb0cf3,
 690     0xb9a1676f, 0x13985f26, 0x2c7d9c58, 0xd3d6b8bb,
 691     0xe76b5cd3, 0x6e57cbdc, 0xc46ef395, 0x03180f06,
 692     0x568a13ac, 0x441a4988, 0x7fdf9efe, 0xa921374f,
 693     0x2a4d8254, 0xbbb16d6b, 0xc146e29f, 0x53a202a6,
 694     0xdcae8ba5, 0x0b582716, 0x9d9cd327, 0x6c47c1d8,
 695     0x3195f562, 0x7487b9e8, 0xf6e309f1, 0x460a438c,
 696     0xac092645, 0x893c970f, 0x14a04428, 0xe15b42df,
 697     0x16b04e2c, 0x3acdd274, 0x696fd0d2, 0x09482d12,
 698     0x70a7ade0, 0xb6d95471, 0xd0ceb7bd, 0xed3b7ec7,
 699     0xcc2edb85, 0x422a5784, 0x98b4c22d, 0xa4490e55,
 700     0x285d8850, 0x5cda31b8, 0xf8933fed, 0x8644a411,
 701   },
 702   {
 703     0x6018c078, 0x8c2305af, 0x3fc67ef9, 0x87e8136f,
 704     0x26874ca1, 0xdab8a962, 0x04010805, 0x214f426e,
 705     0xd836adee, 0xa2a65904, 0x6fd2debd, 0xf3f5fb06,
 706     0xf979ef80, 0xa16f5fce, 0x7e91fcef, 0x5552aa07,
 707     0x9d6027fd, 0xcabc8976, 0x569baccd, 0x028e048c,
 708     0xb6a37115, 0x300c603c, 0xf17bff8a, 0xd435b5e1,
 709     0x741de869, 0xa7e05347, 0x7bd7f6ac, 0x2fc25eed,
 710     0xb82e6d96, 0x314b627a, 0xdffea321, 0x41578216,
 711     0x5415a841, 0xc1779fb6, 0xdc37a5eb, 0xb3e57b56,
 712     0x469f8cd9, 0xe7f0d317, 0x354a6a7f, 0x4fda9e95,
 713     0x7d58fa25, 0x03c906ca, 0xa429558d, 0x280a5022,
 714     0xfeb1e14f, 0xbaa0691a, 0xb16b7fda, 0x2e855cab,
 715     0xcebd8173, 0x695dd234, 0x40108050, 0xf7f4f303,
 716     0x0bcb16c0, 0xf83eedc6, 0x14052811, 0x81671fe6,
 717     0xb7e47353, 0x9c2725bb, 0x19413258, 0x168b2c9d,
 718     0xa6a75101, 0xe97dcf94, 0x6e95dcfb, 0x47d88e9f,
 719     0xcbfb8b30, 0x9fee2371, 0xed7cc791, 0x856617e3,
 720     0x53dda68e, 0x5c17b84b, 0x01470246, 0x429e84dc,
 721     0x0fca1ec5, 0xb42d7599, 0xc6bf9179, 0x1c07381b,
 722     0x8ead0123, 0x755aea2f, 0x36836cb5, 0xcc3385ff,
 723     0x91633ff2, 0x0802100a, 0x92aa3938, 0xd971afa8,
 724     0x07c80ecf, 0x6419c87d, 0x39497270, 0x43d9869a,
 725     0xeff2c31d, 0xabe34b48, 0x715be22a, 0x1a883492,
 726     0x529aa4c8, 0x98262dbe, 0xc8328dfa, 0xfab0e94a,
 727     0x83e91b6a, 0x3c0f7833, 0x73d5e6a6, 0x3a8074ba,
 728     0xc2be997c, 0x13cd26de, 0xd034bde4, 0x3d487a75,
 729     0xdbffab24, 0xf57af78f, 0x7a90f4ea, 0x615fc23e,
 730     0x80201da0, 0xbd6867d5, 0x681ad072, 0x82ae192c,
 731     0xeab4c95e, 0x4d549a19, 0x7693ece5, 0x88220daa,
 732     0x8d6407e9, 0xe3f1db12, 0xd173bfa2, 0x4812905a,
 733     0x1d403a5d, 0x20084028, 0x2bc356e8, 0x97ec337b,
 734     0x4bdb9690, 0xbea1611f, 0x0e8d1c83, 0xf43df5c9,
 735     0x6697ccf1, 0x00000000, 0x1bcf36d4, 0xac2b4587,
 736     0xc57697b3, 0x328264b0, 0x7fd6fea9, 0x6c1bd877,
 737     0xeeb5c15b, 0x86af1129, 0xb56a77df, 0x5d50ba0d,
 738     0x0945124c, 0xebf3cb18, 0xc0309df0, 0x9bef2b74,
 739     0xfc3fe5c3, 0x4955921c, 0xb2a27910, 0x8fea0365,
 740     0x89650fec, 0xd2bab968, 0xbc2f6593, 0x27c04ee7,
 741     0x5fdebe81, 0x701ce06c, 0xd3fdbb2e, 0x294d5264,
 742     0x7292e4e0, 0xc9758fbc, 0x1806301e, 0x128a2498,
 743     0xf2b2f940, 0xbfe66359, 0x380e7036, 0x7c1ff863,
 744     0x956237f7, 0x77d4eea3, 0x9aa82932, 0x6296c4f4,
 745     0xc3f99b3a, 0x33c566f6, 0x942535b1, 0x7959f220,
 746     0x2a8454ae, 0xd572b7a7, 0xe439d5dd, 0x2d4c5a61,
 747     0x655eca3b, 0xfd78e785, 0xe038ddd8, 0x0a8c1486,
 748     0x63d1c6b2, 0xaea5410b, 0xafe2434d, 0x99612ff8,
 749     0xf6b3f145, 0x842115a5, 0x4a9c94d6, 0x781ef066,
 750     0x11432252, 0x3bc776fc, 0xd7fcb32b, 0x10042014,
 751     0x5951b208, 0x5e99bcc7, 0xa96d4fc4, 0x340d6839,
 752     0xcffa8335, 0x5bdfb684, 0xe57ed79b, 0x90243db4,
 753     0xec3bc5d7, 0x96ab313d, 0x1fce3ed1, 0x44118855,
 754     0x068f0c89, 0x254e4a6b, 0xe6b7d151, 0x8beb0b60,
 755     0xf03cfdcc, 0x3e817cbf, 0x6a94d4fe, 0xfbf7eb0c,
 756     0xdeb9a167, 0x4c13985f, 0xb02c7d9c, 0x6bd3d6b8,
 757     0xbbe76b5c, 0xa56e57cb, 0x37c46ef3, 0x0c03180f,
 758     0x45568a13, 0x0d441a49, 0xe17fdf9e, 0x9ea92137,
 759     0xa82a4d82, 0xd6bbb16d, 0x23c146e2, 0x5153a202,
 760     0x57dcae8b, 0x2c0b5827, 0x4e9d9cd3, 0xad6c47c1,
 761     0xc43195f5, 0xcd7487b9, 0xfff6e309, 0x05460a43,
 762     0x8aac0926, 0x1e893c97, 0x5014a044, 0xa3e15b42,
 763     0x5816b04e, 0xe83acdd2, 0xb9696fd0, 0x2409482d,
 764     0xdd70a7ad, 0xe2b6d954, 0x67d0ceb7, 0x93ed3b7e,
 765     0x17cc2edb, 0x15422a57, 0x5a98b4c2, 0xaaa4490e,
 766     0xa0285d88, 0x6d5cda31, 0xc7f8933f, 0x228644a4,
 767   },
 768   {
 769     0x186018c0, 0x238c2305, 0xc63fc67e, 0xe887e813,
 770     0x8726874c, 0xb8dab8a9, 0x01040108, 0x4f214f42,
 771     0x36d836ad, 0xa6a2a659, 0xd26fd2de, 0xf5f3f5fb,
 772     0x79f979ef, 0x6fa16f5f, 0x917e91fc, 0x525552aa,
 773     0x609d6027, 0xbccabc89, 0x9b569bac, 0x8e028e04,
 774     0xa3b6a371, 0x0c300c60, 0x7bf17bff, 0x35d435b5,
 775     0x1d741de8, 0xe0a7e053, 0xd77bd7f6, 0xc22fc25e,
 776     0x2eb82e6d, 0x4b314b62, 0xfedffea3, 0x57415782,
 777     0x155415a8, 0x77c1779f, 0x37dc37a5, 0xe5b3e57b,
 778     0x9f469f8c, 0xf0e7f0d3, 0x4a354a6a, 0xda4fda9e,
 779     0x587d58fa, 0xc903c906, 0x29a42955, 0x0a280a50,
 780     0xb1feb1e1, 0xa0baa069, 0x6bb16b7f, 0x852e855c,
 781     0xbdcebd81, 0x5d695dd2, 0x10401080, 0xf4f7f4f3,
 782     0xcb0bcb16, 0x3ef83eed, 0x05140528, 0x6781671f,
 783     0xe4b7e473, 0x279c2725, 0x41194132, 0x8b168b2c,
 784     0xa7a6a751, 0x7de97dcf, 0x956e95dc, 0xd847d88e,
 785     0xfbcbfb8b, 0xee9fee23, 0x7ced7cc7, 0x66856617,
 786     0xdd53dda6, 0x175c17b8, 0x47014702, 0x9e429e84,
 787     0xca0fca1e, 0x2db42d75, 0xbfc6bf91, 0x071c0738,
 788     0xad8ead01, 0x5a755aea, 0x8336836c, 0x33cc3385,
 789     0x6391633f, 0x02080210, 0xaa92aa39, 0x71d971af,
 790     0xc807c80e, 0x196419c8, 0x49394972, 0xd943d986,
 791     0xf2eff2c3, 0xe3abe34b, 0x5b715be2, 0x881a8834,
 792     0x9a529aa4, 0x2698262d, 0x32c8328d, 0xb0fab0e9,
 793     0xe983e91b, 0x0f3c0f78, 0xd573d5e6, 0x803a8074,
 794     0xbec2be99, 0xcd13cd26, 0x34d034bd, 0x483d487a,
 795     0xffdbffab, 0x7af57af7, 0x907a90f4, 0x5f615fc2,
 796     0x2080201d, 0x68bd6867, 0x1a681ad0, 0xae82ae19,
 797     0xb4eab4c9, 0x544d549a, 0x937693ec, 0x2288220d,
 798     0x648d6407, 0xf1e3f1db, 0x73d173bf, 0x12481290,
 799     0x401d403a, 0x08200840, 0xc32bc356, 0xec97ec33,
 800     0xdb4bdb96, 0xa1bea161, 0x8d0e8d1c, 0x3df43df5,
 801     0x976697cc, 0x00000000, 0xcf1bcf36, 0x2bac2b45,
 802     0x76c57697, 0x82328264, 0xd67fd6fe, 0x1b6c1bd8,
 803     0xb5eeb5c1, 0xaf86af11, 0x6ab56a77, 0x505d50ba,
 804     0x45094512, 0xf3ebf3cb, 0x30c0309d, 0xef9bef2b,
 805     0x3ffc3fe5, 0x55495592, 0xa2b2a279, 0xea8fea03,
 806     0x6589650f, 0xbad2bab9, 0x2fbc2f65, 0xc027c04e,
 807     0xde5fdebe, 0x1c701ce0, 0xfdd3fdbb, 0x4d294d52,
 808     0x927292e4, 0x75c9758f, 0x06180630, 0x8a128a24,
 809     0xb2f2b2f9, 0xe6bfe663, 0x0e380e70, 0x1f7c1ff8,
 810     0x62956237, 0xd477d4ee, 0xa89aa829, 0x966296c4,
 811     0xf9c3f99b, 0xc533c566, 0x25942535, 0x597959f2,
 812     0x842a8454, 0x72d572b7, 0x39e439d5, 0x4c2d4c5a,
 813     0x5e655eca, 0x78fd78e7, 0x38e038dd, 0x8c0a8c14,
 814     0xd163d1c6, 0xa5aea541, 0xe2afe243, 0x6199612f,
 815     0xb3f6b3f1, 0x21842115, 0x9c4a9c94, 0x1e781ef0,
 816     0x43114322, 0xc73bc776, 0xfcd7fcb3, 0x04100420,
 817     0x515951b2, 0x995e99bc, 0x6da96d4f, 0x0d340d68,
 818     0xfacffa83, 0xdf5bdfb6, 0x7ee57ed7, 0x2490243d,
 819     0x3bec3bc5, 0xab96ab31, 0xce1fce3e, 0x11441188,
 820     0x8f068f0c, 0x4e254e4a, 0xb7e6b7d1, 0xeb8beb0b,
 821     0x3cf03cfd, 0x813e817c, 0x946a94d4, 0xf7fbf7eb,
 822     0xb9deb9a1, 0x134c1398, 0x2cb02c7d, 0xd36bd3d6,
 823     0xe7bbe76b, 0x6ea56e57, 0xc437c46e, 0x030c0318,
 824     0x5645568a, 0x440d441a, 0x7fe17fdf, 0xa99ea921,
 825     0x2aa82a4d, 0xbbd6bbb1, 0xc123c146, 0x535153a2,
 826     0xdc57dcae, 0x0b2c0b58, 0x9d4e9d9c, 0x6cad6c47,
 827     0x31c43195, 0x74cd7487, 0xf6fff6e3, 0x4605460a,
 828     0xac8aac09, 0x891e893c, 0x145014a0, 0xe1a3e15b,
 829     0x165816b0, 0x3ae83acd, 0x69b9696f, 0x09240948,
 830     0x70dd70a7, 0xb6e2b6d9, 0xd067d0ce, 0xed93ed3b,
 831     0xcc17cc2e, 0x4215422a, 0x985a98b4, 0xa4aaa449,
 832     0x28a0285d, 0x5c6d5cda, 0xf8c7f893, 0x86228644,
 833   },
 834   {
 835     0x18186018, 0x23238c23, 0xc6c63fc6, 0xe8e887e8,
 836     0x87872687, 0xb8b8dab8, 0x01010401, 0x4f4f214f,
 837     0x3636d836, 0xa6a6a2a6, 0xd2d26fd2, 0xf5f5f3f5,
 838     0x7979f979, 0x6f6fa16f, 0x91917e91, 0x52525552,
 839     0x60609d60, 0xbcbccabc, 0x9b9b569b, 0x8e8e028e,
 840     0xa3a3b6a3, 0x0c0c300c, 0x7b7bf17b, 0x3535d435,
 841     0x1d1d741d, 0xe0e0a7e0, 0xd7d77bd7, 0xc2c22fc2,
 842     0x2e2eb82e, 0x4b4b314b, 0xfefedffe, 0x57574157,
 843     0x15155415, 0x7777c177, 0x3737dc37, 0xe5e5b3e5,
 844     0x9f9f469f, 0xf0f0e7f0, 0x4a4a354a, 0xdada4fda,
 845     0x58587d58, 0xc9c903c9, 0x2929a429, 0x0a0a280a,
 846     0xb1b1feb1, 0xa0a0baa0, 0x6b6bb16b, 0x85852e85,
 847     0xbdbdcebd, 0x5d5d695d, 0x10104010, 0xf4f4f7f4,
 848     0xcbcb0bcb, 0x3e3ef83e, 0x05051405, 0x67678167,
 849     0xe4e4b7e4, 0x27279c27, 0x41411941, 0x8b8b168b,
 850     0xa7a7a6a7, 0x7d7de97d, 0x95956e95, 0xd8d847d8,
 851     0xfbfbcbfb, 0xeeee9fee, 0x7c7ced7c, 0x66668566,
 852     0xdddd53dd, 0x17175c17, 0x47470147, 0x9e9e429e,
 853     0xcaca0fca, 0x2d2db42d, 0xbfbfc6bf, 0x07071c07,
 854     0xadad8ead, 0x5a5a755a, 0x83833683, 0x3333cc33,
 855     0x63639163, 0x02020802, 0xaaaa92aa, 0x7171d971,
 856     0xc8c807c8, 0x19196419, 0x49493949, 0xd9d943d9,
 857     0xf2f2eff2, 0xe3e3abe3, 0x5b5b715b, 0x88881a88,
 858     0x9a9a529a, 0x26269826, 0x3232c832, 0xb0b0fab0,
 859     0xe9e983e9, 0x0f0f3c0f, 0xd5d573d5, 0x80803a80,
 860     0xbebec2be, 0xcdcd13cd, 0x3434d034, 0x48483d48,
 861     0xffffdbff, 0x7a7af57a, 0x90907a90, 0x5f5f615f,
 862     0x20208020, 0x6868bd68, 0x1a1a681a, 0xaeae82ae,
 863     0xb4b4eab4, 0x54544d54, 0x93937693, 0x22228822,
 864     0x64648d64, 0xf1f1e3f1, 0x7373d173, 0x12124812,
 865     0x40401d40, 0x08082008, 0xc3c32bc3, 0xecec97ec,
 866     0xdbdb4bdb, 0xa1a1bea1, 0x8d8d0e8d, 0x3d3df43d,
 867     0x97976697, 0x00000000, 0xcfcf1bcf, 0x2b2bac2b,
 868     0x7676c576, 0x82823282, 0xd6d67fd6, 0x1b1b6c1b,
 869     0xb5b5eeb5, 0xafaf86af, 0x6a6ab56a, 0x50505d50,
 870     0x45450945, 0xf3f3ebf3, 0x3030c030, 0xefef9bef,
 871     0x3f3ffc3f, 0x55554955, 0xa2a2b2a2, 0xeaea8fea,
 872     0x65658965, 0xbabad2ba, 0x2f2fbc2f, 0xc0c027c0,
 873     0xdede5fde, 0x1c1c701c, 0xfdfdd3fd, 0x4d4d294d,
 874     0x92927292, 0x7575c975, 0x06061806, 0x8a8a128a,
 875     0xb2b2f2b2, 0xe6e6bfe6, 0x0e0e380e, 0x1f1f7c1f,
 876     0x62629562, 0xd4d477d4, 0xa8a89aa8, 0x96966296,
 877     0xf9f9c3f9, 0xc5c533c5, 0x25259425, 0x59597959,
 878     0x84842a84, 0x7272d572, 0x3939e439, 0x4c4c2d4c,
 879     0x5e5e655e, 0x7878fd78, 0x3838e038, 0x8c8c0a8c,
 880     0xd1d163d1, 0xa5a5aea5, 0xe2e2afe2, 0x61619961,
 881     0xb3b3f6b3, 0x21218421, 0x9c9c4a9c, 0x1e1e781e,
 882     0x43431143, 0xc7c73bc7, 0xfcfcd7fc, 0x04041004,
 883     0x51515951, 0x99995e99, 0x6d6da96d, 0x0d0d340d,
 884     0xfafacffa, 0xdfdf5bdf, 0x7e7ee57e, 0x24249024,
 885     0x3b3bec3b, 0xabab96ab, 0xcece1fce, 0x11114411,
 886     0x8f8f068f, 0x4e4e254e, 0xb7b7e6b7, 0xebeb8beb,
 887     0x3c3cf03c, 0x81813e81, 0x94946a94, 0xf7f7fbf7,
 888     0xb9b9deb9, 0x13134c13, 0x2c2cb02c, 0xd3d36bd3,
 889     0xe7e7bbe7, 0x6e6ea56e, 0xc4c437c4, 0x03030c03,
 890     0x56564556, 0x44440d44, 0x7f7fe17f, 0xa9a99ea9,
 891     0x2a2aa82a, 0xbbbbd6bb, 0xc1c123c1, 0x53535153,
 892     0xdcdc57dc, 0x0b0b2c0b, 0x9d9d4e9d, 0x6c6cad6c,
 893     0x3131c431, 0x7474cd74, 0xf6f6fff6, 0x46460546,
 894     0xacac8aac, 0x89891e89, 0x14145014, 0xe1e1a3e1,
 895     0x16165816, 0x3a3ae83a, 0x6969b969, 0x09092409,
 896     0x7070dd70, 0xb6b6e2b6, 0xd0d067d0, 0xeded93ed,
 897     0xcccc17cc, 0x42421542, 0x98985a98, 0xa4a4aaa4,
 898     0x2828a028, 0x5c5c6d5c, 0xf8f8c7f8, 0x86862286,
 899   },
 900   {
 901     0xd8181860, 0x2623238c, 0xb8c6c63f, 0xfbe8e887,
 902     0xcb878726, 0x11b8b8da, 0x09010104, 0x0d4f4f21,
 903     0x9b3636d8, 0xffa6a6a2, 0x0cd2d26f, 0x0ef5f5f3,
 904     0x967979f9, 0x306f6fa1, 0x6d91917e, 0xf8525255,
 905     0x4760609d, 0x35bcbcca, 0x379b9b56, 0x8a8e8e02,
 906     0xd2a3a3b6, 0x6c0c0c30, 0x847b7bf1, 0x803535d4,
 907     0xf51d1d74, 0xb3e0e0a7, 0x21d7d77b, 0x9cc2c22f,
 908     0x432e2eb8, 0x294b4b31, 0x5dfefedf, 0xd5575741,
 909     0xbd151554, 0xe87777c1, 0x923737dc, 0x9ee5e5b3,
 910     0x139f9f46, 0x23f0f0e7, 0x204a4a35, 0x44dada4f,
 911     0xa258587d, 0xcfc9c903, 0x7c2929a4, 0x5a0a0a28,
 912     0x50b1b1fe, 0xc9a0a0ba, 0x146b6bb1, 0xd985852e,
 913     0x3cbdbdce, 0x8f5d5d69, 0x90101040, 0x07f4f4f7,
 914     0xddcbcb0b, 0xd33e3ef8, 0x2d050514, 0x78676781,
 915     0x97e4e4b7, 0x0227279c, 0x73414119, 0xa78b8b16,
 916     0xf6a7a7a6, 0xb27d7de9, 0x4995956e, 0x56d8d847,
 917     0x70fbfbcb, 0xcdeeee9f, 0xbb7c7ced, 0x71666685,
 918     0x7bdddd53, 0xaf17175c, 0x45474701, 0x1a9e9e42,
 919     0xd4caca0f, 0x582d2db4, 0x2ebfbfc6, 0x3f07071c,
 920     0xacadad8e, 0xb05a5a75, 0xef838336, 0xb63333cc,
 921     0x5c636391, 0x12020208, 0x93aaaa92, 0xde7171d9,
 922     0xc6c8c807, 0xd1191964, 0x3b494939, 0x5fd9d943,
 923     0x31f2f2ef, 0xa8e3e3ab, 0xb95b5b71, 0xbc88881a,
 924     0x3e9a9a52, 0x0b262698, 0xbf3232c8, 0x59b0b0fa,
 925     0xf2e9e983, 0x770f0f3c, 0x33d5d573, 0xf480803a,
 926     0x27bebec2, 0xebcdcd13, 0x893434d0, 0x3248483d,
 927     0x54ffffdb, 0x8d7a7af5, 0x6490907a, 0x9d5f5f61,
 928     0x3d202080, 0x0f6868bd, 0xca1a1a68, 0xb7aeae82,
 929     0x7db4b4ea, 0xce54544d, 0x7f939376, 0x2f222288,
 930     0x6364648d, 0x2af1f1e3, 0xcc7373d1, 0x82121248,
 931     0x7a40401d, 0x48080820, 0x95c3c32b, 0xdfecec97,
 932     0x4ddbdb4b, 0xc0a1a1be, 0x918d8d0e, 0xc83d3df4,
 933     0x5b979766, 0x00000000, 0xf9cfcf1b, 0x6e2b2bac,
 934     0xe17676c5, 0xe6828232, 0x28d6d67f, 0xc31b1b6c,
 935     0x74b5b5ee, 0xbeafaf86, 0x1d6a6ab5, 0xea50505d,
 936     0x57454509, 0x38f3f3eb, 0xad3030c0, 0xc4efef9b,
 937     0xda3f3ffc, 0xc7555549, 0xdba2a2b2, 0xe9eaea8f,
 938     0x6a656589, 0x03babad2, 0x4a2f2fbc, 0x8ec0c027,
 939     0x60dede5f, 0xfc1c1c70, 0x46fdfdd3, 0x1f4d4d29,
 940     0x76929272, 0xfa7575c9, 0x36060618, 0xae8a8a12,
 941     0x4bb2b2f2, 0x85e6e6bf, 0x7e0e0e38, 0xe71f1f7c,
 942     0x55626295, 0x3ad4d477, 0x81a8a89a, 0x52969662,
 943     0x62f9f9c3, 0xa3c5c533, 0x10252594, 0xab595979,
 944     0xd084842a, 0xc57272d5, 0xec3939e4, 0x164c4c2d,
 945     0x945e5e65, 0x9f7878fd, 0xe53838e0, 0x988c8c0a,
 946     0x17d1d163, 0xe4a5a5ae, 0xa1e2e2af, 0x4e616199,
 947     0x42b3b3f6, 0x34212184, 0x089c9c4a, 0xee1e1e78,
 948     0x61434311, 0xb1c7c73b, 0x4ffcfcd7, 0x24040410,
 949     0xe3515159, 0x2599995e, 0x226d6da9, 0x650d0d34,
 950     0x79fafacf, 0x69dfdf5b, 0xa97e7ee5, 0x19242490,
 951     0xfe3b3bec, 0x9aabab96, 0xf0cece1f, 0x99111144,
 952     0x838f8f06, 0x044e4e25, 0x66b7b7e6, 0xe0ebeb8b,
 953     0xc13c3cf0, 0xfd81813e, 0x4094946a, 0x1cf7f7fb,
 954     0x18b9b9de, 0x8b13134c, 0x512c2cb0, 0x05d3d36b,
 955     0x8ce7e7bb, 0x396e6ea5, 0xaac4c437, 0x1b03030c,
 956     0xdc565645, 0x5e44440d, 0xa07f7fe1, 0x88a9a99e,
 957     0x672a2aa8, 0x0abbbbd6, 0x87c1c123, 0xf1535351,
 958     0x72dcdc57, 0x530b0b2c, 0x019d9d4e, 0x2b6c6cad,
 959     0xa43131c4, 0xf37474cd, 0x15f6f6ff, 0x4c464605,
 960     0xa5acac8a, 0xb589891e, 0xb4141450, 0xbae1e1a3,
 961     0xa6161658, 0xf73a3ae8, 0x066969b9, 0x41090924,
 962     0xd77070dd, 0x6fb6b6e2, 0x1ed0d067, 0xd6eded93,
 963     0xe2cccc17, 0x68424215, 0x2c98985a, 0xeda4a4aa,
 964     0x752828a0, 0x865c5c6d, 0x6bf8f8c7, 0xc2868622,
 965   },
 966   {
 967     0x30d81818, 0x46262323, 0x91b8c6c6, 0xcdfbe8e8,
 968     0x13cb8787, 0x6d11b8b8, 0x02090101, 0x9e0d4f4f,
 969     0x6c9b3636, 0x51ffa6a6, 0xb90cd2d2, 0xf70ef5f5,
 970     0xf2967979, 0xde306f6f, 0x3f6d9191, 0xa4f85252,
 971     0xc0476060, 0x6535bcbc, 0x2b379b9b, 0x018a8e8e,
 972     0x5bd2a3a3, 0x186c0c0c, 0xf6847b7b, 0x6a803535,
 973     0x3af51d1d, 0xddb3e0e0, 0xb321d7d7, 0x999cc2c2,
 974     0x5c432e2e, 0x96294b4b, 0xe15dfefe, 0xaed55757,
 975     0x2abd1515, 0xeee87777, 0x6e923737, 0xd79ee5e5,
 976     0x23139f9f, 0xfd23f0f0, 0x94204a4a, 0xa944dada,
 977     0xb0a25858, 0x8fcfc9c9, 0x527c2929, 0x145a0a0a,
 978     0x7f50b1b1, 0x5dc9a0a0, 0xd6146b6b, 0x17d98585,
 979     0x673cbdbd, 0xba8f5d5d, 0x20901010, 0xf507f4f4,
 980     0x8bddcbcb, 0x7cd33e3e, 0x0a2d0505, 0xce786767,
 981     0xd597e4e4, 0x4e022727, 0x82734141, 0x0ba78b8b,
 982     0x53f6a7a7, 0xfab27d7d, 0x37499595, 0xad56d8d8,
 983     0xeb70fbfb, 0xc1cdeeee, 0xf8bb7c7c, 0xcc716666,
 984     0xa77bdddd, 0x2eaf1717, 0x8e454747, 0x211a9e9e,
 985     0x89d4caca, 0x5a582d2d, 0x632ebfbf, 0x0e3f0707,
 986     0x47acadad, 0xb4b05a5a, 0x1bef8383, 0x66b63333,
 987     0xc65c6363, 0x04120202, 0x4993aaaa, 0xe2de7171,
 988     0x8dc6c8c8, 0x32d11919, 0x923b4949, 0xaf5fd9d9,
 989     0xf931f2f2, 0xdba8e3e3, 0xb6b95b5b, 0x0dbc8888,
 990     0x293e9a9a, 0x4c0b2626, 0x64bf3232, 0x7d59b0b0,
 991     0xcff2e9e9, 0x1e770f0f, 0xb733d5d5, 0x1df48080,
 992     0x6127bebe, 0x87ebcdcd, 0x68893434, 0x90324848,
 993     0xe354ffff, 0xf48d7a7a, 0x3d649090, 0xbe9d5f5f,
 994     0x403d2020, 0xd00f6868, 0x34ca1a1a, 0x41b7aeae,
 995     0x757db4b4, 0xa8ce5454, 0x3b7f9393, 0x442f2222,
 996     0xc8636464, 0xff2af1f1, 0xe6cc7373, 0x24821212,
 997     0x807a4040, 0x10480808, 0x9b95c3c3, 0xc5dfecec,
 998     0xab4ddbdb, 0x5fc0a1a1, 0x07918d8d, 0x7ac83d3d,
 999     0x335b9797, 0x00000000, 0x83f9cfcf, 0x566e2b2b,
1000     0xece17676, 0x19e68282, 0xb128d6d6, 0x36c31b1b,
1001     0x7774b5b5, 0x43beafaf, 0xd41d6a6a, 0xa0ea5050,
1002     0x8a574545, 0xfb38f3f3, 0x60ad3030, 0xc3c4efef,
1003     0x7eda3f3f, 0xaac75555, 0x59dba2a2, 0xc9e9eaea,
1004     0xca6a6565, 0x6903baba, 0x5e4a2f2f, 0x9d8ec0c0,
1005     0xa160dede, 0x38fc1c1c, 0xe746fdfd, 0x9a1f4d4d,
1006     0x39769292, 0xeafa7575, 0x0c360606, 0x09ae8a8a,
1007     0x794bb2b2, 0xd185e6e6, 0x1c7e0e0e, 0x3ee71f1f,
1008     0xc4556262, 0xb53ad4d4, 0x4d81a8a8, 0x31529696,
1009     0xef62f9f9, 0x97a3c5c5, 0x4a102525, 0xb2ab5959,
1010     0x15d08484, 0xe4c57272, 0x72ec3939, 0x98164c4c,
1011     0xbc945e5e, 0xf09f7878, 0x70e53838, 0x05988c8c,
1012     0xbf17d1d1, 0x57e4a5a5, 0xd9a1e2e2, 0xc24e6161,
1013     0x7b42b3b3, 0x42342121, 0x25089c9c, 0x3cee1e1e,
1014     0x86614343, 0x93b1c7c7, 0xe54ffcfc, 0x08240404,
1015     0xa2e35151, 0x2f259999, 0xda226d6d, 0x1a650d0d,
1016     0xe979fafa, 0xa369dfdf, 0xfca97e7e, 0x48192424,
1017     0x76fe3b3b, 0x4b9aabab, 0x81f0cece, 0x22991111,
1018     0x03838f8f, 0x9c044e4e, 0x7366b7b7, 0xcbe0ebeb,
1019     0x78c13c3c, 0x1ffd8181, 0x35409494, 0xf31cf7f7,
1020     0x6f18b9b9, 0x268b1313, 0x58512c2c, 0xbb05d3d3,
1021     0xd38ce7e7, 0xdc396e6e, 0x95aac4c4, 0x061b0303,
1022     0xacdc5656, 0x885e4444, 0xfea07f7f, 0x4f88a9a9,
1023     0x54672a2a, 0x6b0abbbb, 0x9f87c1c1, 0xa6f15353,
1024     0xa572dcdc, 0x16530b0b, 0x27019d9d, 0xd82b6c6c,
1025     0x62a43131, 0xe8f37474, 0xf115f6f6, 0x8c4c4646,
1026     0x45a5acac, 0x0fb58989, 0x28b41414, 0xdfbae1e1,
1027     0x2ca61616, 0x74f73a3a, 0xd2066969, 0x12410909,
1028     0xe0d77070, 0x716fb6b6, 0xbd1ed0d0, 0xc7d6eded,
1029     0x85e2cccc, 0x84684242, 0x2d2c9898, 0x55eda4a4,
1030     0x50752828, 0xb8865c5c, 0xed6bf8f8, 0x11c28686,
1031   },
1032   {
1033     0x7830d818, 0xaf462623, 0xf991b8c6, 0x6fcdfbe8,
1034     0xa113cb87, 0x626d11b8, 0x05020901, 0x6e9e0d4f,
1035     0xee6c9b36, 0x0451ffa6, 0xbdb90cd2, 0x06f70ef5,
1036     0x80f29679, 0xcede306f, 0xef3f6d91, 0x07a4f852,
1037     0xfdc04760, 0x766535bc, 0xcd2b379b, 0x8c018a8e,
1038     0x155bd2a3, 0x3c186c0c, 0x8af6847b, 0xe16a8035,
1039     0x693af51d, 0x47ddb3e0, 0xacb321d7, 0xed999cc2,
1040     0x965c432e, 0x7a96294b, 0x21e15dfe, 0x16aed557,
1041     0x412abd15, 0xb6eee877, 0xeb6e9237, 0x56d79ee5,
1042     0xd923139f, 0x17fd23f0, 0x7f94204a, 0x95a944da,
1043     0x25b0a258, 0xca8fcfc9, 0x8d527c29, 0x22145a0a,
1044     0x4f7f50b1, 0x1a5dc9a0, 0xdad6146b, 0xab17d985,
1045     0x73673cbd, 0x34ba8f5d, 0x50209010, 0x03f507f4,
1046     0xc08bddcb, 0xc67cd33e, 0x110a2d05, 0xe6ce7867,
1047     0x53d597e4, 0xbb4e0227, 0x58827341, 0x9d0ba78b,
1048     0x0153f6a7, 0x94fab27d, 0xfb374995, 0x9fad56d8,
1049     0x30eb70fb, 0x71c1cdee, 0x91f8bb7c, 0xe3cc7166,
1050     0x8ea77bdd, 0x4b2eaf17, 0x468e4547, 0xdc211a9e,
1051     0xc589d4ca, 0x995a582d, 0x79632ebf, 0x1b0e3f07,
1052     0x2347acad, 0x2fb4b05a, 0xb51bef83, 0xff66b633,
1053     0xf2c65c63, 0x0a041202, 0x384993aa, 0xa8e2de71,
1054     0xcf8dc6c8, 0x7d32d119, 0x70923b49, 0x9aaf5fd9,
1055     0x1df931f2, 0x48dba8e3, 0x2ab6b95b, 0x920dbc88,
1056     0xc8293e9a, 0xbe4c0b26, 0xfa64bf32, 0x4a7d59b0,
1057     0x6acff2e9, 0x331e770f, 0xa6b733d5, 0xba1df480,
1058     0x7c6127be, 0xde87ebcd, 0xe4688934, 0x75903248,
1059     0x24e354ff, 0x8ff48d7a, 0xea3d6490, 0x3ebe9d5f,
1060     0xa0403d20, 0xd5d00f68, 0x7234ca1a, 0x2c41b7ae,
1061     0x5e757db4, 0x19a8ce54, 0xe53b7f93, 0xaa442f22,
1062     0xe9c86364, 0x12ff2af1, 0xa2e6cc73, 0x5a248212,
1063     0x5d807a40, 0x28104808, 0xe89b95c3, 0x7bc5dfec,
1064     0x90ab4ddb, 0x1f5fc0a1, 0x8307918d, 0xc97ac83d,
1065     0xf1335b97, 0x00000000, 0xd483f9cf, 0x87566e2b,
1066     0xb3ece176, 0xb019e682, 0xa9b128d6, 0x7736c31b,
1067     0x5b7774b5, 0x2943beaf, 0xdfd41d6a, 0x0da0ea50,
1068     0x4c8a5745, 0x18fb38f3, 0xf060ad30, 0x74c3c4ef,
1069     0xc37eda3f, 0x1caac755, 0x1059dba2, 0x65c9e9ea,
1070     0xecca6a65, 0x686903ba, 0x935e4a2f, 0xe79d8ec0,
1071     0x81a160de, 0x6c38fc1c, 0x2ee746fd, 0x649a1f4d,
1072     0xe0397692, 0xbceafa75, 0x1e0c3606, 0x9809ae8a,
1073     0x40794bb2, 0x59d185e6, 0x361c7e0e, 0x633ee71f,
1074     0xf7c45562, 0xa3b53ad4, 0x324d81a8, 0xf4315296,
1075     0x3aef62f9, 0xf697a3c5, 0xb14a1025, 0x20b2ab59,
1076     0xae15d084, 0xa7e4c572, 0xdd72ec39, 0x6198164c,
1077     0x3bbc945e, 0x85f09f78, 0xd870e538, 0x8605988c,
1078     0xb2bf17d1, 0x0b57e4a5, 0x4dd9a1e2, 0xf8c24e61,
1079     0x457b42b3, 0xa5423421, 0xd625089c, 0x663cee1e,
1080     0x52866143, 0xfc93b1c7, 0x2be54ffc, 0x14082404,
1081     0x08a2e351, 0xc72f2599, 0xc4da226d, 0x391a650d,
1082     0x35e979fa, 0x84a369df, 0x9bfca97e, 0xb4481924,
1083     0xd776fe3b, 0x3d4b9aab, 0xd181f0ce, 0x55229911,
1084     0x8903838f, 0x6b9c044e, 0x517366b7, 0x60cbe0eb,
1085     0xcc78c13c, 0xbf1ffd81, 0xfe354094, 0x0cf31cf7,
1086     0x676f18b9, 0x5f268b13, 0x9c58512c, 0xb8bb05d3,
1087     0x5cd38ce7, 0xcbdc396e, 0xf395aac4, 0x0f061b03,
1088     0x13acdc56, 0x49885e44, 0x9efea07f, 0x374f88a9,
1089     0x8254672a, 0x6d6b0abb, 0xe29f87c1, 0x02a6f153,
1090     0x8ba572dc, 0x2716530b, 0xd327019d, 0xc1d82b6c,
1091     0xf562a431, 0xb9e8f374, 0x09f115f6, 0x438c4c46,
1092     0x2645a5ac, 0x970fb589, 0x4428b414, 0x42dfbae1,
1093     0x4e2ca616, 0xd274f73a, 0xd0d20669, 0x2d124109,
1094     0xade0d770, 0x54716fb6, 0xb7bd1ed0, 0x7ec7d6ed,
1095     0xdb85e2cc, 0x57846842, 0xc22d2c98, 0x0e55eda4,
1096     0x88507528, 0x31b8865c, 0x3fed6bf8, 0xa411c286,
1097   },
1098 };
1099
1100 #ifdef VECT_SIZE1
1101 #define BOX(S,n,i) (u32) ((S)[(n)][(i)])
1102 #endif
1103
1104 __device__ __constant__ u32 rch[R + 1] =
1105 {
1106   0x00000000,
1107   0x1823c6e8,
1108   0x36a6d2f5,
1109   0x60bc9b8e,
1110   0x1de0d7c2,
1111   0x157737e5,
1112   0x58c9290a,
1113   0xbd5d10f4,
1114   0xe427418b,
1115   0xfbee7c66,
1116   0xca2dbf07,
1117 };
1118
1119 __device__ __constant__ u32 rcl[R + 1] =
1120 {
1121   0x00000000,
1122   0x87b8014f,
1123   0x796f9152,
1124   0xa30c7b35,
1125   0x2e4bfe57,
1126   0x9ff04ada,
1127   0xb1a06b85,
1128   0xcb3e0567,
1129   0xa77d95d8,
1130   0xdd17479e,
1131   0xad5a8333,
1132 };
1133
1134 typedef unsigned char uchar;
1135
1136 __device__ static void whirlpool_transform (const u32 w[16], u32 dgst[16], u32 s_Ch[8][256], u32 s_Cl[8][256])
1137 {
1138   u32 Kh[8];
1139   u32 Kl[8];
1140
1141   Kh[0] = dgst[ 0];
1142   Kl[0] = dgst[ 1];
1143   Kh[1] = dgst[ 2];
1144   Kl[1] = dgst[ 3];
1145   Kh[2] = dgst[ 4];
1146   Kl[2] = dgst[ 5];
1147   Kh[3] = dgst[ 6];
1148   Kl[3] = dgst[ 7];
1149   Kh[4] = dgst[ 8];
1150   Kl[4] = dgst[ 9];
1151   Kh[5] = dgst[10];
1152   Kl[5] = dgst[11];
1153   Kh[6] = dgst[12];
1154   Kl[6] = dgst[13];
1155   Kh[7] = dgst[14];
1156   Kl[7] = dgst[15];
1157
1158   u32 stateh[8];
1159   u32 statel[8];
1160
1161   stateh[0] = w[ 0] ^ Kh[0];
1162   statel[0] = w[ 1] ^ Kl[0];
1163   stateh[1] = w[ 2] ^ Kh[1];
1164   statel[1] = w[ 3] ^ Kl[1];
1165   stateh[2] = w[ 4] ^ Kh[2];
1166   statel[2] = w[ 5] ^ Kl[2];
1167   stateh[3] = w[ 6] ^ Kh[3];
1168   statel[3] = w[ 7] ^ Kl[3];
1169   stateh[4] = w[ 8] ^ Kh[4];
1170   statel[4] = w[ 9] ^ Kl[4];
1171   stateh[5] = w[10] ^ Kh[5];
1172   statel[5] = w[11] ^ Kl[5];
1173   stateh[6] = w[12] ^ Kh[6];
1174   statel[6] = w[13] ^ Kl[6];
1175   stateh[7] = w[14] ^ Kh[7];
1176   statel[7] = w[15] ^ Kl[7];
1177
1178   u32 r;
1179
1180   for (r = 1; r <= R; r++)
1181   {
1182     u32 Lh[8];
1183     u32 Ll[8];
1184
1185     u32 i;
1186
1187     #pragma unroll 8
1188     for (i = 0; i < 8; i++)
1189     {
1190       const u8 Lp0 = Kh[(i + 8) & 7] >> 24;
1191       const u8 Lp1 = Kh[(i + 7) & 7] >> 16;
1192       const u8 Lp2 = Kh[(i + 6) & 7] >>  8;
1193       const u8 Lp3 = Kh[(i + 5) & 7] >>  0;
1194       const u8 Lp4 = Kl[(i + 4) & 7] >> 24;
1195       const u8 Lp5 = Kl[(i + 3) & 7] >> 16;
1196       const u8 Lp6 = Kl[(i + 2) & 7] >>  8;
1197       const u8 Lp7 = Kl[(i + 1) & 7] >>  0;
1198
1199       Lh[i] = BOX (s_Ch, 0, Lp0 & 0xff)
1200             ^ BOX (s_Ch, 1, Lp1 & 0xff)
1201             ^ BOX (s_Ch, 2, Lp2 & 0xff)
1202             ^ BOX (s_Ch, 3, Lp3 & 0xff)
1203             ^ BOX (s_Ch, 4, Lp4 & 0xff)
1204             ^ BOX (s_Ch, 5, Lp5 & 0xff)
1205             ^ BOX (s_Ch, 6, Lp6 & 0xff)
1206             ^ BOX (s_Ch, 7, Lp7 & 0xff);
1207
1208       Ll[i] = BOX (s_Cl, 0, Lp0 & 0xff)
1209             ^ BOX (s_Cl, 1, Lp1 & 0xff)
1210             ^ BOX (s_Cl, 2, Lp2 & 0xff)
1211             ^ BOX (s_Cl, 3, Lp3 & 0xff)
1212             ^ BOX (s_Cl, 4, Lp4 & 0xff)
1213             ^ BOX (s_Cl, 5, Lp5 & 0xff)
1214             ^ BOX (s_Cl, 6, Lp6 & 0xff)
1215             ^ BOX (s_Cl, 7, Lp7 & 0xff);
1216     }
1217
1218     Kh[0] = Lh[0] ^ rch[r];
1219     Kl[0] = Ll[0] ^ rcl[r];
1220     Kh[1] = Lh[1];
1221     Kl[1] = Ll[1];
1222     Kh[2] = Lh[2];
1223     Kl[2] = Ll[2];
1224     Kh[3] = Lh[3];
1225     Kl[3] = Ll[3];
1226     Kh[4] = Lh[4];
1227     Kl[4] = Ll[4];
1228     Kh[5] = Lh[5];
1229     Kl[5] = Ll[5];
1230     Kh[6] = Lh[6];
1231     Kl[6] = Ll[6];
1232     Kh[7] = Lh[7];
1233     Kl[7] = Ll[7];
1234
1235     #pragma unroll 8
1236     for (i = 0; i < 8; i++)
1237     {
1238       const u8 Lp0 = stateh[(i + 8) & 7] >> 24;
1239       const u8 Lp1 = stateh[(i + 7) & 7] >> 16;
1240       const u8 Lp2 = stateh[(i + 6) & 7] >>  8;
1241       const u8 Lp3 = stateh[(i + 5) & 7] >>  0;
1242       const u8 Lp4 = statel[(i + 4) & 7] >> 24;
1243       const u8 Lp5 = statel[(i + 3) & 7] >> 16;
1244       const u8 Lp6 = statel[(i + 2) & 7] >>  8;
1245       const u8 Lp7 = statel[(i + 1) & 7] >>  0;
1246
1247       Lh[i] = BOX (s_Ch, 0, Lp0 & 0xff)
1248             ^ BOX (s_Ch, 1, Lp1 & 0xff)
1249             ^ BOX (s_Ch, 2, Lp2 & 0xff)
1250             ^ BOX (s_Ch, 3, Lp3 & 0xff)
1251             ^ BOX (s_Ch, 4, Lp4 & 0xff)
1252             ^ BOX (s_Ch, 5, Lp5 & 0xff)
1253             ^ BOX (s_Ch, 6, Lp6 & 0xff)
1254             ^ BOX (s_Ch, 7, Lp7 & 0xff);
1255
1256       Ll[i] = BOX (s_Cl, 0, Lp0 & 0xff)
1257             ^ BOX (s_Cl, 1, Lp1 & 0xff)
1258             ^ BOX (s_Cl, 2, Lp2 & 0xff)
1259             ^ BOX (s_Cl, 3, Lp3 & 0xff)
1260             ^ BOX (s_Cl, 4, Lp4 & 0xff)
1261             ^ BOX (s_Cl, 5, Lp5 & 0xff)
1262             ^ BOX (s_Cl, 6, Lp6 & 0xff)
1263             ^ BOX (s_Cl, 7, Lp7 & 0xff);
1264     }
1265
1266     stateh[0] = Lh[0] ^ Kh[0];
1267     statel[0] = Ll[0] ^ Kl[0];
1268     stateh[1] = Lh[1] ^ Kh[1];
1269     statel[1] = Ll[1] ^ Kl[1];
1270     stateh[2] = Lh[2] ^ Kh[2];
1271     statel[2] = Ll[2] ^ Kl[2];
1272     stateh[3] = Lh[3] ^ Kh[3];
1273     statel[3] = Ll[3] ^ Kl[3];
1274     stateh[4] = Lh[4] ^ Kh[4];
1275     statel[4] = Ll[4] ^ Kl[4];
1276     stateh[5] = Lh[5] ^ Kh[5];
1277     statel[5] = Ll[5] ^ Kl[5];
1278     stateh[6] = Lh[6] ^ Kh[6];
1279     statel[6] = Ll[6] ^ Kl[6];
1280     stateh[7] = Lh[7] ^ Kh[7];
1281     statel[7] = Ll[7] ^ Kl[7];
1282   }
1283
1284   dgst[ 0] ^= stateh[0] ^ w[ 0];
1285   dgst[ 1] ^= statel[0] ^ w[ 1];
1286   dgst[ 2] ^= stateh[1] ^ w[ 2];
1287   dgst[ 3] ^= statel[1] ^ w[ 3];
1288   dgst[ 4] ^= stateh[2] ^ w[ 4];
1289   dgst[ 5] ^= statel[2] ^ w[ 5];
1290   dgst[ 6] ^= stateh[3] ^ w[ 6];
1291   dgst[ 7] ^= statel[3] ^ w[ 7];
1292   dgst[ 8] ^= stateh[4] ^ w[ 8];
1293   dgst[ 9] ^= statel[4] ^ w[ 9];
1294   dgst[10] ^= stateh[5] ^ w[10];
1295   dgst[11] ^= statel[5] ^ w[11];
1296   dgst[12] ^= stateh[6] ^ w[12];
1297   dgst[13] ^= statel[6] ^ w[13];
1298   dgst[14] ^= stateh[7] ^ w[14];
1299   dgst[15] ^= statel[7] ^ w[15];
1300 }
1301
1302 __device__ static void hmac_run2 (const u32 w1[16], const u32 w2[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], u32 s_Ch[8][256], u32 s_Cl[8][256])
1303 {
1304   dgst[ 0] = ipad[ 0];
1305   dgst[ 1] = ipad[ 1];
1306   dgst[ 2] = ipad[ 2];
1307   dgst[ 3] = ipad[ 3];
1308   dgst[ 4] = ipad[ 4];
1309   dgst[ 5] = ipad[ 5];
1310   dgst[ 6] = ipad[ 6];
1311   dgst[ 7] = ipad[ 7];
1312   dgst[ 8] = ipad[ 8];
1313   dgst[ 9] = ipad[ 9];
1314   dgst[10] = ipad[10];
1315   dgst[11] = ipad[11];
1316   dgst[12] = ipad[12];
1317   dgst[13] = ipad[13];
1318   dgst[14] = ipad[14];
1319   dgst[15] = ipad[15];
1320
1321   whirlpool_transform (w1, dgst, s_Ch, s_Cl);
1322   whirlpool_transform (w2, dgst, s_Ch, s_Cl);
1323
1324   u32 w[16];
1325
1326   w[ 0] = dgst[ 0];
1327   w[ 1] = dgst[ 1];
1328   w[ 2] = dgst[ 2];
1329   w[ 3] = dgst[ 3];
1330   w[ 4] = dgst[ 4];
1331   w[ 5] = dgst[ 5];
1332   w[ 6] = dgst[ 6];
1333   w[ 7] = dgst[ 7];
1334   w[ 8] = dgst[ 8];
1335   w[ 9] = dgst[ 9];
1336   w[10] = dgst[10];
1337   w[11] = dgst[11];
1338   w[12] = dgst[12];
1339   w[13] = dgst[13];
1340   w[14] = dgst[14];
1341   w[15] = dgst[15];
1342
1343   dgst[ 0] = opad[ 0];
1344   dgst[ 1] = opad[ 1];
1345   dgst[ 2] = opad[ 2];
1346   dgst[ 3] = opad[ 3];
1347   dgst[ 4] = opad[ 4];
1348   dgst[ 5] = opad[ 5];
1349   dgst[ 6] = opad[ 6];
1350   dgst[ 7] = opad[ 7];
1351   dgst[ 8] = opad[ 8];
1352   dgst[ 9] = opad[ 9];
1353   dgst[10] = opad[10];
1354   dgst[11] = opad[11];
1355   dgst[12] = opad[12];
1356   dgst[13] = opad[13];
1357   dgst[14] = opad[14];
1358   dgst[15] = opad[15];
1359
1360   whirlpool_transform (w, dgst, s_Ch, s_Cl);
1361
1362   w[ 0] = 0x80000000;
1363   w[ 1] = 0;
1364   w[ 2] = 0;
1365   w[ 3] = 0;
1366   w[ 4] = 0;
1367   w[ 5] = 0;
1368   w[ 6] = 0;
1369   w[ 7] = 0;
1370   w[ 8] = 0;
1371   w[ 9] = 0;
1372   w[10] = 0;
1373   w[11] = 0;
1374   w[12] = 0;
1375   w[13] = 0;
1376   w[14] = 0;
1377   w[15] = (64 + 64) * 8;
1378
1379   whirlpool_transform (w, dgst, s_Ch, s_Cl);
1380 }
1381
1382 __device__ static void hmac_init (u32 w[16], u32 ipad[16], u32 opad[16], u32 s_Ch[8][256], u32 s_Cl[8][256])
1383 {
1384   w[ 0] ^= 0x36363636;
1385   w[ 1] ^= 0x36363636;
1386   w[ 2] ^= 0x36363636;
1387   w[ 3] ^= 0x36363636;
1388   w[ 4] ^= 0x36363636;
1389   w[ 5] ^= 0x36363636;
1390   w[ 6] ^= 0x36363636;
1391   w[ 7] ^= 0x36363636;
1392   w[ 8] ^= 0x36363636;
1393   w[ 9] ^= 0x36363636;
1394   w[10] ^= 0x36363636;
1395   w[11] ^= 0x36363636;
1396   w[12] ^= 0x36363636;
1397   w[13] ^= 0x36363636;
1398   w[14] ^= 0x36363636;
1399   w[15] ^= 0x36363636;
1400
1401   ipad[ 0] = 0;
1402   ipad[ 1] = 0;
1403   ipad[ 2] = 0;
1404   ipad[ 3] = 0;
1405   ipad[ 4] = 0;
1406   ipad[ 5] = 0;
1407   ipad[ 6] = 0;
1408   ipad[ 7] = 0;
1409   ipad[ 8] = 0;
1410   ipad[ 9] = 0;
1411   ipad[10] = 0;
1412   ipad[11] = 0;
1413   ipad[12] = 0;
1414   ipad[13] = 0;
1415   ipad[14] = 0;
1416   ipad[15] = 0;
1417
1418   whirlpool_transform (w, ipad, s_Ch, s_Cl);
1419
1420   w[ 0] ^= 0x6a6a6a6a;
1421   w[ 1] ^= 0x6a6a6a6a;
1422   w[ 2] ^= 0x6a6a6a6a;
1423   w[ 3] ^= 0x6a6a6a6a;
1424   w[ 4] ^= 0x6a6a6a6a;
1425   w[ 5] ^= 0x6a6a6a6a;
1426   w[ 6] ^= 0x6a6a6a6a;
1427   w[ 7] ^= 0x6a6a6a6a;
1428   w[ 8] ^= 0x6a6a6a6a;
1429   w[ 9] ^= 0x6a6a6a6a;
1430   w[10] ^= 0x6a6a6a6a;
1431   w[11] ^= 0x6a6a6a6a;
1432   w[12] ^= 0x6a6a6a6a;
1433   w[13] ^= 0x6a6a6a6a;
1434   w[14] ^= 0x6a6a6a6a;
1435   w[15] ^= 0x6a6a6a6a;
1436
1437   opad[ 0] = 0;
1438   opad[ 1] = 0;
1439   opad[ 2] = 0;
1440   opad[ 3] = 0;
1441   opad[ 4] = 0;
1442   opad[ 5] = 0;
1443   opad[ 6] = 0;
1444   opad[ 7] = 0;
1445   opad[ 8] = 0;
1446   opad[ 9] = 0;
1447   opad[10] = 0;
1448   opad[11] = 0;
1449   opad[12] = 0;
1450   opad[13] = 0;
1451   opad[14] = 0;
1452   opad[15] = 0;
1453
1454   whirlpool_transform (w, opad, s_Ch, s_Cl);
1455 }
1456
1457 __device__ static u32 u8add (const u32 a, const u32 b)
1458 {
1459   const u32 a1 = (a >>  0) & 0xff;
1460   const u32 a2 = (a >>  8) & 0xff;
1461   const u32 a3 = (a >> 16) & 0xff;
1462   const u32 a4 = (a >> 24) & 0xff;
1463
1464   const u32 b1 = (b >>  0) & 0xff;
1465   const u32 b2 = (b >>  8) & 0xff;
1466   const u32 b3 = (b >> 16) & 0xff;
1467   const u32 b4 = (b >> 24) & 0xff;
1468
1469   const u32 r1 = (a1 + b1) & 0xff;
1470   const u32 r2 = (a2 + b2) & 0xff;
1471   const u32 r3 = (a3 + b3) & 0xff;
1472   const u32 r4 = (a4 + b4) & 0xff;
1473
1474   const u32 r = r1 <<  0
1475                | r2 <<  8
1476                | r3 << 16
1477                | r4 << 24;
1478
1479   return r;
1480 }
1481
1482 extern "C" __global__ void __launch_bounds__ (256, 1) m06232_init (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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 tc_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)
1483 {
1484   /**
1485    * base
1486    */
1487
1488   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1489
1490   u32x w0[4];
1491
1492   w0[0] = pws[gid].i[ 0];
1493   w0[1] = pws[gid].i[ 1];
1494   w0[2] = pws[gid].i[ 2];
1495   w0[3] = pws[gid].i[ 3];
1496
1497   u32x w1[4];
1498
1499   w1[0] = pws[gid].i[ 4];
1500   w1[1] = pws[gid].i[ 5];
1501   w1[2] = pws[gid].i[ 6];
1502   w1[3] = pws[gid].i[ 7];
1503
1504   u32x w2[4];
1505
1506   w2[0] = pws[gid].i[ 8];
1507   w2[1] = pws[gid].i[ 9];
1508   w2[2] = pws[gid].i[10];
1509   w2[3] = pws[gid].i[11];
1510
1511   u32x w3[4];
1512
1513   w3[0] = pws[gid].i[12];
1514   w3[1] = pws[gid].i[13];
1515   w3[2] = pws[gid].i[14];
1516   w3[3] = pws[gid].i[15];
1517
1518   /**
1519    * keyfile
1520    */
1521
1522   w0[0] = u8add (w0[0], esalt_bufs[salt_pos].keyfile_buf[ 0]);
1523   w0[1] = u8add (w0[1], esalt_bufs[salt_pos].keyfile_buf[ 1]);
1524   w0[2] = u8add (w0[2], esalt_bufs[salt_pos].keyfile_buf[ 2]);
1525   w0[3] = u8add (w0[3], esalt_bufs[salt_pos].keyfile_buf[ 3]);
1526   w1[0] = u8add (w1[0], esalt_bufs[salt_pos].keyfile_buf[ 4]);
1527   w1[1] = u8add (w1[1], esalt_bufs[salt_pos].keyfile_buf[ 5]);
1528   w1[2] = u8add (w1[2], esalt_bufs[salt_pos].keyfile_buf[ 6]);
1529   w1[3] = u8add (w1[3], esalt_bufs[salt_pos].keyfile_buf[ 7]);
1530   w2[0] = u8add (w2[0], esalt_bufs[salt_pos].keyfile_buf[ 8]);
1531   w2[1] = u8add (w2[1], esalt_bufs[salt_pos].keyfile_buf[ 9]);
1532   w2[2] = u8add (w2[2], esalt_bufs[salt_pos].keyfile_buf[10]);
1533   w2[3] = u8add (w2[3], esalt_bufs[salt_pos].keyfile_buf[11]);
1534   w3[0] = u8add (w3[0], esalt_bufs[salt_pos].keyfile_buf[12]);
1535   w3[1] = u8add (w3[1], esalt_bufs[salt_pos].keyfile_buf[13]);
1536   w3[2] = u8add (w3[2], esalt_bufs[salt_pos].keyfile_buf[14]);
1537   w3[3] = u8add (w3[3], esalt_bufs[salt_pos].keyfile_buf[15]);
1538
1539   /**
1540    * shared mem
1541    */
1542
1543   __shared__ u32 s_Ch[8][256];
1544   __shared__ u32 s_Cl[8][256];
1545
1546   const u32 lid = threadIdx.x;
1547
1548   #pragma unroll 8
1549   for (u32 i = 0; i < 8; i++)
1550   {
1551     s_Ch[i][lid] = Ch[i][lid];
1552     s_Cl[i][lid] = Cl[i][lid];
1553   }
1554
1555   __syncthreads ();
1556
1557   if (gid >= gid_max) return;
1558
1559   /**
1560    * salt
1561    */
1562
1563   u32 salt_buf1[16];
1564
1565   salt_buf1[ 0] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 0]);
1566   salt_buf1[ 1] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 1]);
1567   salt_buf1[ 2] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 2]);
1568   salt_buf1[ 3] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 3]);
1569   salt_buf1[ 4] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 4]);
1570   salt_buf1[ 5] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 5]);
1571   salt_buf1[ 6] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 6]);
1572   salt_buf1[ 7] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 7]);
1573   salt_buf1[ 8] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 8]);
1574   salt_buf1[ 9] = swap_workaround (esalt_bufs[salt_pos].salt_buf[ 9]);
1575   salt_buf1[10] = swap_workaround (esalt_bufs[salt_pos].salt_buf[10]);
1576   salt_buf1[11] = swap_workaround (esalt_bufs[salt_pos].salt_buf[11]);
1577   salt_buf1[12] = swap_workaround (esalt_bufs[salt_pos].salt_buf[12]);
1578   salt_buf1[13] = swap_workaround (esalt_bufs[salt_pos].salt_buf[13]);
1579   salt_buf1[14] = swap_workaround (esalt_bufs[salt_pos].salt_buf[14]);
1580   salt_buf1[15] = swap_workaround (esalt_bufs[salt_pos].salt_buf[15]);
1581
1582   u32 salt_buf2[16];
1583
1584   salt_buf2[ 0] = 0;
1585   salt_buf2[ 1] = 0x80000000;
1586   salt_buf2[ 2] = 0;
1587   salt_buf2[ 3] = 0;
1588   salt_buf2[ 4] = 0;
1589   salt_buf2[ 5] = 0;
1590   salt_buf2[ 6] = 0;
1591   salt_buf2[ 7] = 0;
1592   salt_buf2[ 8] = 0;
1593   salt_buf2[ 9] = 0;
1594   salt_buf2[10] = 0;
1595   salt_buf2[11] = 0;
1596   salt_buf2[12] = 0;
1597   salt_buf2[13] = 0;
1598   salt_buf2[14] = 0;
1599   salt_buf2[15] = (64 + 64 + 4) * 8;
1600
1601   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
1602
1603   u32 w[16];
1604
1605   w[ 0] = swap_workaround (w0[0]);
1606   w[ 1] = swap_workaround (w0[1]);
1607   w[ 2] = swap_workaround (w0[2]);
1608   w[ 3] = swap_workaround (w0[3]);
1609   w[ 4] = swap_workaround (w1[0]);
1610   w[ 5] = swap_workaround (w1[1]);
1611   w[ 6] = swap_workaround (w1[2]);
1612   w[ 7] = swap_workaround (w1[3]);
1613   w[ 8] = swap_workaround (w2[0]);
1614   w[ 9] = swap_workaround (w2[1]);
1615   w[10] = swap_workaround (w2[2]);
1616   w[11] = swap_workaround (w2[3]);
1617   w[12] = swap_workaround (w3[0]);
1618   w[13] = swap_workaround (w3[1]);
1619   w[14] = swap_workaround (w3[2]);
1620   w[15] = swap_workaround (w3[3]);
1621
1622   u32 ipad[16];
1623   u32 opad[16];
1624
1625   hmac_init (w, ipad, opad, s_Ch, s_Cl);
1626
1627   tmps[gid].ipad[ 0] = ipad[ 0];
1628   tmps[gid].ipad[ 1] = ipad[ 1];
1629   tmps[gid].ipad[ 2] = ipad[ 2];
1630   tmps[gid].ipad[ 3] = ipad[ 3];
1631   tmps[gid].ipad[ 4] = ipad[ 4];
1632   tmps[gid].ipad[ 5] = ipad[ 5];
1633   tmps[gid].ipad[ 6] = ipad[ 6];
1634   tmps[gid].ipad[ 7] = ipad[ 7];
1635   tmps[gid].ipad[ 8] = ipad[ 8];
1636   tmps[gid].ipad[ 9] = ipad[ 9];
1637   tmps[gid].ipad[10] = ipad[10];
1638   tmps[gid].ipad[11] = ipad[11];
1639   tmps[gid].ipad[12] = ipad[12];
1640   tmps[gid].ipad[13] = ipad[13];
1641   tmps[gid].ipad[14] = ipad[14];
1642   tmps[gid].ipad[15] = ipad[15];
1643
1644   tmps[gid].opad[ 0] = opad[ 0];
1645   tmps[gid].opad[ 1] = opad[ 1];
1646   tmps[gid].opad[ 2] = opad[ 2];
1647   tmps[gid].opad[ 3] = opad[ 3];
1648   tmps[gid].opad[ 4] = opad[ 4];
1649   tmps[gid].opad[ 5] = opad[ 5];
1650   tmps[gid].opad[ 6] = opad[ 6];
1651   tmps[gid].opad[ 7] = opad[ 7];
1652   tmps[gid].opad[ 8] = opad[ 8];
1653   tmps[gid].opad[ 9] = opad[ 9];
1654   tmps[gid].opad[10] = opad[10];
1655   tmps[gid].opad[11] = opad[11];
1656   tmps[gid].opad[12] = opad[12];
1657   tmps[gid].opad[13] = opad[13];
1658   tmps[gid].opad[14] = opad[14];
1659   tmps[gid].opad[15] = opad[15];
1660
1661   for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 16, j += 1)
1662   {
1663     salt_buf2[0] = j;
1664
1665     u32 dgst[16];
1666
1667     hmac_run2 (salt_buf1, salt_buf2, ipad, opad, dgst, s_Ch, s_Cl);
1668
1669     tmps[gid].dgst[i +  0] = dgst[ 0];
1670     tmps[gid].dgst[i +  1] = dgst[ 1];
1671     tmps[gid].dgst[i +  2] = dgst[ 2];
1672     tmps[gid].dgst[i +  3] = dgst[ 3];
1673     tmps[gid].dgst[i +  4] = dgst[ 4];
1674     tmps[gid].dgst[i +  5] = dgst[ 5];
1675     tmps[gid].dgst[i +  6] = dgst[ 6];
1676     tmps[gid].dgst[i +  7] = dgst[ 7];
1677     tmps[gid].dgst[i +  8] = dgst[ 8];
1678     tmps[gid].dgst[i +  9] = dgst[ 9];
1679     tmps[gid].dgst[i + 10] = dgst[10];
1680     tmps[gid].dgst[i + 11] = dgst[11];
1681     tmps[gid].dgst[i + 12] = dgst[12];
1682     tmps[gid].dgst[i + 13] = dgst[13];
1683     tmps[gid].dgst[i + 14] = dgst[14];
1684     tmps[gid].dgst[i + 15] = dgst[15];
1685
1686     tmps[gid].out[i +  0] = dgst[ 0];
1687     tmps[gid].out[i +  1] = dgst[ 1];
1688     tmps[gid].out[i +  2] = dgst[ 2];
1689     tmps[gid].out[i +  3] = dgst[ 3];
1690     tmps[gid].out[i +  4] = dgst[ 4];
1691     tmps[gid].out[i +  5] = dgst[ 5];
1692     tmps[gid].out[i +  6] = dgst[ 6];
1693     tmps[gid].out[i +  7] = dgst[ 7];
1694     tmps[gid].out[i +  8] = dgst[ 8];
1695     tmps[gid].out[i +  9] = dgst[ 9];
1696     tmps[gid].out[i + 10] = dgst[10];
1697     tmps[gid].out[i + 11] = dgst[11];
1698     tmps[gid].out[i + 12] = dgst[12];
1699     tmps[gid].out[i + 13] = dgst[13];
1700     tmps[gid].out[i + 14] = dgst[14];
1701     tmps[gid].out[i + 15] = dgst[15];
1702   }
1703 }
1704
1705 extern "C" __global__ void __launch_bounds__ (256, 1) m06232_loop (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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 tc_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)
1706 {
1707   const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen;
1708
1709   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1710
1711   __shared__ u32 s_Ch[8][256];
1712   __shared__ u32 s_Cl[8][256];
1713
1714   const u32 lid = threadIdx.x;
1715
1716   #pragma unroll 8
1717   for (u32 i = 0; i < 8; i++)
1718   {
1719     s_Ch[i][lid] = Ch[i][lid];
1720     s_Cl[i][lid] = Cl[i][lid];
1721   }
1722
1723   __syncthreads ();
1724
1725   if (gid >= gid_max) return;
1726
1727   u32 ipad[16];
1728
1729   ipad[ 0] = tmps[gid].ipad[ 0];
1730   ipad[ 1] = tmps[gid].ipad[ 1],
1731   ipad[ 2] = tmps[gid].ipad[ 2];
1732   ipad[ 3] = tmps[gid].ipad[ 3];
1733   ipad[ 4] = tmps[gid].ipad[ 4];
1734   ipad[ 5] = tmps[gid].ipad[ 5];
1735   ipad[ 6] = tmps[gid].ipad[ 6],
1736   ipad[ 7] = tmps[gid].ipad[ 7];
1737   ipad[ 8] = tmps[gid].ipad[ 8];
1738   ipad[ 9] = tmps[gid].ipad[ 9];
1739   ipad[10] = tmps[gid].ipad[10];
1740   ipad[11] = tmps[gid].ipad[11],
1741   ipad[12] = tmps[gid].ipad[12];
1742   ipad[13] = tmps[gid].ipad[13];
1743   ipad[14] = tmps[gid].ipad[14];
1744   ipad[15] = tmps[gid].ipad[15];
1745
1746   u32 opad[16];
1747
1748   opad[ 0] = tmps[gid].opad[ 0];
1749   opad[ 1] = tmps[gid].opad[ 1],
1750   opad[ 2] = tmps[gid].opad[ 2];
1751   opad[ 3] = tmps[gid].opad[ 3];
1752   opad[ 4] = tmps[gid].opad[ 4];
1753   opad[ 5] = tmps[gid].opad[ 5];
1754   opad[ 6] = tmps[gid].opad[ 6],
1755   opad[ 7] = tmps[gid].opad[ 7];
1756   opad[ 8] = tmps[gid].opad[ 8];
1757   opad[ 9] = tmps[gid].opad[ 9];
1758   opad[10] = tmps[gid].opad[10];
1759   opad[11] = tmps[gid].opad[11],
1760   opad[12] = tmps[gid].opad[12];
1761   opad[13] = tmps[gid].opad[13];
1762   opad[14] = tmps[gid].opad[14];
1763   opad[15] = tmps[gid].opad[15];
1764
1765   for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 16)
1766   {
1767     u32 dgst[16];
1768
1769     dgst[ 0] = tmps[gid].dgst[i +  0];
1770     dgst[ 1] = tmps[gid].dgst[i +  1];
1771     dgst[ 2] = tmps[gid].dgst[i +  2];
1772     dgst[ 3] = tmps[gid].dgst[i +  3];
1773     dgst[ 4] = tmps[gid].dgst[i +  4];
1774     dgst[ 5] = tmps[gid].dgst[i +  5];
1775     dgst[ 6] = tmps[gid].dgst[i +  6];
1776     dgst[ 7] = tmps[gid].dgst[i +  7];
1777     dgst[ 8] = tmps[gid].dgst[i +  8];
1778     dgst[ 9] = tmps[gid].dgst[i +  9];
1779     dgst[10] = tmps[gid].dgst[i + 10];
1780     dgst[11] = tmps[gid].dgst[i + 11];
1781     dgst[12] = tmps[gid].dgst[i + 12];
1782     dgst[13] = tmps[gid].dgst[i + 13];
1783     dgst[14] = tmps[gid].dgst[i + 14];
1784     dgst[15] = tmps[gid].dgst[i + 15];
1785
1786     u32 out[16];
1787
1788     out[ 0] = tmps[gid].out[i +  0];
1789     out[ 1] = tmps[gid].out[i +  1];
1790     out[ 2] = tmps[gid].out[i +  2];
1791     out[ 3] = tmps[gid].out[i +  3];
1792     out[ 4] = tmps[gid].out[i +  4];
1793     out[ 5] = tmps[gid].out[i +  5];
1794     out[ 6] = tmps[gid].out[i +  6];
1795     out[ 7] = tmps[gid].out[i +  7];
1796     out[ 8] = tmps[gid].out[i +  8];
1797     out[ 9] = tmps[gid].out[i +  9];
1798     out[10] = tmps[gid].out[i + 10];
1799     out[11] = tmps[gid].out[i + 11];
1800     out[12] = tmps[gid].out[i + 12];
1801     out[13] = tmps[gid].out[i + 13];
1802     out[14] = tmps[gid].out[i + 14];
1803     out[15] = tmps[gid].out[i + 15];
1804
1805     for (u32 j = 0; j < loop_cnt; j++)
1806     {
1807       u32 w1[16];
1808
1809       w1[ 0] = dgst[ 0];
1810       w1[ 1] = dgst[ 1];
1811       w1[ 2] = dgst[ 2];
1812       w1[ 3] = dgst[ 3];
1813       w1[ 4] = dgst[ 4];
1814       w1[ 5] = dgst[ 5];
1815       w1[ 6] = dgst[ 6];
1816       w1[ 7] = dgst[ 7];
1817       w1[ 8] = dgst[ 8];
1818       w1[ 9] = dgst[ 9];
1819       w1[10] = dgst[10];
1820       w1[11] = dgst[11];
1821       w1[12] = dgst[12];
1822       w1[13] = dgst[13];
1823       w1[14] = dgst[14];
1824       w1[15] = dgst[15];
1825
1826       u32 w2[16];
1827
1828       w2[ 0] = 0x80000000;
1829       w2[ 1] = 0;
1830       w2[ 2] = 0;
1831       w2[ 3] = 0;
1832       w2[ 4] = 0;
1833       w2[ 5] = 0;
1834       w2[ 6] = 0;
1835       w2[ 7] = 0;
1836       w2[ 8] = 0;
1837       w2[ 9] = 0;
1838       w2[10] = 0;
1839       w2[11] = 0;
1840       w2[12] = 0;
1841       w2[13] = 0;
1842       w2[14] = 0;
1843       w2[15] = (64 + 64) * 8;
1844
1845       hmac_run2 (w1, w2, ipad, opad, dgst, s_Ch, s_Cl);
1846
1847       out[ 0] ^= dgst[ 0];
1848       out[ 1] ^= dgst[ 1];
1849       out[ 2] ^= dgst[ 2];
1850       out[ 3] ^= dgst[ 3];
1851       out[ 4] ^= dgst[ 4];
1852       out[ 5] ^= dgst[ 5];
1853       out[ 6] ^= dgst[ 6];
1854       out[ 7] ^= dgst[ 7];
1855       out[ 8] ^= dgst[ 8];
1856       out[ 9] ^= dgst[ 9];
1857       out[10] ^= dgst[10];
1858       out[11] ^= dgst[11];
1859       out[12] ^= dgst[12];
1860       out[13] ^= dgst[13];
1861       out[14] ^= dgst[14];
1862       out[15] ^= dgst[15];
1863     }
1864
1865     tmps[gid].dgst[i +  0] = dgst[ 0];
1866     tmps[gid].dgst[i +  1] = dgst[ 1];
1867     tmps[gid].dgst[i +  2] = dgst[ 2];
1868     tmps[gid].dgst[i +  3] = dgst[ 3];
1869     tmps[gid].dgst[i +  4] = dgst[ 4];
1870     tmps[gid].dgst[i +  5] = dgst[ 5];
1871     tmps[gid].dgst[i +  6] = dgst[ 6];
1872     tmps[gid].dgst[i +  7] = dgst[ 7];
1873     tmps[gid].dgst[i +  8] = dgst[ 8];
1874     tmps[gid].dgst[i +  9] = dgst[ 9];
1875     tmps[gid].dgst[i + 10] = dgst[10];
1876     tmps[gid].dgst[i + 11] = dgst[11];
1877     tmps[gid].dgst[i + 12] = dgst[12];
1878     tmps[gid].dgst[i + 13] = dgst[13];
1879     tmps[gid].dgst[i + 14] = dgst[14];
1880     tmps[gid].dgst[i + 15] = dgst[15];
1881
1882     tmps[gid].out[i +  0] = out[ 0];
1883     tmps[gid].out[i +  1] = out[ 1];
1884     tmps[gid].out[i +  2] = out[ 2];
1885     tmps[gid].out[i +  3] = out[ 3];
1886     tmps[gid].out[i +  4] = out[ 4];
1887     tmps[gid].out[i +  5] = out[ 5];
1888     tmps[gid].out[i +  6] = out[ 6];
1889     tmps[gid].out[i +  7] = out[ 7];
1890     tmps[gid].out[i +  8] = out[ 8];
1891     tmps[gid].out[i +  9] = out[ 9];
1892     tmps[gid].out[i + 10] = out[10];
1893     tmps[gid].out[i + 11] = out[11];
1894     tmps[gid].out[i + 12] = out[12];
1895     tmps[gid].out[i + 13] = out[13];
1896     tmps[gid].out[i + 14] = out[14];
1897     tmps[gid].out[i + 15] = out[15];
1898   }
1899 }
1900
1901 extern "C" __global__ void __launch_bounds__ (256, 1) m06232_comp (const pw_t *pws, const gpu_rule_t *rules_buf, const comb_t *combs_buf, const bf_t *bfs_buf, tc_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 tc_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)
1902 {
1903   /**
1904    * base
1905    */
1906
1907   const u32 gid = (blockIdx.x * blockDim.x) + threadIdx.x;
1908   const u32 lid = threadIdx.x;
1909
1910   if (gid >= gid_max) return;
1911
1912   u32 ukey1[8];
1913
1914   ukey1[0] = swap_workaround (tmps[gid].out[ 0]);
1915   ukey1[1] = swap_workaround (tmps[gid].out[ 1]);
1916   ukey1[2] = swap_workaround (tmps[gid].out[ 2]);
1917   ukey1[3] = swap_workaround (tmps[gid].out[ 3]);
1918   ukey1[4] = swap_workaround (tmps[gid].out[ 4]);
1919   ukey1[5] = swap_workaround (tmps[gid].out[ 5]);
1920   ukey1[6] = swap_workaround (tmps[gid].out[ 6]);
1921   ukey1[7] = swap_workaround (tmps[gid].out[ 7]);
1922
1923   u32 ukey2[8];
1924
1925   ukey2[0] = swap_workaround (tmps[gid].out[ 8]);
1926   ukey2[1] = swap_workaround (tmps[gid].out[ 9]);
1927   ukey2[2] = swap_workaround (tmps[gid].out[10]);
1928   ukey2[3] = swap_workaround (tmps[gid].out[11]);
1929   ukey2[4] = swap_workaround (tmps[gid].out[12]);
1930   ukey2[5] = swap_workaround (tmps[gid].out[13]);
1931   ukey2[6] = swap_workaround (tmps[gid].out[14]);
1932   ukey2[7] = swap_workaround (tmps[gid].out[15]);
1933
1934   u32 data[4];
1935
1936   data[0] = esalt_bufs[0].data_buf[0];
1937   data[1] = esalt_bufs[0].data_buf[1];
1938   data[2] = esalt_bufs[0].data_buf[2];
1939   data[3] = esalt_bufs[0].data_buf[3];
1940
1941   u32 tmp[4];
1942
1943   {
1944     tmp[0] = data[0];
1945     tmp[1] = data[1];
1946     tmp[2] = data[2];
1947     tmp[3] = data[3];
1948
1949     aes256_decrypt_xts (ukey1, ukey2, tmp, tmp);
1950
1951     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
1952     {
1953       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
1954
1955       d_return_buf[lid] = 1;
1956     }
1957   }
1958
1959   {
1960     tmp[0] = data[0];
1961     tmp[1] = data[1];
1962     tmp[2] = data[2];
1963     tmp[3] = data[3];
1964
1965     serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp);
1966
1967     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
1968     {
1969       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
1970
1971       d_return_buf[lid] = 1;
1972     }
1973   }
1974
1975   {
1976     tmp[0] = data[0];
1977     tmp[1] = data[1];
1978     tmp[2] = data[2];
1979     tmp[3] = data[3];
1980
1981     twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp);
1982
1983     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
1984     {
1985       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
1986
1987       d_return_buf[lid] = 1;
1988     }
1989   }
1990
1991   u32 ukey3[8];
1992
1993   ukey3[0] = swap_workaround (tmps[gid].out[16]);
1994   ukey3[1] = swap_workaround (tmps[gid].out[17]);
1995   ukey3[2] = swap_workaround (tmps[gid].out[18]);
1996   ukey3[3] = swap_workaround (tmps[gid].out[19]);
1997   ukey3[4] = swap_workaround (tmps[gid].out[20]);
1998   ukey3[5] = swap_workaround (tmps[gid].out[21]);
1999   ukey3[6] = swap_workaround (tmps[gid].out[22]);
2000   ukey3[7] = swap_workaround (tmps[gid].out[23]);
2001
2002   u32 ukey4[8];
2003
2004   ukey4[0] = swap_workaround (tmps[gid].out[24]);
2005   ukey4[1] = swap_workaround (tmps[gid].out[25]);
2006   ukey4[2] = swap_workaround (tmps[gid].out[26]);
2007   ukey4[3] = swap_workaround (tmps[gid].out[27]);
2008   ukey4[4] = swap_workaround (tmps[gid].out[28]);
2009   ukey4[5] = swap_workaround (tmps[gid].out[29]);
2010   ukey4[6] = swap_workaround (tmps[gid].out[30]);
2011   ukey4[7] = swap_workaround (tmps[gid].out[31]);
2012
2013   {
2014     tmp[0] = data[0];
2015     tmp[1] = data[1];
2016     tmp[2] = data[2];
2017     tmp[3] = data[3];
2018
2019     aes256_decrypt_xts     (ukey2, ukey4, tmp, tmp);
2020     twofish256_decrypt_xts (ukey1, ukey3, tmp, tmp);
2021
2022     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
2023     {
2024       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
2025
2026       d_return_buf[lid] = 1;
2027     }
2028   }
2029
2030   {
2031     tmp[0] = data[0];
2032     tmp[1] = data[1];
2033     tmp[2] = data[2];
2034     tmp[3] = data[3];
2035
2036     serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp);
2037     aes256_decrypt_xts     (ukey1, ukey3, tmp, tmp);
2038
2039     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
2040     {
2041       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
2042
2043       d_return_buf[lid] = 1;
2044     }
2045   }
2046
2047   {
2048     tmp[0] = data[0];
2049     tmp[1] = data[1];
2050     tmp[2] = data[2];
2051     tmp[3] = data[3];
2052
2053     twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp);
2054     serpent256_decrypt_xts (ukey1, ukey3, tmp, tmp);
2055
2056     if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5)))
2057     {
2058       mark_hash_s0 (plains_buf, hashes_shown, 0, gid, 0);
2059
2060       d_return_buf[lid] = 1;
2061     }
2062   }
2063 }