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 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)
727 u32 temp = rek[j + 7];
729 rek[j + 8] = rek[j + 0]
730 ^ (s_te2[(temp >> 16) & 0xff] & 0xff000000)
731 ^ (s_te3[(temp >> 8) & 0xff] & 0x00ff0000)
732 ^ (s_te0[(temp >> 0) & 0xff] & 0x0000ff00)
733 ^ (s_te1[(temp >> 24) & 0xff] & 0x000000ff)
736 rek[j + 9] = rek[j + 1] ^ rek[j + 8];
737 rek[j + 10] = rek[j + 2] ^ rek[j + 9];
738 rek[j + 11] = rek[j + 3] ^ rek[j + 10];
748 rek[j + 12] = rek[j + 4]
749 ^ (s_te2[(temp >> 24) & 0xff] & 0xff000000)
750 ^ (s_te3[(temp >> 16) & 0xff] & 0x00ff0000)
751 ^ (s_te0[(temp >> 8) & 0xff] & 0x0000ff00)
752 ^ (s_te1[(temp >> 0) & 0xff] & 0x000000ff);
754 rek[j + 13] = rek[j + 5] ^ rek[j + 12];
755 rek[j + 14] = rek[j + 6] ^ rek[j + 13];
756 rek[j + 15] = rek[j + 7] ^ rek[j + 14];
762 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)
764 for (u32 i = 0, j = 56; i < j; i += 4, j -= 4)
768 temp = rdk[i + 0]; rdk[i + 0] = rdk[j + 0]; rdk[j + 0] = temp;
769 temp = rdk[i + 1]; rdk[i + 1] = rdk[j + 1]; rdk[j + 1] = temp;
770 temp = rdk[i + 2]; rdk[i + 2] = rdk[j + 2]; rdk[j + 2] = temp;
771 temp = rdk[i + 3]; rdk[i + 3] = rdk[j + 3]; rdk[j + 3] = temp;
774 for (u32 i = 1, j = 4; i < 14; i += 1, j += 4)
777 s_td0[s_te1[(rdk[j + 0] >> 24) & 0xff] & 0xff] ^
778 s_td1[s_te1[(rdk[j + 0] >> 16) & 0xff] & 0xff] ^
779 s_td2[s_te1[(rdk[j + 0] >> 8) & 0xff] & 0xff] ^
780 s_td3[s_te1[(rdk[j + 0] >> 0) & 0xff] & 0xff];
783 s_td0[s_te1[(rdk[j + 1] >> 24) & 0xff] & 0xff] ^
784 s_td1[s_te1[(rdk[j + 1] >> 16) & 0xff] & 0xff] ^
785 s_td2[s_te1[(rdk[j + 1] >> 8) & 0xff] & 0xff] ^
786 s_td3[s_te1[(rdk[j + 1] >> 0) & 0xff] & 0xff];
789 s_td0[s_te1[(rdk[j + 2] >> 24) & 0xff] & 0xff] ^
790 s_td1[s_te1[(rdk[j + 2] >> 16) & 0xff] & 0xff] ^
791 s_td2[s_te1[(rdk[j + 2] >> 8) & 0xff] & 0xff] ^
792 s_td3[s_te1[(rdk[j + 2] >> 0) & 0xff] & 0xff];
795 s_td0[s_te1[(rdk[j + 3] >> 24) & 0xff] & 0xff] ^
796 s_td1[s_te1[(rdk[j + 3] >> 16) & 0xff] & 0xff] ^
797 s_td2[s_te1[(rdk[j + 3] >> 8) & 0xff] & 0xff] ^
798 s_td3[s_te1[(rdk[j + 3] >> 0) & 0xff] & 0xff];
802 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)
804 u32 s0 = in[0] ^ rdk[0];
805 u32 s1 = in[1] ^ rdk[1];
806 u32 s2 = in[2] ^ rdk[2];
807 u32 s3 = in[3] ^ rdk[3];
814 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[ 4];
815 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[ 5];
816 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[ 6];
817 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[ 7];
818 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[ 8];
819 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[ 9];
820 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[10];
821 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[11];
822 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[12];
823 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[13];
824 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[14];
825 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[15];
826 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[16];
827 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[17];
828 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[18];
829 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[19];
830 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[20];
831 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[21];
832 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[22];
833 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[23];
834 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[24];
835 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[25];
836 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[26];
837 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[27];
838 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[28];
839 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[29];
840 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[30];
841 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[31];
842 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[32];
843 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[33];
844 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[34];
845 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[35];
846 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[36];
847 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[37];
848 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[38];
849 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[39];
850 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[40];
851 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[41];
852 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[42];
853 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[43];
854 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[44];
855 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[45];
856 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[46];
857 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[47];
858 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[48];
859 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[49];
860 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[50];
861 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[51];
862 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[52];
863 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[53];
864 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[54];
865 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[55];
867 out[0] = (s_td4[(t0 >> 24) & 0xff] & 0xff000000)
868 ^ (s_td4[(t3 >> 16) & 0xff] & 0x00ff0000)
869 ^ (s_td4[(t2 >> 8) & 0xff] & 0x0000ff00)
870 ^ (s_td4[(t1 >> 0) & 0xff] & 0x000000ff)
873 out[1] = (s_td4[(t1 >> 24) & 0xff] & 0xff000000)
874 ^ (s_td4[(t0 >> 16) & 0xff] & 0x00ff0000)
875 ^ (s_td4[(t3 >> 8) & 0xff] & 0x0000ff00)
876 ^ (s_td4[(t2 >> 0) & 0xff] & 0x000000ff)
879 out[2] = (s_td4[(t2 >> 24) & 0xff] & 0xff000000)
880 ^ (s_td4[(t1 >> 16) & 0xff] & 0x00ff0000)
881 ^ (s_td4[(t0 >> 8) & 0xff] & 0x0000ff00)
882 ^ (s_td4[(t3 >> 0) & 0xff] & 0x000000ff)
885 out[3] = (s_td4[(t3 >> 24) & 0xff] & 0xff000000)
886 ^ (s_td4[(t2 >> 16) & 0xff] & 0x00ff0000)
887 ^ (s_td4[(t1 >> 8) & 0xff] & 0x0000ff00)
888 ^ (s_td4[(t0 >> 0) & 0xff] & 0x000000ff)
892 __constant u32 k_sha256[64] =
894 SHA256C00, SHA256C01, SHA256C02, SHA256C03,
895 SHA256C04, SHA256C05, SHA256C06, SHA256C07,
896 SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
897 SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
898 SHA256C10, SHA256C11, SHA256C12, SHA256C13,
899 SHA256C14, SHA256C15, SHA256C16, SHA256C17,
900 SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
901 SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
902 SHA256C20, SHA256C21, SHA256C22, SHA256C23,
903 SHA256C24, SHA256C25, SHA256C26, SHA256C27,
904 SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
905 SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
906 SHA256C30, SHA256C31, SHA256C32, SHA256C33,
907 SHA256C34, SHA256C35, SHA256C36, SHA256C37,
908 SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
909 SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
912 void sha256_transform (const u32 w[16], u32 digest[8])
940 #define ROUND_EXPAND() \
942 w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
943 w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
944 w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
945 w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
946 w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
947 w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
948 w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
949 w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
950 w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
951 w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
952 wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
953 wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
954 wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
955 wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
956 we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
957 wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
960 #define ROUND_STEP(i) \
962 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
963 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
964 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
965 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
966 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
967 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
968 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
969 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
970 SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
971 SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
972 SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
973 SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
974 SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
975 SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
976 SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
977 SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
985 for (int i = 16; i < 64; i += 16)
987 ROUND_EXPAND (); ROUND_STEP (i);
1000 __constant u32 crc32tab[0x100] =
1002 0x00000000, 0x77073096, 0xee0e612c, 0x990951ba,
1003 0x076dc419, 0x706af48f, 0xe963a535, 0x9e6495a3,
1004 0x0edb8832, 0x79dcb8a4, 0xe0d5e91e, 0x97d2d988,
1005 0x09b64c2b, 0x7eb17cbd, 0xe7b82d07, 0x90bf1d91,
1006 0x1db71064, 0x6ab020f2, 0xf3b97148, 0x84be41de,
1007 0x1adad47d, 0x6ddde4eb, 0xf4d4b551, 0x83d385c7,
1008 0x136c9856, 0x646ba8c0, 0xfd62f97a, 0x8a65c9ec,
1009 0x14015c4f, 0x63066cd9, 0xfa0f3d63, 0x8d080df5,
1010 0x3b6e20c8, 0x4c69105e, 0xd56041e4, 0xa2677172,
1011 0x3c03e4d1, 0x4b04d447, 0xd20d85fd, 0xa50ab56b,
1012 0x35b5a8fa, 0x42b2986c, 0xdbbbc9d6, 0xacbcf940,
1013 0x32d86ce3, 0x45df5c75, 0xdcd60dcf, 0xabd13d59,
1014 0x26d930ac, 0x51de003a, 0xc8d75180, 0xbfd06116,
1015 0x21b4f4b5, 0x56b3c423, 0xcfba9599, 0xb8bda50f,
1016 0x2802b89e, 0x5f058808, 0xc60cd9b2, 0xb10be924,
1017 0x2f6f7c87, 0x58684c11, 0xc1611dab, 0xb6662d3d,
1018 0x76dc4190, 0x01db7106, 0x98d220bc, 0xefd5102a,
1019 0x71b18589, 0x06b6b51f, 0x9fbfe4a5, 0xe8b8d433,
1020 0x7807c9a2, 0x0f00f934, 0x9609a88e, 0xe10e9818,
1021 0x7f6a0dbb, 0x086d3d2d, 0x91646c97, 0xe6635c01,
1022 0x6b6b51f4, 0x1c6c6162, 0x856530d8, 0xf262004e,
1023 0x6c0695ed, 0x1b01a57b, 0x8208f4c1, 0xf50fc457,
1024 0x65b0d9c6, 0x12b7e950, 0x8bbeb8ea, 0xfcb9887c,
1025 0x62dd1ddf, 0x15da2d49, 0x8cd37cf3, 0xfbd44c65,
1026 0x4db26158, 0x3ab551ce, 0xa3bc0074, 0xd4bb30e2,
1027 0x4adfa541, 0x3dd895d7, 0xa4d1c46d, 0xd3d6f4fb,
1028 0x4369e96a, 0x346ed9fc, 0xad678846, 0xda60b8d0,
1029 0x44042d73, 0x33031de5, 0xaa0a4c5f, 0xdd0d7cc9,
1030 0x5005713c, 0x270241aa, 0xbe0b1010, 0xc90c2086,
1031 0x5768b525, 0x206f85b3, 0xb966d409, 0xce61e49f,
1032 0x5edef90e, 0x29d9c998, 0xb0d09822, 0xc7d7a8b4,
1033 0x59b33d17, 0x2eb40d81, 0xb7bd5c3b, 0xc0ba6cad,
1034 0xedb88320, 0x9abfb3b6, 0x03b6e20c, 0x74b1d29a,
1035 0xead54739, 0x9dd277af, 0x04db2615, 0x73dc1683,
1036 0xe3630b12, 0x94643b84, 0x0d6d6a3e, 0x7a6a5aa8,
1037 0xe40ecf0b, 0x9309ff9d, 0x0a00ae27, 0x7d079eb1,
1038 0xf00f9344, 0x8708a3d2, 0x1e01f268, 0x6906c2fe,
1039 0xf762575d, 0x806567cb, 0x196c3671, 0x6e6b06e7,
1040 0xfed41b76, 0x89d32be0, 0x10da7a5a, 0x67dd4acc,
1041 0xf9b9df6f, 0x8ebeeff9, 0x17b7be43, 0x60b08ed5,
1042 0xd6d6a3e8, 0xa1d1937e, 0x38d8c2c4, 0x4fdff252,
1043 0xd1bb67f1, 0xa6bc5767, 0x3fb506dd, 0x48b2364b,
1044 0xd80d2bda, 0xaf0a1b4c, 0x36034af6, 0x41047a60,
1045 0xdf60efc3, 0xa867df55, 0x316e8eef, 0x4669be79,
1046 0xcb61b38c, 0xbc66831a, 0x256fd2a0, 0x5268e236,
1047 0xcc0c7795, 0xbb0b4703, 0x220216b9, 0x5505262f,
1048 0xc5ba3bbe, 0xb2bd0b28, 0x2bb45a92, 0x5cb36a04,
1049 0xc2d7ffa7, 0xb5d0cf31, 0x2cd99e8b, 0x5bdeae1d,
1050 0x9b64c2b0, 0xec63f226, 0x756aa39c, 0x026d930a,
1051 0x9c0906a9, 0xeb0e363f, 0x72076785, 0x05005713,
1052 0x95bf4a82, 0xe2b87a14, 0x7bb12bae, 0x0cb61b38,
1053 0x92d28e9b, 0xe5d5be0d, 0x7cdcefb7, 0x0bdbdf21,
1054 0x86d3d2d4, 0xf1d4e242, 0x68ddb3f8, 0x1fda836e,
1055 0x81be16cd, 0xf6b9265b, 0x6fb077e1, 0x18b74777,
1056 0x88085ae6, 0xff0f6a70, 0x66063bca, 0x11010b5c,
1057 0x8f659eff, 0xf862ae69, 0x616bffd3, 0x166ccf45,
1058 0xa00ae278, 0xd70dd2ee, 0x4e048354, 0x3903b3c2,
1059 0xa7672661, 0xd06016f7, 0x4969474d, 0x3e6e77db,
1060 0xaed16a4a, 0xd9d65adc, 0x40df0b66, 0x37d83bf0,
1061 0xa9bcae53, 0xdebb9ec5, 0x47b2cf7f, 0x30b5ffe9,
1062 0xbdbdf21c, 0xcabac28a, 0x53b39330, 0x24b4a3a6,
1063 0xbad03605, 0xcdd70693, 0x54de5729, 0x23d967bf,
1064 0xb3667a2e, 0xc4614ab8, 0x5d681b02, 0x2a6f2b94,
1065 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d
1068 u32 round_crc32 (u32 a, const u32 v)
1070 const u32 k = (a ^ v) & 0xff;
1072 const u32 s = a >> 8;
1081 u32 crc32 (const u32 w[16], const u32 pw_len, const u32 iv)
1085 if (pw_len >= 1) a = round_crc32 (a, w[0] >> 0);
1086 if (pw_len >= 2) a = round_crc32 (a, w[0] >> 8);
1087 if (pw_len >= 3) a = round_crc32 (a, w[0] >> 16);
1088 if (pw_len >= 4) a = round_crc32 (a, w[0] >> 24);
1090 for (u32 i = 4, j = 1; i < pw_len; i += 4, j += 1)
1092 if (pw_len >= (i + 1)) a = round_crc32 (a, w[j] >> 0);
1093 if (pw_len >= (i + 2)) a = round_crc32 (a, w[j] >> 8);
1094 if (pw_len >= (i + 3)) a = round_crc32 (a, w[j] >> 16);
1095 if (pw_len >= (i + 4)) a = round_crc32 (a, w[j] >> 24);
1101 u32 memcat8c_be (u32 block[16], const u32 block_len, const u32 append, const u32 append_len, u32 digest[8])
1103 const u32 mod = block_len & 3;
1104 const u32 div = block_len / 4;
1110 const int selector = (0x76543210 >> ((block_len & 3) * 4)) & 0xffff;
1112 tmp0 = __byte_perm (append, 0, selector);
1113 tmp1 = __byte_perm (0, append, selector);
1116 #if defined IS_AMD || defined IS_GENERIC
1117 tmp0 = amd_bytealign (0, append, block_len);
1118 tmp1 = amd_bytealign (append, 0, block_len);
1125 case 0: block[ 0] |= tmp0;
1128 case 1: block[ 1] |= tmp0;
1131 case 2: block[ 2] |= tmp0;
1134 case 3: block[ 3] |= tmp0;
1137 case 4: block[ 4] |= tmp0;
1140 case 5: block[ 5] |= tmp0;
1143 case 6: block[ 6] |= tmp0;
1146 case 7: block[ 7] |= tmp0;
1149 case 8: block[ 8] |= tmp0;
1152 case 9: block[ 9] |= tmp0;
1155 case 10: block[10] |= tmp0;
1158 case 11: block[11] |= tmp0;
1161 case 12: block[12] |= tmp0;
1164 case 13: block[13] |= tmp0;
1167 case 14: block[14] |= tmp0;
1170 case 15: block[15] |= tmp0;
1175 u32 new_len = block_len + append_len;
1181 sha256_transform (block, digest);
1204 u32 memcat64c_be (u32 block[16], const u32 block_len, const u32 append[16], const u32 append_len, u32 digest[8])
1206 const u32 mod = block_len & 3;
1207 const u32 div = block_len / 4;
1228 const int selector = (0x76543210 >> ((block_len & 3) * 4)) & 0xffff;
1230 tmp00 = __byte_perm (append[ 0], 0, selector);
1231 tmp01 = __byte_perm (append[ 1], append[ 0], selector);
1232 tmp02 = __byte_perm (append[ 2], append[ 1], selector);
1233 tmp03 = __byte_perm (append[ 3], append[ 2], selector);
1234 tmp04 = __byte_perm (append[ 4], append[ 3], selector);
1235 tmp05 = __byte_perm (append[ 5], append[ 4], selector);
1236 tmp06 = __byte_perm (append[ 6], append[ 5], selector);
1237 tmp07 = __byte_perm (append[ 7], append[ 6], selector);
1238 tmp08 = __byte_perm (append[ 8], append[ 7], selector);
1239 tmp09 = __byte_perm (append[ 9], append[ 8], selector);
1240 tmp10 = __byte_perm (append[10], append[ 9], selector);
1241 tmp11 = __byte_perm (append[11], append[10], selector);
1242 tmp12 = __byte_perm (append[12], append[11], selector);
1243 tmp13 = __byte_perm (append[13], append[12], selector);
1244 tmp14 = __byte_perm (append[14], append[13], selector);
1245 tmp15 = __byte_perm (append[15], append[14], selector);
1246 tmp16 = __byte_perm ( 0, append[15], selector);
1249 #if defined IS_AMD || defined IS_GENERIC
1250 tmp00 = amd_bytealign ( 0, append[ 0], block_len);
1251 tmp01 = amd_bytealign (append[ 0], append[ 1], block_len);
1252 tmp02 = amd_bytealign (append[ 1], append[ 2], block_len);
1253 tmp03 = amd_bytealign (append[ 2], append[ 3], block_len);
1254 tmp04 = amd_bytealign (append[ 3], append[ 4], block_len);
1255 tmp05 = amd_bytealign (append[ 4], append[ 5], block_len);
1256 tmp06 = amd_bytealign (append[ 5], append[ 6], block_len);
1257 tmp07 = amd_bytealign (append[ 6], append[ 7], block_len);
1258 tmp08 = amd_bytealign (append[ 7], append[ 8], block_len);
1259 tmp09 = amd_bytealign (append[ 8], append[ 9], block_len);
1260 tmp10 = amd_bytealign (append[ 9], append[10], block_len);
1261 tmp11 = amd_bytealign (append[10], append[11], block_len);
1262 tmp12 = amd_bytealign (append[11], append[12], block_len);
1263 tmp13 = amd_bytealign (append[12], append[13], block_len);
1264 tmp14 = amd_bytealign (append[13], append[14], block_len);
1265 tmp15 = amd_bytealign (append[14], append[15], block_len);
1266 tmp16 = amd_bytealign (append[15], 0, block_len);
1269 u32 carry[16] = { 0 };
1273 case 0: block[ 0] |= tmp00;
1291 case 1: block[ 1] |= tmp00;
1309 case 2: block[ 2] |= tmp00;
1327 case 3: block[ 3] |= tmp00;
1345 case 4: block[ 4] |= tmp00;
1363 case 5: block[ 5] |= tmp00;
1381 case 6: block[ 6] |= tmp00;
1399 case 7: block[ 7] |= tmp00;
1417 case 8: block[ 8] |= tmp00;
1435 case 9: block[ 9] |= tmp00;
1453 case 10: block[10] |= tmp00;
1471 case 11: block[11] |= tmp00;
1489 case 12: block[12] |= tmp00;
1507 case 13: block[13] |= tmp00;
1525 case 14: block[14] |= tmp00;
1543 case 15: block[15] |= tmp00;
1563 u32 new_len = block_len + append_len;
1569 sha256_transform (block, digest);
1571 block[ 0] = carry[ 0];
1572 block[ 1] = carry[ 1];
1573 block[ 2] = carry[ 2];
1574 block[ 3] = carry[ 3];
1575 block[ 4] = carry[ 4];
1576 block[ 5] = carry[ 5];
1577 block[ 6] = carry[ 6];
1578 block[ 7] = carry[ 7];
1579 block[ 8] = carry[ 8];
1580 block[ 9] = carry[ 9];
1581 block[10] = carry[10];
1582 block[11] = carry[11];
1583 block[12] = carry[12];
1584 block[13] = carry[13];
1585 block[14] = carry[14];
1586 block[15] = carry[15];
1592 __kernel void m11600_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_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 seven_zip_t *esalt_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)
1598 const u32 gid = get_global_id (0);
1600 if (gid >= gid_max) return;
1606 tmps[gid].dgst[0] = SHA256M_A;
1607 tmps[gid].dgst[1] = SHA256M_B;
1608 tmps[gid].dgst[2] = SHA256M_C;
1609 tmps[gid].dgst[3] = SHA256M_D;
1610 tmps[gid].dgst[4] = SHA256M_E;
1611 tmps[gid].dgst[5] = SHA256M_F;
1612 tmps[gid].dgst[6] = SHA256M_G;
1613 tmps[gid].dgst[7] = SHA256M_H;
1615 tmps[gid].block[ 0] = 0;
1616 tmps[gid].block[ 1] = 0;
1617 tmps[gid].block[ 2] = 0;
1618 tmps[gid].block[ 3] = 0;
1619 tmps[gid].block[ 4] = 0;
1620 tmps[gid].block[ 5] = 0;
1621 tmps[gid].block[ 6] = 0;
1622 tmps[gid].block[ 7] = 0;
1623 tmps[gid].block[ 8] = 0;
1624 tmps[gid].block[ 9] = 0;
1625 tmps[gid].block[10] = 0;
1626 tmps[gid].block[11] = 0;
1627 tmps[gid].block[12] = 0;
1628 tmps[gid].block[13] = 0;
1629 tmps[gid].block[14] = 0;
1630 tmps[gid].block[15] = 0;
1632 tmps[gid].block_len = 0;
1633 tmps[gid].final_len = 0;
1636 __kernel void m11600_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_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 seven_zip_t *esalt_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)
1642 const u32 gid = get_global_id (0);
1644 if (gid >= gid_max) return;
1648 pw[0] = pws[gid].i[0];
1649 pw[1] = pws[gid].i[1];
1650 pw[2] = pws[gid].i[2];
1651 pw[3] = pws[gid].i[3];
1652 pw[4] = pws[gid].i[4];
1653 pw[5] = pws[gid].i[5];
1654 pw[6] = pws[gid].i[6];
1655 pw[7] = pws[gid].i[7];
1657 u32 pw_len = pws[gid].pw_len;
1659 make_unicode (&pw[ 4], &pw[ 8], &pw[12]);
1660 make_unicode (&pw[ 0], &pw[ 0], &pw[ 4]);
1664 pw[ 0] = swap32 (pw[ 0]);
1665 pw[ 1] = swap32 (pw[ 1]);
1666 pw[ 2] = swap32 (pw[ 2]);
1667 pw[ 3] = swap32 (pw[ 3]);
1668 pw[ 4] = swap32 (pw[ 4]);
1669 pw[ 5] = swap32 (pw[ 5]);
1670 pw[ 6] = swap32 (pw[ 6]);
1671 pw[ 7] = swap32 (pw[ 7]);
1672 pw[ 8] = swap32 (pw[ 8]);
1673 pw[ 9] = swap32 (pw[ 9]);
1674 pw[10] = swap32 (pw[10]);
1675 pw[11] = swap32 (pw[11]);
1676 pw[12] = swap32 (pw[12]);
1677 pw[13] = swap32 (pw[13]);
1678 pw[14] = swap32 (pw[14]);
1679 pw[15] = swap32 (pw[15]);
1687 dgst[0] = tmps[gid].dgst[0];
1688 dgst[1] = tmps[gid].dgst[1];
1689 dgst[2] = tmps[gid].dgst[2];
1690 dgst[3] = tmps[gid].dgst[3];
1691 dgst[4] = tmps[gid].dgst[4];
1692 dgst[5] = tmps[gid].dgst[5];
1693 dgst[6] = tmps[gid].dgst[6];
1694 dgst[7] = tmps[gid].dgst[7];
1698 block[ 0] = tmps[gid].block[ 0];
1699 block[ 1] = tmps[gid].block[ 1];
1700 block[ 2] = tmps[gid].block[ 2];
1701 block[ 3] = tmps[gid].block[ 3];
1702 block[ 4] = tmps[gid].block[ 4];
1703 block[ 5] = tmps[gid].block[ 5];
1704 block[ 6] = tmps[gid].block[ 6];
1705 block[ 7] = tmps[gid].block[ 7];
1706 block[ 8] = tmps[gid].block[ 8];
1707 block[ 9] = tmps[gid].block[ 9];
1708 block[10] = tmps[gid].block[10];
1709 block[11] = tmps[gid].block[11];
1710 block[12] = tmps[gid].block[12];
1711 block[13] = tmps[gid].block[13];
1712 block[14] = tmps[gid].block[14];
1713 block[15] = tmps[gid].block[15];
1715 u32 block_len = tmps[gid].block_len;
1716 u32 final_len = tmps[gid].final_len;
1722 for (u32 i = 0, j = loop_pos; i < loop_cnt; i++, j++)
1724 const u32 j_swap = swap32 (j);
1726 block_len = memcat64c_be (block, block_len, pw, pw_len, dgst); final_len += pw_len;
1727 block_len = memcat8c_be (block, block_len, j_swap, 8, dgst); final_len += 8;
1734 tmps[gid].dgst[0] = dgst[0];
1735 tmps[gid].dgst[1] = dgst[1];
1736 tmps[gid].dgst[2] = dgst[2];
1737 tmps[gid].dgst[3] = dgst[3];
1738 tmps[gid].dgst[4] = dgst[4];
1739 tmps[gid].dgst[5] = dgst[5];
1740 tmps[gid].dgst[6] = dgst[6];
1741 tmps[gid].dgst[7] = dgst[7];
1743 tmps[gid].block[ 0] = block[ 0];
1744 tmps[gid].block[ 1] = block[ 1];
1745 tmps[gid].block[ 2] = block[ 2];
1746 tmps[gid].block[ 3] = block[ 3];
1747 tmps[gid].block[ 4] = block[ 4];
1748 tmps[gid].block[ 5] = block[ 5];
1749 tmps[gid].block[ 6] = block[ 6];
1750 tmps[gid].block[ 7] = block[ 7];
1751 tmps[gid].block[ 8] = block[ 8];
1752 tmps[gid].block[ 9] = block[ 9];
1753 tmps[gid].block[10] = block[10];
1754 tmps[gid].block[11] = block[11];
1755 tmps[gid].block[12] = block[12];
1756 tmps[gid].block[13] = block[13];
1757 tmps[gid].block[14] = block[14];
1758 tmps[gid].block[15] = block[15];
1760 tmps[gid].block_len = block_len;
1761 tmps[gid].final_len = final_len;
1764 __kernel void m11600_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global seven_zip_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 seven_zip_t *esalt_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)
1770 const u32 gid = get_global_id (0);
1771 const u32 lid = get_local_id (0);
1772 const u32 lsz = get_local_size (0);
1778 __local u32 s_td0[256];
1779 __local u32 s_td1[256];
1780 __local u32 s_td2[256];
1781 __local u32 s_td3[256];
1782 __local u32 s_td4[256];
1784 __local u32 s_te0[256];
1785 __local u32 s_te1[256];
1786 __local u32 s_te2[256];
1787 __local u32 s_te3[256];
1788 __local u32 s_te4[256];
1790 for (u32 i = lid; i < 256; i += lsz)
1805 barrier (CLK_LOCAL_MEM_FENCE);
1807 if (gid >= gid_max) return;
1815 dgst[0] = tmps[gid].dgst[0];
1816 dgst[1] = tmps[gid].dgst[1];
1817 dgst[2] = tmps[gid].dgst[2];
1818 dgst[3] = tmps[gid].dgst[3];
1819 dgst[4] = tmps[gid].dgst[4];
1820 dgst[5] = tmps[gid].dgst[5];
1821 dgst[6] = tmps[gid].dgst[6];
1822 dgst[7] = tmps[gid].dgst[7];
1824 u32 block_len = tmps[gid].block_len;
1825 u32 final_len = tmps[gid].final_len;
1827 // this optimization should work as long as we have an iteration 6 or higher
1831 block[ 0] = 0x80000000;
1846 block[15] = final_len * 8;
1848 sha256_transform (block, dgst);
1851 * final key operations
1856 iv[0] = esalt_bufs[salt_pos].iv_buf[0];
1857 iv[1] = esalt_bufs[salt_pos].iv_buf[1];
1858 iv[2] = esalt_bufs[salt_pos].iv_buf[2];
1859 iv[3] = esalt_bufs[salt_pos].iv_buf[3];
1876 AES256_ExpandKey (ukey, rk, s_te0, s_te1, s_te2, s_te3, s_te4);
1878 AES256_InvertKey (rk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
1882 int data_len = esalt_bufs[salt_pos].data_len;
1883 int unpack_size = esalt_bufs[salt_pos].unpack_size;
1888 for (i = 0, j = 0; i < data_len - 16; i += 16, j += 4)
1892 data[0] = swap32 (esalt_bufs[salt_pos].data_buf[j + 0]);
1893 data[1] = swap32 (esalt_bufs[salt_pos].data_buf[j + 1]);
1894 data[2] = swap32 (esalt_bufs[salt_pos].data_buf[j + 2]);
1895 data[3] = swap32 (esalt_bufs[salt_pos].data_buf[j + 3]);
1899 AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1911 out[0] = swap32 (out[0]);
1912 out[1] = swap32 (out[1]);
1913 out[2] = swap32 (out[2]);
1914 out[3] = swap32 (out[3]);
1916 crc = crc32 (out, 16, crc);
1921 data[0] = swap32 (esalt_bufs[salt_pos].data_buf[j + 0]);
1922 data[1] = swap32 (esalt_bufs[salt_pos].data_buf[j + 1]);
1923 data[2] = swap32 (esalt_bufs[salt_pos].data_buf[j + 2]);
1924 data[3] = swap32 (esalt_bufs[salt_pos].data_buf[j + 3]);
1928 AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1940 out[0] = swap32 (out[0]);
1941 out[1] = swap32 (out[1]);
1942 out[2] = swap32 (out[2]);
1943 out[3] = swap32 (out[3]);
1945 const u32 margin = data_len - unpack_size;
1947 const u32 left = 16 - margin;
1949 crc = crc32 (out, left, crc);
1951 // use padding attack in that case
1957 case 15: out[0] &= 0xffffff00;
1959 case 14: out[0] &= 0xffff0000;
1961 case 13: out[0] &= 0xff000000;
1963 case 12: out[0] = 0;
1965 case 11: out[0] = 0;
1966 out[1] &= 0xffffff00;
1968 case 10: out[0] = 0;
1969 out[1] &= 0xffff0000;
1972 out[1] &= 0xff000000;
1979 out[2] &= 0xffffff00;
1983 out[2] &= 0xffff0000;
1987 out[2] &= 0xff000000;
1996 out[3] &= 0xffffff00;
2001 out[3] &= 0xffff0000;
2006 out[3] &= 0xff000000;
2010 if ((out[0] == 0) && (out[1] == 0) && (out[2] == 0) && (out[3] == 0))
2012 mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, 0);