forked from hashcat/hashcat
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathinc_common.h
383 lines (343 loc) · 26.8 KB
/
inc_common.h
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
/**
* Author......: See docs/credits.txt
* License.....: MIT
*/
#ifndef INC_COMMON_H
#define INC_COMMON_H
/*
* Prototype kernel function that fits all kernel macros
*
* There are four variables where major differences occur:
*
* - P2: Address space of kernel_rules_t struct.
* If the kernel uses rules_buf, it will be stored in CONSTANT_AS.
* If it does not, cheaper GLOBAL_AS space is used.
*
* - P4: Innerloop word buffer:
* Most kernels use a bf_t structure in GLOBAL_AS address space (_BASIC).
* Some use u32x pointer to a vector in CONSTANT_AS address space (_VECTOR).
* A few use a specific bs_word_t struct (_BITSLICE).
*
* - P5: Type of the tmps structure with additional data, or void.
* Used with slow hash types (ATTACK_EXEC_OUTSIDE_KERNEL) only.
*
* - P19: Type of the esalt_bufs structure with additional data, or void.
*/
#if defined IS_CUDA || defined IS_HIP
#define KERN_ATTR(p2,p4,p5,p6,p19) \
MAYBE_UNUSED GLOBAL_AS pw_t *pws, \
MAYBE_UNUSED p2 const kernel_rule_t *g_rules_buf, \
MAYBE_UNUSED GLOBAL_AS const pw_t *combs_buf, \
MAYBE_UNUSED p4, \
MAYBE_UNUSED GLOBAL_AS p5 *tmps, \
MAYBE_UNUSED GLOBAL_AS p6 *hooks, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_a, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_b, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_c, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_d, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_a, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_b, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_c, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_d, \
MAYBE_UNUSED GLOBAL_AS plain_t *plains_buf, \
MAYBE_UNUSED GLOBAL_AS const digest_t *digests_buf, \
MAYBE_UNUSED GLOBAL_AS u32 *hashes_shown, \
MAYBE_UNUSED GLOBAL_AS const salt_t *salt_bufs, \
MAYBE_UNUSED GLOBAL_AS const p19 *esalt_bufs, \
MAYBE_UNUSED GLOBAL_AS u32 *d_return_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra0_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra1_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra2_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra3_buf, \
MAYBE_UNUSED GLOBAL_AS const kernel_param_t *kernel_param
#else
#define KERN_ATTR(p2,p4,p5,p6,p19) \
MAYBE_UNUSED GLOBAL_AS pw_t *pws, \
MAYBE_UNUSED p2 const kernel_rule_t *rules_buf, \
MAYBE_UNUSED GLOBAL_AS const pw_t *combs_buf, \
MAYBE_UNUSED p4, \
MAYBE_UNUSED GLOBAL_AS p5 *tmps, \
MAYBE_UNUSED GLOBAL_AS p6 *hooks, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_a, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_b, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_c, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s1_d, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_a, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_b, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_c, \
MAYBE_UNUSED GLOBAL_AS const u32 *bitmaps_buf_s2_d, \
MAYBE_UNUSED GLOBAL_AS plain_t *plains_buf, \
MAYBE_UNUSED GLOBAL_AS const digest_t *digests_buf, \
MAYBE_UNUSED GLOBAL_AS u32 *hashes_shown, \
MAYBE_UNUSED GLOBAL_AS const salt_t *salt_bufs, \
MAYBE_UNUSED GLOBAL_AS const p19 *esalt_bufs, \
MAYBE_UNUSED GLOBAL_AS u32 *d_return_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra0_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra1_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra2_buf, \
MAYBE_UNUSED GLOBAL_AS void *d_extra3_buf, \
MAYBE_UNUSED GLOBAL_AS const kernel_param_t *kernel_param
#endif
/*
* Shortcut macros for usage in the actual kernels
*
* Not all possible combinations are needed. E.g. all kernels that use rules
* do not use the tmps pointer, all kernels that use a vector pointer in P4
* do not use rules or tmps, etc.
*/
#if defined IS_CUDA || defined IS_HIP
#define _KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, void)
#define _KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bs_word_t *g_words_buf_s, void, void, void)
#define _KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, e)
#define _KERN_ATTR_RULES() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, void)
#define _KERN_ATTR_RULES_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, void, void, e)
#define _KERN_ATTR_TMPS(t) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, void, void)
#define _KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, void, e)
#define _KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, h, void)
#define _KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const bf_t *g_bfs_buf, t, h, e)
#define _KERN_ATTR_VECTOR() KERN_ATTR (GLOBAL_AS, GLOBAL_AS const u32x *g_words_buf_r, void, void, void)
#define _KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (GLOBAL_AS, GLOBAL_AS const u32x *g_words_buf_r, void, void, e)
#else
#define _KERN_ATTR_BASIC() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, void, void, void)
#define _KERN_ATTR_BITSLICE() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bs_word_t *words_buf_s, void, void, void)
#define _KERN_ATTR_ESALT(e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, void, void, e)
#define _KERN_ATTR_RULES() KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, void)
#define _KERN_ATTR_RULES_ESALT(e) KERN_ATTR (CONSTANT_AS, GLOBAL_AS const bf_t *bfs_buf, void, void, e)
#define _KERN_ATTR_TMPS(t) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, void, void)
#define _KERN_ATTR_TMPS_ESALT(t,e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, void, e)
#define _KERN_ATTR_TMPS_HOOKS(t,h) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, h, void)
#define _KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const bf_t *bfs_buf, t, h, e)
#define _KERN_ATTR_VECTOR() KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, void)
#define _KERN_ATTR_VECTOR_ESALT(e) KERN_ATTR (GLOBAL_AS, CONSTANT_AS const u32x *words_buf_r, void, void, e)
#endif
/*
* with Metal we have additional parameters for the 'modifiers' (gid, lid, lsize, etc...)
* they should not be declared in the host code
* the compiler make the magic for us :)
*/
#if defined IS_METAL
#define KERN_ATTR_MAIN_PARAMS \
uint hc_gid [[ thread_position_in_grid ]], \
uint hc_lid [[ thread_position_in_threadgroup ]], \
uint hc_lsz [[ threads_per_threadgroup ]]
#endif // IS_METAL
/*
* TM kernel function parameters
*/
#define _KERN_ATTR_TM \
GLOBAL_AS u32 *mod, \
GLOBAL_AS bs_word_t *words_buf_b
/*
* Below are the macros to be used in the declarations of the KERNEL_FQ functions (pure and optimized kernels)
* With Metal we need to add KERN_ATTR_MAIN_PARAMS
* With CUDA, HIP and OpenCL there are no additional parameters other than those declared in the host code
*/
#if defined IS_METAL
#define KERN_ATTR_BASIC() _KERN_ATTR_BASIC(), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_BITSLICE() _KERN_ATTR_BITSLICE(), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_ESALT(e) _KERN_ATTR_ESALT(e), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_RULES() _KERN_ATTR_RULES(), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_RULES_ESALT(e) _KERN_ATTR_RULES_ESALT(e), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_TMPS(t) _KERN_ATTR_TMPS(t), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_TMPS_ESALT(t,e) _KERN_ATTR_TMPS_ESALT(t,e), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_TMPS_HOOKS(t,h) _KERN_ATTR_TMPS_HOOKS(t,h), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e) _KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_VECTOR() _KERN_ATTR_VECTOR(), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_VECTOR_ESALT(e) _KERN_ATTR_VECTOR_ESALT(e), KERN_ATTR_MAIN_PARAMS
#define KERN_ATTR_TM _KERN_ATTR_TM, KERN_ATTR_MAIN_PARAMS
#else // CUDA, HIP and OpenCL
#define KERN_ATTR_BASIC() _KERN_ATTR_BASIC()
#define KERN_ATTR_BITSLICE() _KERN_ATTR_BITSLICE()
#define KERN_ATTR_ESALT(e) _KERN_ATTR_ESALT(e)
#define KERN_ATTR_RULES() _KERN_ATTR_RULES()
#define KERN_ATTR_RULES_ESALT(e) _KERN_ATTR_RULES_ESALT(e)
#define KERN_ATTR_TMPS(t) _KERN_ATTR_TMPS(t)
#define KERN_ATTR_TMPS_ESALT(t,e) _KERN_ATTR_TMPS_ESALT(t,e)
#define KERN_ATTR_TMPS_HOOKS(t,h) _KERN_ATTR_TMPS_HOOKS(t,h)
#define KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e) _KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e)
#define KERN_ATTR_VECTOR() _KERN_ATTR_VECTOR()
#define KERN_ATTR_VECTOR_ESALT(e) _KERN_ATTR_VECTOR_ESALT(e)
#define KERN_ATTR_TM _KERN_ATTR_TM
#endif // IS_METAL
/*
* Below are the macros to be used in the declarations of the DECLSPEC functions
* that reuse the same parameters of KERNEL_FQ functions (-a3 optimized kernels)
* They are shared with CUDA, HIP, Metal and OpenCL runtime
* With these we can reuse 'modifier' (gid, lid, lsz, etc...) if we got it inside KERNEL functions
*/
#define KERN_ATTR_FUNC_PARAMS \
const u64 gid, \
MAYBE_UNUSED const u64 lid, \
MAYBE_UNUSED const u64 lsz
#define KERN_ATTR_FUNC_BASIC() _KERN_ATTR_BASIC(), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_BITSLICE() _KERN_ATTR_BITSLICE(), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_ESALT(e) _KERN_ATTR_ESALT(e), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_RULES() _KERN_ATTR_RULES(), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_RULES_ESALT(e) _KERN_ATTR_RULES_ESALT(e), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_TMPS(t) _KERN_ATTR_TMPS(t), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_TMPS_ESALT(t,e) _KERN_ATTR_TMPS_ESALT(t,e), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_TMPS_HOOKS(t,h) _KERN_ATTR_TMPS_HOOKS(t,h), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_TMPS_HOOKS_ESALT(t,h,e) _KERN_ATTR_TMPS_HOOKS_ESALT(t,h,e), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_VECTOR() _KERN_ATTR_VECTOR(), KERN_ATTR_FUNC_PARAMS
#define KERN_ATTR_FUNC_VECTOR_ESALT(e) _KERN_ATTR_VECTOR_ESALT(e), KERN_ATTR_FUNC_PARAMS
// union based packing
DECLSPEC u8 v8a_from_v32_S (const u32 v32);
DECLSPEC u8 v8b_from_v32_S (const u32 v32);
DECLSPEC u8 v8c_from_v32_S (const u32 v32);
DECLSPEC u8 v8d_from_v32_S (const u32 v32);
DECLSPEC u8 v8a_from_v64_S (const u64 v64);
DECLSPEC u8 v8b_from_v64_S (const u64 v64);
DECLSPEC u8 v8c_from_v64_S (const u64 v64);
DECLSPEC u8 v8d_from_v64_S (const u64 v64);
DECLSPEC u8 v8e_from_v64_S (const u64 v64);
DECLSPEC u8 v8f_from_v64_S (const u64 v64);
DECLSPEC u8 v8g_from_v64_S (const u64 v64);
DECLSPEC u8 v8h_from_v64_S (const u64 v64);
DECLSPEC u8x v8a_from_v64 (const u64x v64);
DECLSPEC u8x v8b_from_v64 (const u64x v64);
DECLSPEC u8x v8c_from_v64 (const u64x v64);
DECLSPEC u8x v8d_from_v64 (const u64x v64);
DECLSPEC u8x v8e_from_v64 (const u64x v64);
DECLSPEC u8x v8f_from_v64 (const u64x v64);
DECLSPEC u8x v8g_from_v64 (const u64x v64);
DECLSPEC u8x v8h_from_v64 (const u64x v64);
DECLSPEC u16 v16a_from_v32_S (const u32 v32);
DECLSPEC u16 v16b_from_v32_S (const u32 v32);
DECLSPEC u32 v32a_from_v64_S (const u64 v64);
DECLSPEC u32 v32b_from_v64_S (const u64 v64);
DECLSPEC u32 v32_from_v16ab_S (const u16 v16a, const u16 v16b);
DECLSPEC u64 v64_from_v32ab_S (const u32 v32a, const u32 v32b);
// inline asm packing
DECLSPEC u32x unpack_v8a_from_v32 (const u32x v32);
DECLSPEC u32x unpack_v8b_from_v32 (const u32x v32);
DECLSPEC u32x unpack_v8c_from_v32 (const u32x v32);
DECLSPEC u32x unpack_v8d_from_v32 (const u32x v32);
DECLSPEC u32 unpack_v8a_from_v32_S (const u32 v32);
DECLSPEC u32 unpack_v8b_from_v32_S (const u32 v32);
DECLSPEC u32 unpack_v8c_from_v32_S (const u32 v32);
DECLSPEC u32 unpack_v8d_from_v32_S (const u32 v32);
// opencl intern based packing
DECLSPEC u32x l32_from_64 (u64x a);
DECLSPEC u32x h32_from_64 (u64x a);
DECLSPEC u32 l32_from_64_S (u64 a);
DECLSPEC u32 h32_from_64_S (u64 a);
DECLSPEC u64x hl32_to_64 (const u32x a, const u32x b);
DECLSPEC u64 hl32_to_64_S (const u32 a, const u32 b);
// bit operations
DECLSPEC u32x hc_rotl32 (const u32x a, const int n);
DECLSPEC u32x hc_rotr32 (const u32x a, const int n);
DECLSPEC u32 hc_rotl32_S (const u32 a, const int n);
DECLSPEC u32 hc_rotr32_S (const u32 a, const int n);
DECLSPEC u64x hc_rotl64 (const u64x a, const int n);
DECLSPEC u64x hc_rotr64 (const u64x a, const int n);
DECLSPEC u64 hc_rotl64_S (const u64 a, const int n);
DECLSPEC u64 hc_rotr64_S (const u64 a, const int n);
DECLSPEC u32x hc_swap32 (const u32x v);
DECLSPEC u32 hc_swap32_S (const u32 v);
DECLSPEC u64x hc_swap64 (const u64x v);
DECLSPEC u64 hc_swap64_S (const u64 v);
// byte operations
DECLSPEC u32x hc_bytealign (const u32x a, const u32x b, const int c);
DECLSPEC u32 hc_bytealign_S (const u32 a, const u32 b, const int c);
DECLSPEC u32x hc_bytealign_be (const u32x a, const u32x b, const int c);
DECLSPEC u32 hc_bytealign_be_S (const u32 a, const u32 b, const int c);
DECLSPEC u32x hc_byte_perm (const u32x a, const u32x b, const int c);
DECLSPEC u32 hc_byte_perm_S (const u32 a, const u32 b, const int c);
DECLSPEC u32x hc_add3 (const u32x a, const u32x b, const u32x c);
DECLSPEC u32 hc_add3_S (const u32 a, const u32 b, const u32 c);
DECLSPEC u32x hc_bfe (const u32x a, const u32x b, const u32x c);
DECLSPEC u32 hc_bfe_S (const u32 a, const u32 b, const u32 c);
DECLSPEC u32x hc_lop_0x96 (const u32x a, const u32x b, const u32x c);
DECLSPEC u32 hc_lop_0x96_S (const u32 a, const u32 b, const u32 c);
// legacy common code
DECLSPEC int ffz (const u32 v);
#ifdef KERNEL_STATIC
DECLSPEC int hash_comp (PRIVATE_AS const u32 *d1, GLOBAL_AS const u32 *d2);
DECLSPEC int find_hash (PRIVATE_AS const u32 *digest, const u32 digests_cnt, GLOBAL_AS const digest_t *digests_buf);
#endif
DECLSPEC int hc_enc_scan (PRIVATE_AS const u32 *buf, const int len);
DECLSPEC int hc_enc_scan_global (GLOBAL_AS const u32 *buf, const int len);
DECLSPEC void hc_enc_init (PRIVATE_AS hc_enc_t *hc_enc);
DECLSPEC int hc_enc_has_next (PRIVATE_AS hc_enc_t *hc_enc, const int sz);
DECLSPEC int hc_enc_next (PRIVATE_AS hc_enc_t *hc_enc, PRIVATE_AS const u32 *src_buf, const int src_len, const int src_sz, PRIVATE_AS u32 *dst_buf, const int dst_sz);
DECLSPEC int hc_enc_next_global (PRIVATE_AS hc_enc_t *hc_enc, GLOBAL_AS const u32 *src_buf, const int src_len, const int src_sz, PRIVATE_AS u32 *dst_buf, const int dst_sz);
DECLSPEC int pkcs_padding_bs8 (PRIVATE_AS const u32 *data_buf, const int data_len);
DECLSPEC int pkcs_padding_bs16 (PRIVATE_AS const u32 *data_buf, const int data_len);
DECLSPEC int asn1_detect (PRIVATE_AS const u32 *buf, const int len);
DECLSPEC u32 check_bitmap (GLOBAL_AS const u32 *bitmap, const u32 bitmap_mask, const u32 bitmap_shift, const u32 digest);
DECLSPEC u32 check (PRIVATE_AS const u32 *digest, GLOBAL_AS const u32 *bitmap_s1_a, GLOBAL_AS const u32 *bitmap_s1_b, GLOBAL_AS const u32 *bitmap_s1_c, GLOBAL_AS const u32 *bitmap_s1_d, GLOBAL_AS const u32 *bitmap_s2_a, GLOBAL_AS const u32 *bitmap_s2_b, GLOBAL_AS const u32 *bitmap_s2_c, GLOBAL_AS const u32 *bitmap_s2_d, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2);
DECLSPEC void mark_hash (GLOBAL_AS plain_t *plains_buf, GLOBAL_AS u32 *d_result, const u32 salt_pos, const u32 digests_cnt, const u32 digest_pos, const u32 hash_pos, const u64 gid, const u32 il_pos, const u32 extra1, const u32 extra2);
DECLSPEC int hc_count_char (PRIVATE_AS const u32 *buf, const int elems, const u32 c);
DECLSPEC float hc_get_entropy (PRIVATE_AS const u32 *buf, const int elems);
DECLSPEC int is_valid_hex_8 (const u8 v);
DECLSPEC int is_valid_hex_32 (const u32 v);
DECLSPEC int is_valid_base58_8 (const u8 v);
DECLSPEC int is_valid_base58_32 (const u32 v);
DECLSPEC int is_valid_printable_8 (const u8 v);
DECLSPEC int is_valid_printable_32 (const u32 v);
DECLSPEC int hc_find_keyboard_layout_map (const u32 search, const int search_len, LOCAL_AS keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt);
DECLSPEC int hc_execute_keyboard_layout_mapping (PRIVATE_AS u32 *w, const int pw_len, LOCAL_AS keyboard_layout_mapping_t *s_keyboard_layout_mapping_buf, const int keyboard_layout_mapping_cnt);
DECLSPEC int count_bits_32 (const u32 v0, const u32 v1);
DECLSPEC void make_utf16be (PRIVATE_AS const u32x *in, PRIVATE_AS u32x *out1, PRIVATE_AS u32x *out2);
DECLSPEC void make_utf16beN (PRIVATE_AS const u32x *in, PRIVATE_AS u32x *out1, PRIVATE_AS u32x *out2);
DECLSPEC void make_utf16beN_S (PRIVATE_AS const u32 *in, PRIVATE_AS u32 *out1, PRIVATE_AS u32 *out2);
DECLSPEC void make_utf16le (PRIVATE_AS const u32x *in, PRIVATE_AS u32x *out1, PRIVATE_AS u32x *out2);
DECLSPEC void make_utf16leN (PRIVATE_AS const u32x *in, PRIVATE_AS u32x *out1, PRIVATE_AS u32x *out2);
DECLSPEC void undo_utf16be (PRIVATE_AS const u32x *in1, PRIVATE_AS const u32x *in2, PRIVATE_AS u32x *out);
DECLSPEC void undo_utf16le (PRIVATE_AS const u32x *in1, PRIVATE_AS const u32x *in2, PRIVATE_AS u32x *out);
DECLSPEC void set_mark_1x4 (PRIVATE_AS u32 *v, const u32 offset);
DECLSPEC void append_helper_1x4 (PRIVATE_AS u32x *r, const u32 v, PRIVATE_AS const u32 *m);
DECLSPEC void append_0x80_1x4 (PRIVATE_AS u32x *w0, const u32 offset);
DECLSPEC void append_0x80_2x4 (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, const u32 offset);
DECLSPEC void append_0x80_3x4 (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, const u32 offset);
DECLSPEC void append_0x80_4x4 (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32 offset);
DECLSPEC void append_0x80_8x4 (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, const u32 offset);
DECLSPEC void append_0x80_1x16 (PRIVATE_AS u32x *w, const u32 offset);
DECLSPEC void switch_buffer_by_offset_le (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_carry_le (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *c0, PRIVATE_AS u32x *c1, PRIVATE_AS u32x *c2, PRIVATE_AS u32x *c3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_be (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_carry_be (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *c0, PRIVATE_AS u32x *c1, PRIVATE_AS u32x *c2, PRIVATE_AS u32x *c3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_le (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_carry_le (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, PRIVATE_AS u32x *c0, PRIVATE_AS u32x *c1, PRIVATE_AS u32x *c2, PRIVATE_AS u32x *c3, PRIVATE_AS u32x *c4, PRIVATE_AS u32x *c5, PRIVATE_AS u32x *c6, PRIVATE_AS u32x *c7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_be (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_carry_be (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, PRIVATE_AS u32x *c0, PRIVATE_AS u32x *c1, PRIVATE_AS u32x *c2, PRIVATE_AS u32x *c3, PRIVATE_AS u32x *c4, PRIVATE_AS u32x *c5, PRIVATE_AS u32x *c6, PRIVATE_AS u32x *c7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_1x64_le (PRIVATE_AS u32x *w, const u32 offset);
DECLSPEC void switch_buffer_by_offset_1x64_be (PRIVATE_AS u32x *w, const u32 offset);
DECLSPEC void truncate_block_4x4_le_S (PRIVATE_AS u32 *w0, const u32 len);
DECLSPEC void truncate_block_4x4_be_S (PRIVATE_AS u32 *w0, const u32 len);
DECLSPEC void truncate_block_16x4_le_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 len);
DECLSPEC void truncate_block_16x4_be_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 len);
DECLSPEC void set_mark_1x4_S (PRIVATE_AS u32 *v, const u32 offset);
DECLSPEC void append_helper_1x4_S (PRIVATE_AS u32 *r, const u32 v, PRIVATE_AS const u32 *m);
DECLSPEC void append_0x01_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u32 offset);
DECLSPEC void append_0x06_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u32 offset);
DECLSPEC void append_0x01_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x2d_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x80_1x4_S (PRIVATE_AS u32 *w0, const u32 offset);
DECLSPEC void append_0x80_2x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, const u32 offset);
DECLSPEC void append_0x80_3x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, const u32 offset);
DECLSPEC void append_0x80_4x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void append_0x80_8x4_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *w4, PRIVATE_AS u32 *w5, PRIVATE_AS u32 *w6, PRIVATE_AS u32 *w7, const u32 offset);
DECLSPEC void make_utf16be_S (PRIVATE_AS const u32 *in, PRIVATE_AS u32 *out1, PRIVATE_AS u32 *out2);
DECLSPEC void make_utf16le_S (PRIVATE_AS const u32 *in, PRIVATE_AS u32 *out1, PRIVATE_AS u32 *out2);
DECLSPEC void undo_utf16be_S (PRIVATE_AS const u32 *in1, PRIVATE_AS const u32 *in2, PRIVATE_AS u32 *out);
DECLSPEC void undo_utf16le_S (PRIVATE_AS const u32 *in1, PRIVATE_AS const u32 *in2, PRIVATE_AS u32 *out);
DECLSPEC void switch_buffer_by_offset_le_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_carry_le_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *c0, PRIVATE_AS u32 *c1, PRIVATE_AS u32 *c2, PRIVATE_AS u32 *c3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_be_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_carry_be_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *c0, PRIVATE_AS u32 *c1, PRIVATE_AS u32 *c2, PRIVATE_AS u32 *c3, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_le_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *w4, PRIVATE_AS u32 *w5, PRIVATE_AS u32 *w6, PRIVATE_AS u32 *w7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_carry_le_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *w4, PRIVATE_AS u32 *w5, PRIVATE_AS u32 *w6, PRIVATE_AS u32 *w7, PRIVATE_AS u32 *c0, PRIVATE_AS u32 *c1, PRIVATE_AS u32 *c2, PRIVATE_AS u32 *c3, PRIVATE_AS u32 *c4, PRIVATE_AS u32 *c5, PRIVATE_AS u32 *c6, PRIVATE_AS u32 *c7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_be_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *w4, PRIVATE_AS u32 *w5, PRIVATE_AS u32 *w6, PRIVATE_AS u32 *w7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_8x4_carry_be_S (PRIVATE_AS u32 *w0, PRIVATE_AS u32 *w1, PRIVATE_AS u32 *w2, PRIVATE_AS u32 *w3, PRIVATE_AS u32 *w4, PRIVATE_AS u32 *w5, PRIVATE_AS u32 *w6, PRIVATE_AS u32 *w7, PRIVATE_AS u32 *c0, PRIVATE_AS u32 *c1, PRIVATE_AS u32 *c2, PRIVATE_AS u32 *c3, PRIVATE_AS u32 *c4, PRIVATE_AS u32 *c5, PRIVATE_AS u32 *c6, PRIVATE_AS u32 *c7, const u32 offset);
DECLSPEC void switch_buffer_by_offset_1x64_le_S (PRIVATE_AS u32 *w, const u32 offset);
DECLSPEC void switch_buffer_by_offset_1x64_be_S (PRIVATE_AS u32 *w, const u32 offset);
DECLSPEC void switch_buffer_by_offset_le_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
DECLSPEC void switch_buffer_by_offset_8x4_le_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, PRIVATE_AS u32x *w4, PRIVATE_AS u32x *w5, PRIVATE_AS u32x *w6, PRIVATE_AS u32x *w7, const u32x offset);
DECLSPEC void append_0x01_2x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, const u32x offset);
DECLSPEC void append_0x01_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
DECLSPEC void append_0x06_2x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, const u32x offset);
DECLSPEC void append_0x80_2x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, const u32x offset);
DECLSPEC void append_0x80_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
DECLSPEC void append_0x2d_4x4_VV (PRIVATE_AS u32x *w0, PRIVATE_AS u32x *w1, PRIVATE_AS u32x *w2, PRIVATE_AS u32x *w3, const u32x offset);
#endif // INC_COMMON_H