From cf23a254899de61fd39e60f614bbecb654c8af20 Mon Sep 17 00:00:00 2001 From: magnum Date: Mon, 30 Jan 2023 08:46:17 +0100 Subject: [PATCH] NT-long-OpenCL: Support lengths up to 125 as a separate format This way we still got optimal speed for the original nt-opencl format that support lengths up to 27. 1 block of MD4 is up to 27 characters. 2 blocks is 59, 3 is 91, 4 is 123 and 5 is 125 (due to core max). As long as we bump it over 27 we don't seem to gain any speed by limiting it to less than 125 characters. See #5245 --- run/opencl/nt_kernel.cl | 48 +++++++++++++-- src/opencl_nt_fmt_plug.c | 129 ++++++++++++++++++++++++++++++++++++--- 2 files changed, 166 insertions(+), 11 deletions(-) diff --git a/run/opencl/nt_kernel.cl b/run/opencl/nt_kernel.cl index 31c9fcca4d8..d3351083492 100644 --- a/run/opencl/nt_kernel.cl +++ b/run/opencl/nt_kernel.cl @@ -8,7 +8,7 @@ * Copyright (c) 2010 Alain Espinosa * Copyright (c) 2011 Samuele Giovanni Tonon * Copyright (c) 2015 Sayantan Datta - * Copyright (c) 2015 magnum + * Copyright (c) 2015-2023 magnum * and it is hereby released to the general public under the following terms: * * Redistribution and use in source and binary forms, with or without @@ -43,6 +43,8 @@ inline void nt_crypt(uint *hash, uint *nt_buffer, uint md4_size) { uint tmp; + md4_size <<= 4; + /* Round 1 */ hash[0] = 0xFFFFFFFF + nt_buffer[0]; hash[0]=rotate(hash[0], 3u); hash[3] = INIT_D + (INIT_C ^ (hash[0] & 0x77777777)) + nt_buffer[1]; hash[3]=rotate(hash[3], 7u); @@ -110,6 +112,40 @@ inline void nt_crypt(uint *hash, uint *nt_buffer, uint md4_size) { hash[2] += MD4_H(tmp, hash[0], hash[3]) + nt_buffer[7] + SQRT_3; hash[2] = rotate(hash[2] , 11u); } +#if PLAINTEXT_LENGTH > 27 +inline void md4_reverse(uint *hash) +{ + hash[0] -= INIT_A; + hash[1] -= INIT_B; + hash[2] -= INIT_C; + hash[3] -= INIT_D; + hash[1] = (hash[1] >> 15) | (hash[1] << 17); + hash[1] -= SQRT_3 + (hash[2] ^ hash[3] ^ hash[0]); + hash[1] = (hash[1] >> 15) | (hash[1] << 17); + hash[1] -= SQRT_3; +} + +inline void nt_crypt_long(uint *hash, uint *nt_buffer, uint md4_size) +{ + md4_init(hash); + + uint blocks = (md4_size + 5 + 31) / 32; + while (--blocks) { + md4_block(uint, nt_buffer, hash); + nt_buffer += 16; + } + + nt_buffer[14] = md4_size << 4; + md4_block(uint, nt_buffer, hash); + + /* + * This *adds* a little work to long crypts instead + * of losing the real reverse for single block crypts. + */ + md4_reverse(hash); +} +#endif + #if __OS_X__ && (cpu(DEVICE_INFO) || gpu_nvidia(DEVICE_INFO)) /* This is a workaround for driver/runtime bugs */ #define MAYBE_VOLATILE volatile @@ -341,7 +377,7 @@ __kernel void nt(__global uint *keys, uint i; uint gid = get_global_id(0); uint base = index[gid]; - uint nt_buffer[14] = { 0 }; + uint nt_buffer[(PLAINTEXT_LENGTH + 5 + 31) / 32 * 16] = { 0 }; uint md4_size = base & 127; uint hash[4]; @@ -390,7 +426,6 @@ __kernel void nt(__global uint *keys, keys += base >> 7; md4_size = prepare_key(keys, md4_size, nt_buffer); - md4_size = md4_size << 4; for (i = 0; i < NUM_INT_KEYS; i++) { #if NUM_INT_KEYS > 1 @@ -411,7 +446,12 @@ __kernel void nt(__global uint *keys, #endif #endif #endif - nt_crypt(hash, nt_buffer, md4_size); +#if PLAINTEXT_LENGTH > 27 + if (md4_size > 27) + nt_crypt_long(hash, nt_buffer, md4_size); + else +#endif + nt_crypt(hash, nt_buffer, md4_size); cmp(gid, i, hash, #if USE_LOCAL_BITMAPS s_bitmaps diff --git a/src/opencl_nt_fmt_plug.c b/src/opencl_nt_fmt_plug.c index 75f2bd6f6d1..3389216e8d4 100644 --- a/src/opencl_nt_fmt_plug.c +++ b/src/opencl_nt_fmt_plug.c @@ -9,7 +9,7 @@ * Copyright (c) 2010 Alain Espinosa * Copyright (c) 2011 Samuele Giovanni Tonon * Copyright (c) 2015 Sayantan Datta - * Copyright (c) 2015 magnum + * Copyright (c) 2015-2023 magnum * and it is hereby released to the general public under the following terms: * * Redistribution and use in source and binary forms, with or without @@ -24,8 +24,10 @@ #if FMT_EXTERNS_H extern struct fmt_main fmt_opencl_NT; +extern struct fmt_main fmt_opencl_NT_long; #elif FMT_REGISTERS_H john_register_one(&fmt_opencl_NT); +john_register_one(&fmt_opencl_NT_long); #else #include @@ -44,6 +46,7 @@ john_register_one(&fmt_opencl_NT); #include "opencl_hash_check_128.h" #define FORMAT_LABEL "NT-opencl" +#define FORMAT_LABEL_LONG "NT-long-opencl" #define FORMAT_NAME "" #define FORMAT_TAG "$NT$" #define FORMAT_TAG_LEN (sizeof(FORMAT_TAG)-1) @@ -51,8 +54,9 @@ john_register_one(&fmt_opencl_NT); #define BENCHMARK_COMMENT "" #define BENCHMARK_LENGTH 0x107 #define PLAINTEXT_LENGTH 27 +#define PLAINTEXT_LEN_LONG 125 /* 59, 91, 123, 125 are supported */ /* At most 3 bytes of UTF-8 needed per character */ -#define UTF8_MAX_LENGTH (3 * PLAINTEXT_LENGTH) +#define UTF8_MAX_LENGTH MIN(125, 3 * utf16len) #define BUFSIZE ((UTF8_MAX_LENGTH + 3) / 4 * 4) #define AUTOTUNE_LENGTH 8 #define CIPHERTEXT_LENGTH 32 @@ -110,9 +114,55 @@ static struct fmt_tests tests[] = { {"$NT$dd555241a4321657e8b827a40b67dd4a", "jordan"}, {"$NT$bb53a477af18526ada697ce2e51f76b3", "michael"}, {"$NT$92b7b06bb313bf666640c5a1e75e0c18", "michelle"}, + {"$NT$0ae2ac07ba42fb76e0d9e5852d00e83f", "xxxxxxxxxxxxxxxxxxxxxxxxxxx"}, {NULL} }; +static struct fmt_tests tests_long[] = { + {"8846f7eaee8fb117ad06bdd830b7586c", "password"}, + {"$NT$31d6cfe0d16ae931b73c59d7e0c089c0", ""}, + {"$NT$31d6cfe0d16ae931b73c59d7e0c089c0", ""}, + {"$NT$31d6cfe0d16ae931b73c59d7e0c089c0", ""}, + {"$NT$31d6cfe0d16ae931b73c59d7e0c089c0", ""}, + {"$NT$31d6cfe0d16ae931b73c59d7e0c089c0", ""}, + {"$NT$7a21990fcd3d759941e45c490f143d5f", "12345"}, + {"$NT$0ae2ac07ba42fb76e0d9e5852d00e83f", "xxxxxxxxxxxxxxxxxxxxxxxxxxx"}, +#if PLAINTEXT_LEN_LONG > 27 + {"$NT$e4e10a22597efd64ad85ec18c948cbf2", "xxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$ea1be9a74e6a7ca800ba932293aa2d6d", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$ec1814e21f7f5bed537fbab2e357bb60", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$59c7fe1adebc59dcb0f513cbed87eb92", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$59c1f6430d9d1aea6d9212f4cb6ea3ea", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$ce4e33c64e8fa0084ef3974a8c8ece59", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$4ccdaf5b8534ffc158b96e55669314a3", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, +#endif /* 28..59 */ +#if PLAINTEXT_LEN_LONG > 59 + {"$NT$889359447c0a6a784f8736e76326ce51", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$9774270d26d6ff5539326a7a39ae4b7a", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$15f9da1d28df9e7088bba11c0977a201", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$f99dddb5d218f258fe041d3a9079ff2a", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$b1b3356e5e05046ffe6d3f87ae2f8c12", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$5a4880f7cfcaa3bbfc7c6b4ca920970d", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$6c68a921eba1cce568cbe543ecf106ad", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, +#endif /* 60..91 */ +#if PLAINTEXT_LEN_LONG > 91 + {"$NT$2e99a04f4626ca16d1e40879ece1977e", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$8f83db5a44550b592b7de15b16939a30", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$d756304967ef98371509fa3150b018aa", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$6a66155ad4ebaab10267ddffca3cf2d3", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$01d3b7a7958e26e914ea851a2a3e882c", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$8924aa73dd0ce16a37bdec6edb3c8802", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$afb0648a73bd1e9662ff672251f80f63", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, +#endif /* 92..123 */ +#if PLAINTEXT_LEN_LONG > 123 + {"$NT$8f637d62e1d14cb97ca150c89222937a", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, + {"$NT$9d10efd08eb95db46f9a2badb2a71fcb", "xxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx"}, +#endif /* 124..125 */ + {NULL} +}; + +static int utf16len; /* We mess with plaintext_length in case of UTF-8 so have to keep track of the original */ + //Init values #define INIT_A 0x67452301 #define INIT_B 0xefcdab89 @@ -303,7 +353,7 @@ static void init_kernel(unsigned int num_ld_hashes, char *bitmap_para) num_ld_hashes, mask_int_cand.num_int_cand, bitmap_para, mask_gpu_is_static, (unsigned long long)const_cache_size, cp_id2macro(options.target_enc), options.internal_cp == UTF_8 ? cp_id2macro(ENC_RAW) : - cp_id2macro(options.internal_cp), PLAINTEXT_LENGTH, + cp_id2macro(options.internal_cp), utf16len, static_gpu_locations[0] #if MASK_FMT_INT_PLHDR > 1 , static_gpu_locations[1] @@ -323,6 +373,9 @@ static void init_kernel(unsigned int num_ld_hashes, char *bitmap_para) static void init(struct fmt_main *_self) { + if (!utf16len || _self != self) + utf16len = _self->params.plaintext_length; + self = _self; num_loaded_hashes = 0; @@ -331,7 +384,7 @@ static void init(struct fmt_main *_self) opencl_prepare_dev(gpu_id); mask_int_cand_target = opencl_speed_index(gpu_id) / 300; if (options.target_enc == UTF_8) { - self->params.plaintext_length = MIN(125, UTF8_MAX_LENGTH); + self->params.plaintext_length = UTF8_MAX_LENGTH; tests[1].plaintext = "\xC3\xBC"; // German u-umlaut in UTF-8 tests[1].ciphertext = "$NT$8bd6e4fb88e01009818749c5443ea712"; tests[2].plaintext = "\xC3\xBC\xC3\xBC"; // two of them @@ -341,7 +394,7 @@ static void init(struct fmt_main *_self) tests[4].plaintext = "\xE2\x82\xAC\xE2\x82\xAC"; tests[4].ciphertext = "$NT$682467b963bb4e61943e170a04f7db46"; } else if (CP_to_Unicode[0xfc] == 0x00fc) { - tests[1].plaintext = "\xFC"; // German u-umlaut in UTF-8 + tests[1].plaintext = "\xFC"; // u-umlaut in many Latin codepages tests[1].ciphertext = "$NT$8bd6e4fb88e01009818749c5443ea712"; tests[2].plaintext = "\xFC\xFC"; // two of them tests[2].ciphertext = "$NT$cc1260adb6985ca749f150c7e0b22063"; @@ -425,6 +478,7 @@ static void *get_binary(char *ciphertext) out[i]=temp; } + /* Reverse MD4 steps */ out[0] -= INIT_A; out[1] -= INIT_B; out[2] -= INIT_C; @@ -491,7 +545,7 @@ static void set_key(char *_key, int index) static char *get_key(int index) { - static char out[UTF8_MAX_LENGTH + 1]; + static char out[PLAINTEXT_BUFFER_SIZE]; int i, len, int_index, t; char *key; @@ -530,7 +584,7 @@ static char *get_key(int index) /* Ensure truncation due to over-length or invalid UTF-8 is made like in GPU code. */ if (options.target_enc == UTF_8) - truncate_utf8((UTF8*)out, PLAINTEXT_LENGTH); + truncate_utf8((UTF8*)out, utf16len); return out; } @@ -644,6 +698,67 @@ struct fmt_main fmt_opencl_NT = { } }; +struct fmt_main fmt_opencl_NT_long = { + { + FORMAT_LABEL_LONG, + FORMAT_NAME, + ALGORITHM_NAME, + BENCHMARK_COMMENT, + BENCHMARK_LENGTH, + 0, + PLAINTEXT_LEN_LONG, + BINARY_SIZE, + BINARY_ALIGN, + SALT_SIZE, + SALT_ALIGN, + MIN_KEYS_PER_CRYPT, + MAX_KEYS_PER_CRYPT, + FMT_CASE | FMT_8_BIT | FMT_SPLIT_UNIFIES_CASE | FMT_UNICODE | FMT_ENC | FMT_REMOVE | FMT_MASK, + { NULL }, + { FORMAT_TAG }, + tests_long + }, { + init, + done, + reset, + prepare, + valid, + split, + get_binary, + fmt_default_salt, + { NULL }, + fmt_default_source, + { + binary_hash_0, + binary_hash_1, + binary_hash_2, + binary_hash_3, + binary_hash_4, + binary_hash_5, + binary_hash_6 + }, + fmt_default_salt_hash, + NULL, + fmt_default_set_salt, + set_key, + get_key, + clear_keys, + crypt_all, + { + get_hash_0, + get_hash_1, + get_hash_2, + get_hash_3, + get_hash_4, + get_hash_5, + get_hash_6 + }, + ocl_hc_128_cmp_all, + ocl_hc_128_cmp_one, + ocl_hc_128_cmp_exact + } +}; + #endif /* plugin stanza */ #endif /* HAVE_OPENCL */