2 * Authors.....: Jens Steube <jens.steube@gmail.com>
3 * Gabriele Gristina <matrix@hashcat.net>
10 #include "include/constants.h"
11 #include "include/kernel_vendor.h"
18 #include "include/kernel_functions.c"
19 #include "OpenCL/types_ocl.c"
20 #include "OpenCL/common.c"
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 static 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)
730 u32 temp = rek[j + 7];
732 rek[j + 8] = rek[j + 0]
733 ^ (s_te2[(temp >> 16) & 0xff] & 0xff000000)
734 ^ (s_te3[(temp >> 8) & 0xff] & 0x00ff0000)
735 ^ (s_te0[(temp >> 0) & 0xff] & 0x0000ff00)
736 ^ (s_te1[(temp >> 24) & 0xff] & 0x000000ff)
739 rek[j + 9] = rek[j + 1] ^ rek[j + 8];
740 rek[j + 10] = rek[j + 2] ^ rek[j + 9];
741 rek[j + 11] = rek[j + 3] ^ rek[j + 10];
751 rek[j + 12] = rek[j + 4]
752 ^ (s_te2[(temp >> 24) & 0xff] & 0xff000000)
753 ^ (s_te3[(temp >> 16) & 0xff] & 0x00ff0000)
754 ^ (s_te0[(temp >> 8) & 0xff] & 0x0000ff00)
755 ^ (s_te1[(temp >> 0) & 0xff] & 0x000000ff);
757 rek[j + 13] = rek[j + 5] ^ rek[j + 12];
758 rek[j + 14] = rek[j + 6] ^ rek[j + 13];
759 rek[j + 15] = rek[j + 7] ^ rek[j + 14];
765 static 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)
767 for (u32 i = 0, j = 56; i < j; i += 4, j -= 4)
771 temp = rdk[i + 0]; rdk[i + 0] = rdk[j + 0]; rdk[j + 0] = temp;
772 temp = rdk[i + 1]; rdk[i + 1] = rdk[j + 1]; rdk[j + 1] = temp;
773 temp = rdk[i + 2]; rdk[i + 2] = rdk[j + 2]; rdk[j + 2] = temp;
774 temp = rdk[i + 3]; rdk[i + 3] = rdk[j + 3]; rdk[j + 3] = temp;
777 for (u32 i = 1, j = 4; i < 14; i += 1, j += 4)
780 s_td0[s_te1[(rdk[j + 0] >> 24) & 0xff] & 0xff] ^
781 s_td1[s_te1[(rdk[j + 0] >> 16) & 0xff] & 0xff] ^
782 s_td2[s_te1[(rdk[j + 0] >> 8) & 0xff] & 0xff] ^
783 s_td3[s_te1[(rdk[j + 0] >> 0) & 0xff] & 0xff];
786 s_td0[s_te1[(rdk[j + 1] >> 24) & 0xff] & 0xff] ^
787 s_td1[s_te1[(rdk[j + 1] >> 16) & 0xff] & 0xff] ^
788 s_td2[s_te1[(rdk[j + 1] >> 8) & 0xff] & 0xff] ^
789 s_td3[s_te1[(rdk[j + 1] >> 0) & 0xff] & 0xff];
792 s_td0[s_te1[(rdk[j + 2] >> 24) & 0xff] & 0xff] ^
793 s_td1[s_te1[(rdk[j + 2] >> 16) & 0xff] & 0xff] ^
794 s_td2[s_te1[(rdk[j + 2] >> 8) & 0xff] & 0xff] ^
795 s_td3[s_te1[(rdk[j + 2] >> 0) & 0xff] & 0xff];
798 s_td0[s_te1[(rdk[j + 3] >> 24) & 0xff] & 0xff] ^
799 s_td1[s_te1[(rdk[j + 3] >> 16) & 0xff] & 0xff] ^
800 s_td2[s_te1[(rdk[j + 3] >> 8) & 0xff] & 0xff] ^
801 s_td3[s_te1[(rdk[j + 3] >> 0) & 0xff] & 0xff];
805 static 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)
807 u32 s0 = in[0] ^ rdk[0];
808 u32 s1 = in[1] ^ rdk[1];
809 u32 s2 = in[2] ^ rdk[2];
810 u32 s3 = in[3] ^ rdk[3];
817 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[ 4];
818 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[ 5];
819 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[ 6];
820 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[ 7];
821 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[ 8];
822 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[ 9];
823 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[10];
824 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[11];
825 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[12];
826 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[13];
827 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[14];
828 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[15];
829 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[16];
830 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[17];
831 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[18];
832 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[19];
833 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[20];
834 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[21];
835 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[22];
836 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[23];
837 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[24];
838 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[25];
839 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[26];
840 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[27];
841 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[28];
842 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[29];
843 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[30];
844 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[31];
845 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[32];
846 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[33];
847 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[34];
848 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[35];
849 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[36];
850 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[37];
851 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[38];
852 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[39];
853 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[40];
854 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[41];
855 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[42];
856 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[43];
857 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[44];
858 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[45];
859 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[46];
860 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[47];
861 s0 = s_td0[t0 >> 24] ^ s_td1[(t3 >> 16) & 0xff] ^ s_td2[(t2 >> 8) & 0xff] ^ s_td3[t1 & 0xff] ^ rdk[48];
862 s1 = s_td0[t1 >> 24] ^ s_td1[(t0 >> 16) & 0xff] ^ s_td2[(t3 >> 8) & 0xff] ^ s_td3[t2 & 0xff] ^ rdk[49];
863 s2 = s_td0[t2 >> 24] ^ s_td1[(t1 >> 16) & 0xff] ^ s_td2[(t0 >> 8) & 0xff] ^ s_td3[t3 & 0xff] ^ rdk[50];
864 s3 = s_td0[t3 >> 24] ^ s_td1[(t2 >> 16) & 0xff] ^ s_td2[(t1 >> 8) & 0xff] ^ s_td3[t0 & 0xff] ^ rdk[51];
865 t0 = s_td0[s0 >> 24] ^ s_td1[(s3 >> 16) & 0xff] ^ s_td2[(s2 >> 8) & 0xff] ^ s_td3[s1 & 0xff] ^ rdk[52];
866 t1 = s_td0[s1 >> 24] ^ s_td1[(s0 >> 16) & 0xff] ^ s_td2[(s3 >> 8) & 0xff] ^ s_td3[s2 & 0xff] ^ rdk[53];
867 t2 = s_td0[s2 >> 24] ^ s_td1[(s1 >> 16) & 0xff] ^ s_td2[(s0 >> 8) & 0xff] ^ s_td3[s3 & 0xff] ^ rdk[54];
868 t3 = s_td0[s3 >> 24] ^ s_td1[(s2 >> 16) & 0xff] ^ s_td2[(s1 >> 8) & 0xff] ^ s_td3[s0 & 0xff] ^ rdk[55];
870 out[0] = (s_td4[(t0 >> 24) & 0xff] & 0xff000000)
871 ^ (s_td4[(t3 >> 16) & 0xff] & 0x00ff0000)
872 ^ (s_td4[(t2 >> 8) & 0xff] & 0x0000ff00)
873 ^ (s_td4[(t1 >> 0) & 0xff] & 0x000000ff)
876 out[1] = (s_td4[(t1 >> 24) & 0xff] & 0xff000000)
877 ^ (s_td4[(t0 >> 16) & 0xff] & 0x00ff0000)
878 ^ (s_td4[(t3 >> 8) & 0xff] & 0x0000ff00)
879 ^ (s_td4[(t2 >> 0) & 0xff] & 0x000000ff)
882 out[2] = (s_td4[(t2 >> 24) & 0xff] & 0xff000000)
883 ^ (s_td4[(t1 >> 16) & 0xff] & 0x00ff0000)
884 ^ (s_td4[(t0 >> 8) & 0xff] & 0x0000ff00)
885 ^ (s_td4[(t3 >> 0) & 0xff] & 0x000000ff)
888 out[3] = (s_td4[(t3 >> 24) & 0xff] & 0xff000000)
889 ^ (s_td4[(t2 >> 16) & 0xff] & 0x00ff0000)
890 ^ (s_td4[(t1 >> 8) & 0xff] & 0x0000ff00)
891 ^ (s_td4[(t0 >> 0) & 0xff] & 0x000000ff)
895 __constant u64 k_sha512[80] =
897 SHA512C00, SHA512C01, SHA512C02, SHA512C03,
898 SHA512C04, SHA512C05, SHA512C06, SHA512C07,
899 SHA512C08, SHA512C09, SHA512C0a, SHA512C0b,
900 SHA512C0c, SHA512C0d, SHA512C0e, SHA512C0f,
901 SHA512C10, SHA512C11, SHA512C12, SHA512C13,
902 SHA512C14, SHA512C15, SHA512C16, SHA512C17,
903 SHA512C18, SHA512C19, SHA512C1a, SHA512C1b,
904 SHA512C1c, SHA512C1d, SHA512C1e, SHA512C1f,
905 SHA512C20, SHA512C21, SHA512C22, SHA512C23,
906 SHA512C24, SHA512C25, SHA512C26, SHA512C27,
907 SHA512C28, SHA512C29, SHA512C2a, SHA512C2b,
908 SHA512C2c, SHA512C2d, SHA512C2e, SHA512C2f,
909 SHA512C30, SHA512C31, SHA512C32, SHA512C33,
910 SHA512C34, SHA512C35, SHA512C36, SHA512C37,
911 SHA512C38, SHA512C39, SHA512C3a, SHA512C3b,
912 SHA512C3c, SHA512C3d, SHA512C3e, SHA512C3f,
913 SHA512C40, SHA512C41, SHA512C42, SHA512C43,
914 SHA512C44, SHA512C45, SHA512C46, SHA512C47,
915 SHA512C48, SHA512C49, SHA512C4a, SHA512C4b,
916 SHA512C4c, SHA512C4d, SHA512C4e, SHA512C4f,
919 static void sha512_transform (const u64 w[16], u64 dgst[8])
947 #define ROUND_EXPAND() \
949 w0_t = SHA512_EXPAND (we_t, w9_t, w1_t, w0_t); \
950 w1_t = SHA512_EXPAND (wf_t, wa_t, w2_t, w1_t); \
951 w2_t = SHA512_EXPAND (w0_t, wb_t, w3_t, w2_t); \
952 w3_t = SHA512_EXPAND (w1_t, wc_t, w4_t, w3_t); \
953 w4_t = SHA512_EXPAND (w2_t, wd_t, w5_t, w4_t); \
954 w5_t = SHA512_EXPAND (w3_t, we_t, w6_t, w5_t); \
955 w6_t = SHA512_EXPAND (w4_t, wf_t, w7_t, w6_t); \
956 w7_t = SHA512_EXPAND (w5_t, w0_t, w8_t, w7_t); \
957 w8_t = SHA512_EXPAND (w6_t, w1_t, w9_t, w8_t); \
958 w9_t = SHA512_EXPAND (w7_t, w2_t, wa_t, w9_t); \
959 wa_t = SHA512_EXPAND (w8_t, w3_t, wb_t, wa_t); \
960 wb_t = SHA512_EXPAND (w9_t, w4_t, wc_t, wb_t); \
961 wc_t = SHA512_EXPAND (wa_t, w5_t, wd_t, wc_t); \
962 wd_t = SHA512_EXPAND (wb_t, w6_t, we_t, wd_t); \
963 we_t = SHA512_EXPAND (wc_t, w7_t, wf_t, we_t); \
964 wf_t = SHA512_EXPAND (wd_t, w8_t, w0_t, wf_t); \
967 #define ROUND_STEP(i) \
969 SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha512[i + 0]); \
970 SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha512[i + 1]); \
971 SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha512[i + 2]); \
972 SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha512[i + 3]); \
973 SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha512[i + 4]); \
974 SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha512[i + 5]); \
975 SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha512[i + 6]); \
976 SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha512[i + 7]); \
977 SHA512_STEP (SHA512_F0o, SHA512_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha512[i + 8]); \
978 SHA512_STEP (SHA512_F0o, SHA512_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha512[i + 9]); \
979 SHA512_STEP (SHA512_F0o, SHA512_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha512[i + 10]); \
980 SHA512_STEP (SHA512_F0o, SHA512_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha512[i + 11]); \
981 SHA512_STEP (SHA512_F0o, SHA512_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha512[i + 12]); \
982 SHA512_STEP (SHA512_F0o, SHA512_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha512[i + 13]); \
983 SHA512_STEP (SHA512_F0o, SHA512_F1o, c, d, e, f, g, h, a, b, we_t, k_sha512[i + 14]); \
984 SHA512_STEP (SHA512_F0o, SHA512_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha512[i + 15]); \
990 for (int i = 16; i < 80; i += 16)
992 ROUND_EXPAND (); ROUND_STEP (i);
1005 __kernel void m11300_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global bitcoin_wallet_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 bitcoin_wallet_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1011 const u32 gid = get_global_id (0);
1013 if (gid >= gid_max) return;
1017 w0[0] = pws[gid].i[ 0];
1018 w0[1] = pws[gid].i[ 1];
1019 w0[2] = pws[gid].i[ 2];
1020 w0[3] = pws[gid].i[ 3];
1024 w1[0] = pws[gid].i[ 4];
1025 w1[1] = pws[gid].i[ 5];
1026 w1[2] = pws[gid].i[ 6];
1027 w1[3] = pws[gid].i[ 7];
1031 w2[0] = pws[gid].i[ 8];
1032 w2[1] = pws[gid].i[ 9];
1043 const u32 pw_len = pws[gid].pw_len;
1051 salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
1052 salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
1053 salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
1054 salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
1058 salt_buf1[0] = salt_bufs[salt_pos].salt_buf[4];
1059 salt_buf1[1] = salt_bufs[salt_pos].salt_buf[5];
1060 salt_buf1[2] = salt_bufs[salt_pos].salt_buf[6];
1061 salt_buf1[3] = salt_bufs[salt_pos].salt_buf[7];
1077 u32 salt_len = salt_bufs[salt_pos].salt_len;
1079 switch_buffer_by_offset_le (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len);
1081 w0[0] |= salt_buf0[0];
1082 w0[1] |= salt_buf0[1];
1083 w0[2] |= salt_buf0[2];
1084 w0[3] |= salt_buf0[3];
1086 w1[0] |= salt_buf1[0];
1087 w1[1] |= salt_buf1[1];
1088 w1[2] |= salt_buf1[2];
1089 w1[3] |= salt_buf1[3];
1091 w2[0] |= salt_buf2[0];
1092 w2[1] |= salt_buf2[1];
1093 w2[2] |= salt_buf2[2];
1094 w2[3] |= salt_buf2[3];
1096 w3[0] |= salt_buf3[0];
1097 w3[1] |= salt_buf3[1];
1098 w3[2] |= salt_buf3[2];
1099 w3[3] |= salt_buf3[3];
1101 const u32 block_len = pw_len + salt_len;
1107 w0[0] = swap32 (w0[0]);
1108 w0[1] = swap32 (w0[1]);
1109 w0[2] = swap32 (w0[2]);
1110 w0[3] = swap32 (w0[3]);
1111 w1[0] = swap32 (w1[0]);
1112 w1[1] = swap32 (w1[1]);
1113 w1[2] = swap32 (w1[2]);
1114 w1[3] = swap32 (w1[3]);
1115 w2[0] = swap32 (w2[0]);
1116 w2[1] = swap32 (w2[1]);
1117 w2[2] = swap32 (w2[2]);
1118 w2[3] = swap32 (w2[3]);
1119 w3[0] = swap32 (w3[0]);
1120 w3[1] = swap32 (w3[1]);
1121 w3[2] = swap32 (w3[2]);
1122 w3[3] = swap32 (w3[3]);
1126 w[ 0] = hl32_to_64 (w0[0], w0[1]);
1127 w[ 1] = hl32_to_64 (w0[2], w0[3]);
1128 w[ 2] = hl32_to_64 (w1[0], w1[1]);
1129 w[ 3] = hl32_to_64 (w1[2], w1[3]);
1130 w[ 4] = hl32_to_64 (w2[0], w2[1]);
1131 w[ 5] = hl32_to_64 (w2[2], w2[3]);
1132 w[ 6] = hl32_to_64 (w3[0], w3[1]);
1133 w[ 7] = hl32_to_64 (w3[2], w3[3]);
1141 w[15] = block_len * 8;
1145 dgst[0] = SHA512M_A;
1146 dgst[1] = SHA512M_B;
1147 dgst[2] = SHA512M_C;
1148 dgst[3] = SHA512M_D;
1149 dgst[4] = SHA512M_E;
1150 dgst[5] = SHA512M_F;
1151 dgst[6] = SHA512M_G;
1152 dgst[7] = SHA512M_H;
1154 sha512_transform (w, dgst);
1156 tmps[gid].dgst[0] = dgst[0];
1157 tmps[gid].dgst[1] = dgst[1];
1158 tmps[gid].dgst[2] = dgst[2];
1159 tmps[gid].dgst[3] = dgst[3];
1160 tmps[gid].dgst[4] = dgst[4];
1161 tmps[gid].dgst[5] = dgst[5];
1162 tmps[gid].dgst[6] = dgst[6];
1163 tmps[gid].dgst[7] = dgst[7];
1166 __kernel void m11300_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global bitcoin_wallet_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 bitcoin_wallet_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1168 const u32 gid = get_global_id (0);
1170 if (gid >= gid_max) return;
1174 dgst[0] = tmps[gid].dgst[0];
1175 dgst[1] = tmps[gid].dgst[1];
1176 dgst[2] = tmps[gid].dgst[2];
1177 dgst[3] = tmps[gid].dgst[3];
1178 dgst[4] = tmps[gid].dgst[4];
1179 dgst[5] = tmps[gid].dgst[5];
1180 dgst[6] = tmps[gid].dgst[6];
1181 dgst[7] = tmps[gid].dgst[7];
1193 w[ 8] = 0x8000000000000000;
1202 for (u32 i = 0; i < loop_cnt; i++)
1213 dgst[0] = SHA512M_A;
1214 dgst[1] = SHA512M_B;
1215 dgst[2] = SHA512M_C;
1216 dgst[3] = SHA512M_D;
1217 dgst[4] = SHA512M_E;
1218 dgst[5] = SHA512M_F;
1219 dgst[6] = SHA512M_G;
1220 dgst[7] = SHA512M_H;
1222 sha512_transform (w, dgst);
1225 tmps[gid].dgst[0] = dgst[0];
1226 tmps[gid].dgst[1] = dgst[1];
1227 tmps[gid].dgst[2] = dgst[2];
1228 tmps[gid].dgst[3] = dgst[3];
1229 tmps[gid].dgst[4] = dgst[4];
1230 tmps[gid].dgst[5] = dgst[5];
1231 tmps[gid].dgst[6] = dgst[6];
1232 tmps[gid].dgst[7] = dgst[7];
1235 __kernel void m11300_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global bitcoin_wallet_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 bitcoin_wallet_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
1241 const u32 gid = get_global_id (0);
1242 const u32 lid = get_local_id (0);
1243 const u32 lsz = get_local_size (0);
1249 __local u32 s_td0[256];
1250 __local u32 s_td1[256];
1251 __local u32 s_td2[256];
1252 __local u32 s_td3[256];
1253 __local u32 s_td4[256];
1255 __local u32 s_te0[256];
1256 __local u32 s_te1[256];
1257 __local u32 s_te2[256];
1258 __local u32 s_te3[256];
1259 __local u32 s_te4[256];
1261 for (u32 i = lid; i < 256; i += lsz)
1276 barrier (CLK_LOCAL_MEM_FENCE);
1278 if (gid >= gid_max) return;
1286 dgst[0] = tmps[gid].dgst[0];
1287 dgst[1] = tmps[gid].dgst[1];
1288 dgst[2] = tmps[gid].dgst[2];
1289 dgst[3] = tmps[gid].dgst[3];
1290 dgst[4] = tmps[gid].dgst[4];
1291 dgst[5] = tmps[gid].dgst[5];
1292 dgst[6] = tmps[gid].dgst[6];
1293 dgst[7] = tmps[gid].dgst[7];
1297 key[0] = h32_from_64 (dgst[0]);
1298 key[1] = l32_from_64 (dgst[0]);
1299 key[2] = h32_from_64 (dgst[1]);
1300 key[3] = l32_from_64 (dgst[1]);
1301 key[4] = h32_from_64 (dgst[2]);
1302 key[5] = l32_from_64 (dgst[2]);
1303 key[6] = h32_from_64 (dgst[3]);
1304 key[7] = l32_from_64 (dgst[3]);
1308 iv[0] = h32_from_64 (dgst[4]);
1309 iv[1] = l32_from_64 (dgst[4]);
1310 iv[2] = h32_from_64 (dgst[5]);
1311 iv[3] = l32_from_64 (dgst[5]);
1317 AES256_ExpandKey (key, rk, s_te0, s_te1, s_te2, s_te3, s_te4);
1319 AES256_InvertKey (rk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4);
1323 for (u32 i = 0; i < esalt_bufs[salt_pos].cry_master_len; i += 16)
1327 data[0] = swap32 (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 0]);
1328 data[1] = swap32 (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 1]);
1329 data[2] = swap32 (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 2]);
1330 data[3] = swap32 (esalt_bufs[salt_pos].cry_master_buf[(i / 4) + 3]);
1332 AES256_decrypt (data, out, rk, s_td0, s_td1, s_td2, s_td3, s_td4);
1345 if ((out[0] == 0x10101010)
1346 && (out[1] == 0x10101010)
1347 && (out[2] == 0x10101010)
1348 && (out[3] == 0x10101010))
1350 mark_hash (plains_buf, hashes_shown, digests_offset + 0, gid, 0);
1352 d_return_buf[lid] = 1;