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