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