Iterated sha1

run/opencl/iterated_sha1_kernel.cl

@@ -0,0 +1,259 @@
+/*
+ * Copyright (c) 2025, magnum
+ * This software is hereby released to the general public under
+ * the following terms: Redistribution and use in source and binary
+ * forms, with or without modification, are permitted.
+ */
+
+#include "opencl_device_info.h"
+#define AMD_PUTCHAR_NOCAST
+#include "opencl_misc.h"
+#include "opencl_mask.h"
+#include "opencl_sha1.h"
+
+/* This handles an input of 0xffffffffU correctly */
+#define BITMAP_SHIFT ((BITMAP_MASK >> 5) + 1)
+
+INLINE void cmp_final(uint gid,
+		uint iter,
+		uint *hash,
+		__global uint *offset_table,
+		__global uint *hash_table,
+		__global uint *return_hashes,
+		volatile __global uint *output,
+		volatile __global uint *bitmap_dupe) {
+
+	uint t, offset_table_index, hash_table_index;
+	ulong LO, MI, HI;
+	ulong p;
+
+	HI = (ulong)hash[4];
+	MI = ((ulong)hash[3] << 32) | (ulong)hash[2];
+	LO = ((ulong)hash[1] << 32) | (ulong)hash[0];
+
+	p = (HI % OFFSET_TABLE_SIZE) * SHIFT128_OT_SZ;
+	p += (MI % OFFSET_TABLE_SIZE) * SHIFT64_OT_SZ;
+	p += LO % OFFSET_TABLE_SIZE;
+	p %= OFFSET_TABLE_SIZE;
+	offset_table_index = (uint)p;
+
+	//error: chances of overflow is extremely low.
+	LO += (ulong)offset_table[offset_table_index];
+
+	p = (HI % HASH_TABLE_SIZE) * SHIFT128_HT_SZ;
+	p += (MI % HASH_TABLE_SIZE) * SHIFT64_HT_SZ;
+	p += LO % HASH_TABLE_SIZE;
+	p %= HASH_TABLE_SIZE;
+	hash_table_index = (uint)p;
+
+	if (hash_table[hash_table_index] == hash[0])
+	if (hash_table[HASH_TABLE_SIZE + hash_table_index] == hash[1])
+	{
+/*
+ * Prevent duplicate keys from cracking same hash
+ */
+		if (!(atomic_or(&bitmap_dupe[hash_table_index/32], (1U << (hash_table_index % 32))) & (1U << (hash_table_index % 32)))) {
+			t = atomic_inc(&output[0]);
+			output[1 + 3 * t] = gid;
+			output[2 + 3 * t] = iter;
+			output[3 + 3 * t] = hash_table_index;
+			return_hashes[2 * t] = hash[2];
+			return_hashes[2 * t + 1] = hash[3];
+		}
+	}
+}
+
+INLINE void cmp(uint gid,
+		uint iter,
+		uint *hash,
+#if USE_LOCAL_BITMAPS
+		__local
+#else
+		__global
+#endif
+		uint *bitmaps,
+		__global uint *offset_table,
+		__global uint *hash_table,
+		__global uint *return_hashes,
+		volatile __global uint *output,
+		volatile __global uint *bitmap_dupe) {
+	uint bitmap_index, tmp = 1;
+
+	hash[0] = hash[0];
+	hash[1] = hash[1];
+	hash[2] = hash[2];
+	hash[3] = hash[3];
+	hash[4] = hash[4];
+
+#if SELECT_CMP_STEPS > 4
+	bitmap_index = hash[0] & BITMAP_MASK;
+	tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = (hash[0] >> 16) & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[1] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = (hash[1] >> 16) & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[2] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 4 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = (hash[2] >> 16) & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 5 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[3] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 6 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = (hash[3] >> 16) & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 7 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+#elif SELECT_CMP_STEPS > 2
+	bitmap_index = hash[3] & BITMAP_MASK;
+	tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[2] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[1] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 2 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[0] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT * 3 + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+#elif SELECT_CMP_STEPS > 1
+	bitmap_index = hash[3] & BITMAP_MASK;
+	tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
+	bitmap_index = hash[2] & BITMAP_MASK;
+	tmp &= (bitmaps[BITMAP_SHIFT + (bitmap_index >> 5)] >> (bitmap_index & 31)) & 1U;
+#else
+	bitmap_index = hash[3] & BITMAP_MASK;
+	tmp &= (bitmaps[bitmap_index >> 5] >> (bitmap_index & 31)) & 1U;
+#endif
+
+	if (tmp)
+		cmp_final(gid, iter, hash, offset_table, hash_table, return_hashes, output, bitmap_dupe);
+}
+
+#define USE_CONST_CACHE \
+	(CONST_CACHE_SIZE >= (NUM_INT_KEYS * 4))
+
+typedef struct {
+	uint iter;
+	uint len;
+	uchar salt[MAX_SALT_SIZE];
+} salt_t;
+
+__kernel
+void sha1(__global uint *keys,
+          __global uint *index,
+          __global uint *int_key_loc,
+#if USE_CONST_CACHE
+          constant
+#else
+          __global
+#endif
+          uint *int_keys,
+          __constant salt_t *salt,
+          __global uint *bitmaps,
+          __global uint *offset_table,
+          __global uint *hash_table,
+          __global uint *return_hashes,
+          volatile __global uint *out_hash_ids,
+          volatile __global uint *bitmap_dupe)
+{
+	uint i;
+	uint gid = get_global_id(0);
+	uint base = index[gid];
+	uint T[16] = { 0 };
+	uint len = base & 63;
+
+#if NUM_INT_KEYS > 1 && !IS_STATIC_GPU_MASK
+	uint ikl = int_key_loc[gid];
+	uint loc0 = ikl & 0xff;
+#if MASK_FMT_INT_PLHDR > 1
+#if LOC_1 >= 0
+	uint loc1 = (ikl & 0xff00) >> 8;
+#endif
+#endif
+#if MASK_FMT_INT_PLHDR > 2
+#if LOC_2 >= 0
+	uint loc2 = (ikl & 0xff0000) >> 16;
+#endif
+#endif
+#if MASK_FMT_INT_PLHDR > 3
+#if LOC_3 >= 0
+	uint loc3 = (ikl & 0xff000000) >> 24;
+#endif
+#endif
+#endif
+
+#if !IS_STATIC_GPU_MASK
+#define GPU_LOC_0 loc0
+#define GPU_LOC_1 loc1
+#define GPU_LOC_2 loc2
+#define GPU_LOC_3 loc3
+#else
+#define GPU_LOC_0 LOC_0
+#define GPU_LOC_1 LOC_1
+#define GPU_LOC_2 LOC_2
+#define GPU_LOC_3 LOC_3
+#endif
+
+#if USE_LOCAL_BITMAPS
+	uint lid = get_local_id(0);
+	uint lws = get_local_size(0);
+	__local uint s_bitmaps[BITMAP_SHIFT * SELECT_CMP_STEPS];
+
+	for (i = lid; i < BITMAP_SHIFT * SELECT_CMP_STEPS; i+= lws)
+		s_bitmaps[i] = bitmaps[i];
+
+	barrier(CLK_LOCAL_MEM_FENCE);
+#endif
+
+	keys += base >> 6;
+
+	for (i = 0; i < salt->len; i++)
+		PUTCHAR_BE(T, i, salt->salt[i]);
+
+	__global uchar *key = (__global uchar*)keys;
+	for (; i < salt->len + len; i++)
+		PUTCHAR_BE(T, i, *key++);
+
+	PUTCHAR_BE(T, (salt->len + len), 0x80);
+	T[15] = (salt->len + len) << 3;
+
+	for (uint idx = 0; idx < NUM_INT_KEYS; idx++) {
+#if NUM_INT_KEYS > 1
+		PUTCHAR_BE(T, salt->len + GPU_LOC_0, (int_keys[idx] & 0xff));
+
+#if MASK_FMT_INT_PLHDR > 1
+#if LOC_1 >= 0
+		PUTCHAR_BE(T, salt->len + GPU_LOC_1, ((int_keys[idx] & 0xff00) >> 8));
+#endif
+#endif
+#if MASK_FMT_INT_PLHDR > 2
+#if LOC_2 >= 0
+		PUTCHAR_BE(T, salt->len + GPU_LOC_2, ((int_keys[idx] & 0xff0000) >> 16));
+#endif
+#endif
+#if MASK_FMT_INT_PLHDR > 3
+#if LOC_3 >= 0
+		PUTCHAR_BE(T, salt->len + GPU_LOC_3, ((int_keys[idx] & 0xff000000) >> 24));
+#endif
+#endif
+#endif
+		uint W[16];
+		uint hash[5];
+
+		memcpy_macro(W, T, 16);
+
+		sha1_single(uint, W, hash);
+
+		uint iter = salt->iter;
+		while (--iter) {
+			memcpy_macro(W, hash, 5);
+			W[5] = 0x80000000;
+			W[15] = 20 << 3;
+			sha1_single_160Z(uint, W, hash);
+		}
+
+		cmp(gid, idx, hash,
+#if USE_LOCAL_BITMAPS
+		    s_bitmaps
+#else
+		    bitmaps
+#endif
+		    , offset_table, hash_table, return_hashes, out_hash_ids, bitmap_dupe);
+	}
+}

src/XSHA512_fmt_plug.c

@@ -24,7 +24,7 @@ john_register_one(&fmt_XSHA512);
 
 #define FORMAT_LABEL			"xsha512"
 #define FORMAT_NAME			"Mac OS X 10.7"
-#define ALGORITHM_NAME			"SHA512 " SHA512_ALGORITHM_NAME
+#define ALGORITHM_NAME			"SHA512 ($s.$p) " SHA512_ALGORITHM_NAME
 
 #define PLAINTEXT_LENGTH		107
 

src/XSHA_fmt_plug.c

@@ -37,7 +37,7 @@ static unsigned int threads = 1;
 #define FORMAT_LABEL			"xsha"
 #define FORMAT_NAME			"Mac OS X 10.4 - 10.6"
 
-#define ALGORITHM_NAME			"SHA1 " SHA1_ALGORITHM_NAME
+#define ALGORITHM_NAME			"SHA1 ($s.$p)" SHA1_ALGORITHM_NAME
 
 #define BENCHMARK_COMMENT		""
 #define BENCHMARK_LENGTH		7

src/common.c

@@ -225,3 +225,23 @@ int isdecu(const char *q)
 		return 0;
 	return isdec_len(q, "4294967295");
 }
+
+int getdec(const char *string, char separator)
+{
+	int i = 0;
+	char buf[11];
+
+	while (string[i] >= '0' && string[i] <= '9') {
+		buf[i] = string[i];
+		if (++i >= sizeof(buf) - 1)
+			break;
+	}
+	if (i == 0 || string[i] != separator)
+		return -1;
+	buf[i] = 0;
+
+	if (!isdec(buf))
+		return -1;
+
+	return atoi(buf);
+}

src/common.h

@@ -108,38 +108,48 @@ extern void common_init(void);
  **************************************************************/
 
 /* is string full valid hex string */
-int ishex(const char *q);
+extern int ishex(const char *q);
 /* Same as ishex(), BUT will still return true for an odd length string */
-int ishex_oddOK(const char *q);
+extern int ishex_oddOK(const char *q);
 /* is string full valid hex string (only upper case letters) */
-int ishexuc(const char *q);
+extern int ishexuc(const char *q);
 /* is string full valid hex string (only lower case letters) */
-int ishexlc(const char *q);
+extern int ishexlc(const char *q);
 /* same as ishexuc/lc except odd length is ok */
-int ishexuc_oddOK(const char *q);
-int ishexlc_oddOK(const char *q);
+extern int ishexuc_oddOK(const char *q);
+extern int ishexlc_oddOK(const char *q);
 /* provide a length field, so return true if 'n' bytes of the string are hex */
 /* the n is length q, so there is no need for a 'odd' field. If checking for */
 /* a 49 byte string, simply specify 49 */
-int ishexn(const char *q, int n);
-int ishexucn(const char *q, int n);
-int ishexlcn(const char *q, int n);
+extern int ishexn(const char *q, int n);
+extern int ishexucn(const char *q, int n);
+extern int ishexlcn(const char *q, int n);
 /* length of hex. if extra_chars not null, it will be 1 if there are more
  * non-hex characters after the length of valid hex chars returned.
  * NOTE, the return will always be an even number (rounded down). so if we
  * want the length of "ABCDE", it will be 4 not 5.
  */
-size_t hexlen(const char *q, int *extra_chars);
-size_t hexlenl(const char *q, int *extra_chars); /* lower cased only */
-size_t hexlenu(const char *q, int *extra_chars); /* upper cased only */
+extern size_t hexlen(const char *q, int *extra_chars);
+extern size_t hexlenl(const char *q, int *extra_chars); /* lower cased only */
+extern size_t hexlenu(const char *q, int *extra_chars); /* upper cased only */
 /* Is this a valid number <=10digits and in the range [0 .... <= 0x7fffffff]
  * ONLY positive numbers are valid. */
-int isdec(const char *q);
+extern int isdec(const char *q);
 /* Is this a valid number <=10digits.
  * Positive [0..<= 0x7fffffff] and negative [ <= 0x80000000] numbers are valid */
-int isdec_negok(const char *q);
+extern int isdec_negok(const char *q);
 /* Is this a valid number <=10digits.ONLY positive [0..<=0xffffffff] numbers are valid */
-int isdecu(const char *q);
+extern int isdecu(const char *q);
+
+/*
+ * Read integer from string, eg. "1024$deadcafe", up to the separator.
+ *
+ * Returns the positive integer, eg. 1024.
+ * If separator doesn't match or decimal is not positive int32_t, return -1.
+ * Note that an end of string instead of the separator is a non-match but if
+ * you do want to read to an expected end of a string, separator can be '\0'!
+ */
+extern int getdec(const char *string, char separator);
 
 #endif