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