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