Skip to content

Commit 82ec652

Browse files
magnumrippermagnumripper
committed
Add iterated-sha1 format(s)
This is optionally salted ($s.$p) and optionally iterated. Canonical ciphertext: $sisha1$iter$<salt><hash> where salt length can be 1-16 bytes and inferred from total length (note: no field delimiter). Raw ciphertext: The last of the ciphertext is always the SHA-1 hash and anything before that is the salt. 4 bytes salt defaults to 1 iteration (this is XSHA). Other salt lengths default to 1024 iterations. No salt still means 1024 iterations but you can force this format to do single raw sha1 using $sisha1$1$<hash> because, well, why not?
1 parent c1bfb4c commit 82ec652

File tree

7 files changed

+1755
-15
lines changed

7 files changed

+1755
-15
lines changed

‎run/opencl/iterated_sha1_kernel.cl‎

Lines changed: 259 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,259 @@
1+
/*
2+
* Copyright (c) 2025, magnum
3+
* This software is hereby released to the general public under
4+
* the following terms: Redistribution and use in source and binary
5+
* forms, with or without modification, are permitted.
6+
*/
7+
8+
#include "opencl_device_info.h"
9+
#define AMD_PUTCHAR_NOCAST
10+
#include "opencl_misc.h"
11+
#include "opencl_mask.h"
12+
#include "opencl_sha1.h"
13+
14+
/* This handles an input of 0xffffffffU correctly */
15+
#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)
16+
17+
INLINE void cmp_final(uint gid,
18+
uint iter,
19+
uint *hash,
20+
__global uint *offset_table,
21+
__global uint *hash_table,
22+
__global uint *return_hashes,
23+
volatile __global uint *output,
24+
volatile __global uint *bitmap_dupe) {
25+
26+
uint t, offset_table_index, hash_table_index;
27+
ulong LO, MI, HI;
28+
ulong p;
29+
30+
HI = (ulong)hash[4];
31+
MI = ((ulong)hash[3] << 32) | (ulong)hash[2];
32+
LO = ((ulong)hash[1] << 32) | (ulong)hash[0];
33+
34+
p = (HI % OFFSET_TABLE_SIZE) * SHIFT128_OT_SZ;
35+
p += (MI % OFFSET_TABLE_SIZE) * SHIFT64_OT_SZ;
36+
p += LO % OFFSET_TABLE_SIZE;
37+
p %= OFFSET_TABLE_SIZE;
38+
offset_table_index = (uint)p;
39+
40+
//error: chances of overflow is extremely low.
41+
LO += (ulong)offset_table[offset_table_index];
42+
43+
p = (HI % HASH_TABLE_SIZE) * SHIFT128_HT_SZ;
44+
p += (MI % HASH_TABLE_SIZE) * SHIFT64_HT_SZ;
45+
p += LO % HASH_TABLE_SIZE;
46+
p %= HASH_TABLE_SIZE;
47+
hash_table_index = (uint)p;
48+
49+
if (hash_table[hash_table_index] == hash[0])
50+
if (hash_table[HASH_TABLE_SIZE + hash_table_index] == hash[1])
51+
{
52+
/*
53+
* Prevent duplicate keys from cracking same hash
54+
*/
55+
if (!(atomic_or(&bitmap_dupe[hash_table_index/32], (1U << (hash_table_index % 32))) & (1U << (hash_table_index % 32)))) {
56+
t = atomic_inc(&output[0]);
57+
output[1 + 3 * t] = gid;
58+
output[2 + 3 * t] = iter;
59+
output[3 + 3 * t] = hash_table_index;
60+
return_hashes[2 * t] = hash[2];
61+
return_hashes[2 * t + 1] = hash[3];
62+
}
63+
}
64+
}
65+
66+
INLINE void cmp(uint gid,
67+
uint iter,
68+
uint *hash,
69+
#if USE_LOCAL_BITMAPS
70+
__local
71+
#else
72+
__global
73+
#endif
74+
uint *bitmaps,
75+
__global uint *offset_table,
76+
__global uint *hash_table,
77+
__global uint *return_hashes,
78+
volatile __global uint *output,
79+
volatile __global uint *bitmap_dupe) {
80+
uint bitmap_index, tmp = 1;
81+
82+
hash[0] = hash[0];
83+
hash[1] = hash[1];
84+
hash[2] = hash[2];
85+
hash[3] = hash[3];
86+
hash[4] = hash[4];
87+
88+
#if SELECT_CMP_STEPS > 4
89+
bitmap_index = hash[0] & BITMAP_MASK;
90+
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
91+
bitmap_index = (hash[0] >> 16) & BITMAP_MASK;
92+
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
93+
bitmap_index = hash[1] & BITMAP_MASK;
94+
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
95+
bitmap_index = (hash[1] >> 16) & BITMAP_MASK;
96+
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
97+
bitmap_index = hash[2] & BITMAP_MASK;
98+
tmp &= (bitmaps[BITMAP_SHIFT * 4 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
99+
bitmap_index = (hash[2] >> 16) & BITMAP_MASK;
100+
tmp &= (bitmaps[BITMAP_SHIFT * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
101+
bitmap_index = hash[3] & BITMAP_MASK;
102+
tmp &= (bitmaps[BITMAP_SHIFT * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
103+
bitmap_index = (hash[3] >> 16) & BITMAP_MASK;
104+
tmp &= (bitmaps[BITMAP_SHIFT * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
105+
#elif SELECT_CMP_STEPS > 2
106+
bitmap_index = hash[3] & BITMAP_MASK;
107+
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
108+
bitmap_index = hash[2] & BITMAP_MASK;
109+
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
110+
bitmap_index = hash[1] & BITMAP_MASK;
111+
tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
112+
bitmap_index = hash[0] & BITMAP_MASK;
113+
tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
114+
#elif SELECT_CMP_STEPS > 1
115+
bitmap_index = hash[3] & BITMAP_MASK;
116+
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
117+
bitmap_index = hash[2] & BITMAP_MASK;
118+
tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
119+
#else
120+
bitmap_index = hash[3] & BITMAP_MASK;
121+
tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
122+
#endif
123+
124+
if (tmp)
125+
cmp_final(gid, iter, hash, offset_table, hash_table, return_hashes, output, bitmap_dupe);
126+
}
127+
128+
#define USE_CONST_CACHE \
129+
(CONST_CACHE_SIZE >= (NUM_INT_KEYS * 4))
130+
131+
typedef struct {
132+
uint iter;
133+
uint len;
134+
uchar salt[MAX_SALT_SIZE];
135+
} salt_t;
136+
137+
__kernel
138+
void sha1(__global uint *keys,
139+
__global uint *index,
140+
__global uint *int_key_loc,
141+
#if USE_CONST_CACHE
142+
constant
143+
#else
144+
__global
145+
#endif
146+
uint *int_keys,
147+
__constant salt_t *salt,
148+
__global uint *bitmaps,
149+
__global uint *offset_table,
150+
__global uint *hash_table,
151+
__global uint *return_hashes,
152+
volatile __global uint *out_hash_ids,
153+
volatile __global uint *bitmap_dupe)
154+
{
155+
uint i;
156+
uint gid = get_global_id(0);
157+
uint base = index[gid];
158+
uint T[16] = { 0 };
159+
uint len = base & 63;
160+
161+
#if NUM_INT_KEYS > 1 && !IS_STATIC_GPU_MASK
162+
uint ikl = int_key_loc[gid];
163+
uint loc0 = ikl & 0xff;
164+
#if MASK_FMT_INT_PLHDR > 1
165+
#if LOC_1 >= 0
166+
uint loc1 = (ikl & 0xff00) >> 8;
167+
#endif
168+
#endif
169+
#if MASK_FMT_INT_PLHDR > 2
170+
#if LOC_2 >= 0
171+
uint loc2 = (ikl & 0xff0000) >> 16;
172+
#endif
173+
#endif
174+
#if MASK_FMT_INT_PLHDR > 3
175+
#if LOC_3 >= 0
176+
uint loc3 = (ikl & 0xff000000) >> 24;
177+
#endif
178+
#endif
179+
#endif
180+
181+
#if !IS_STATIC_GPU_MASK
182+
#define GPU_LOC_0 loc0
183+
#define GPU_LOC_1 loc1
184+
#define GPU_LOC_2 loc2
185+
#define GPU_LOC_3 loc3
186+
#else
187+
#define GPU_LOC_0 LOC_0
188+
#define GPU_LOC_1 LOC_1
189+
#define GPU_LOC_2 LOC_2
190+
#define GPU_LOC_3 LOC_3
191+
#endif
192+
193+
#if USE_LOCAL_BITMAPS
194+
uint lid = get_local_id(0);
195+
uint lws = get_local_size(0);
196+
__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];
197+
198+
for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
199+
s_bitmaps[i] = bitmaps[i];
200+
201+
barrier(CLK_LOCAL_MEM_FENCE);
202+
#endif
203+
204+
keys += base >> 6;
205+
206+
for (i = 0; i < salt->len; i++)
207+
PUTCHAR_BE(T, i, salt->salt[i]);
208+
209+
__global uchar *key = (__global uchar*)keys;
210+
for (; i < salt->len + len; i++)
211+
PUTCHAR_BE(T, i, *key++);
212+
213+
PUTCHAR_BE(T, (salt->len + len), 0x80);
214+
T[15] = (salt->len + len) << 3;
215+
216+
for (uint idx = 0; idx < NUM_INT_KEYS; idx++) {
217+
#if NUM_INT_KEYS > 1
218+
PUTCHAR_BE(T, salt->len + GPU_LOC_0, (int_keys[idx] & 0xff));
219+
220+
#if MASK_FMT_INT_PLHDR > 1
221+
#if LOC_1 >= 0
222+
PUTCHAR_BE(T, salt->len + GPU_LOC_1, ((int_keys[idx] & 0xff00) >> 8));
223+
#endif
224+
#endif
225+
#if MASK_FMT_INT_PLHDR > 2
226+
#if LOC_2 >= 0
227+
PUTCHAR_BE(T, salt->len + GPU_LOC_2, ((int_keys[idx] & 0xff0000) >> 16));
228+
#endif
229+
#endif
230+
#if MASK_FMT_INT_PLHDR > 3
231+
#if LOC_3 >= 0
232+
PUTCHAR_BE(T, salt->len + GPU_LOC_3, ((int_keys[idx] & 0xff000000) >> 24));
233+
#endif
234+
#endif
235+
#endif
236+
uint W[16];
237+
uint hash[5];
238+
239+
memcpy_macro(W, T, 16);
240+
241+
sha1_single(uint, W, hash);
242+
243+
uint iter = salt->iter;
244+
while (--iter) {
245+
memcpy_macro(W, hash, 5);
246+
W[5] = 0x80000000;
247+
W[15] = 20 << 3;
248+
sha1_single_160Z(uint, W, hash);
249+
}
250+
251+
cmp(gid, idx, hash,
252+
#if USE_LOCAL_BITMAPS
253+
s_bitmaps
254+
#else
255+
bitmaps
256+
#endif
257+
, offset_table, hash_table, return_hashes, out_hash_ids, bitmap_dupe);
258+
}
259+
}

‎src/common.c‎

Lines changed: 20 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -225,3 +225,23 @@ int isdecu(const char *q)
225225
return 0;
226226
return isdec_len(q, "4294967295");
227227
}
228+
229+
int getdec(const char *string, char separator)
230+
{
231+
int i = 0;
232+
char buf[11];
233+
234+
while (string[i] >= '0' && string[i] <= '9') {
235+
buf[i] = string[i];
236+
if (++i >= sizeof(buf) - 1)
237+
break;
238+
}
239+
if (i == 0 || string[i] != separator)
240+
return -1;
241+
buf[i] = 0;
242+
243+
if (!isdec(buf))
244+
return -1;
245+
246+
return atoi(buf);
247+
}

‎src/common.h‎

Lines changed: 25 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -108,38 +108,48 @@ extern void common_init(void);
108108
**************************************************************/
109109

110110
/* is string full valid hex string */
111-
int ishex(const char *q);
111+
extern int ishex(const char *q);
112112
/* Same as ishex(), BUT will still return true for an odd length string */
113-
int ishex_oddOK(const char *q);
113+
extern int ishex_oddOK(const char *q);
114114
/* is string full valid hex string (only upper case letters) */
115-
int ishexuc(const char *q);
115+
extern int ishexuc(const char *q);
116116
/* is string full valid hex string (only lower case letters) */
117-
int ishexlc(const char *q);
117+
extern int ishexlc(const char *q);
118118
/* same as ishexuc/lc except odd length is ok */
119-
int ishexuc_oddOK(const char *q);
120-
int ishexlc_oddOK(const char *q);
119+
extern int ishexuc_oddOK(const char *q);
120+
extern int ishexlc_oddOK(const char *q);
121121
/* provide a length field, so return true if 'n' bytes of the string are hex */
122122
/* the n is length q, so there is no need for a 'odd' field. If checking for */
123123
/* a 49 byte string, simply specify 49 */
124-
int ishexn(const char *q, int n);
125-
int ishexucn(const char *q, int n);
126-
int ishexlcn(const char *q, int n);
124+
extern int ishexn(const char *q, int n);
125+
extern int ishexucn(const char *q, int n);
126+
extern int ishexlcn(const char *q, int n);
127127
/* length of hex. if extra_chars not null, it will be 1 if there are more
128128
* non-hex characters after the length of valid hex chars returned.
129129
* NOTE, the return will always be an even number (rounded down). so if we
130130
* want the length of "ABCDE", it will be 4 not 5.
131131
*/
132-
size_t hexlen(const char *q, int *extra_chars);
133-
size_t hexlenl(const char *q, int *extra_chars); /* lower cased only */
134-
size_t hexlenu(const char *q, int *extra_chars); /* upper cased only */
132+
extern size_t hexlen(const char *q, int *extra_chars);
133+
extern size_t hexlenl(const char *q, int *extra_chars); /* lower cased only */
134+
extern size_t hexlenu(const char *q, int *extra_chars); /* upper cased only */
135135
/* Is this a valid number <=10digits and in the range [0 .... <= 0x7fffffff]
136136
* ONLY positive numbers are valid. */
137-
int isdec(const char *q);
137+
extern int isdec(const char *q);
138138
/* Is this a valid number <=10digits.
139139
* Positive [0..<= 0x7fffffff] and negative [ <= 0x80000000] numbers are valid */
140-
int isdec_negok(const char *q);
140+
extern int isdec_negok(const char *q);
141141
/* Is this a valid number <=10digits.ONLY positive [0..<=0xffffffff] numbers are valid */
142-
int isdecu(const char *q);
142+
extern int isdecu(const char *q);
143+
144+
/*
145+
* Read integer from string, eg. "1024$deadcafe", up to the separator.
146+
*
147+
* Returns the positive integer, eg. 1024.
148+
* If separator doesn't match or decimal is not positive int32_t, return -1.
149+
* Note that an end of string instead of the separator is a non-match but if
150+
* you do want to read to an expected end of a string, separator can be '\0'!
151+
*/
152+
extern int getdec(const char *string, char separator);
143153

144154
#endif
145155

‎src/iterated_sha1_common.h‎

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,47 @@
1+
/*
2+
* Copyright (c) 2025, magnum
3+
* Redistribution and use in source and binary forms, with or without
4+
* modification, are permitted.
5+
*/
6+
7+
#ifndef ITERATED_SHA1_COMMON_H__
8+
#define ITERATED_SHA1_COMMON_H__
9+
10+
#include "formats.h"
11+
12+
#define FORMAT_NAME "salted"
13+
14+
#define FORMAT_TAG "$sisha1$"
15+
#define FORMAT_TAG_LEN (sizeof(FORMAT_TAG) - 1)
16+
17+
#define BINARY_SIZE 20
18+
#define BINARY_ALIGN 4
19+
#define SALT_SIZE sizeof(salt_t)
20+
#define SALT_ALIGN 4
21+
22+
#define BENCHMARK_COMMENT ""
23+
#define BENCHMARK_LENGTH 0x507
24+
25+
#define MAX_SALT_SIZE 16 // 16 bytes, 32 hex chars
26+
#define PLAINTEXT_LENGTH (55 - MAX_SALT_SIZE)
27+
#define MIN_CIPHERTEXT_LEN (160 / 4)
28+
#define MAX_CIPHERTEXT_LEN (160 / 4 + 2 * MAX_SALT_SIZE)
29+
30+
#define MAX_ITER (1 << 20)
31+
32+
typedef struct {
33+
uint32_t iter;
34+
uint32_t len;
35+
uint32_t salt[MAX_SALT_SIZE / sizeof(uint32_t)];
36+
} salt_t;
37+
38+
extern struct fmt_tests iterated_sha1_tests[];
39+
40+
extern char *iterated_sha1_prepare(char *fields[10], struct fmt_main *self);
41+
extern int iterated_sha1_valid(char *ciphertext, struct fmt_main *self);
42+
extern void* iterated_sha1_get_binary(char* ciphertext);
43+
extern void* iterated_sha1_get_salt(char* ciphertext);
44+
extern int iterated_sha1_salt_hash(void *salt);
45+
extern unsigned int iterated_sha1_iterations(void *salt);
46+
47+
#endif // ITERATED_SHA1_COMMON_H__

0 commit comments

Comments
 (0)