Skip to content

Commit

Permalink
NT-long-OpenCL: Support lengths up to 125 as a separate format
Browse files Browse the repository at this point in the history 
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 openwall#5245
  • Loading branch information
@magnumripper
magnumripper committed Mar 19, 2023
1 parent 7da254b commit 5176095
Show file tree
Hide file tree
Showing 2 changed files with 182 additions and 10 deletions.
48 changes: 44 additions & 4 deletions run/opencl/nt_kernel.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@
* Copyright (c) 2010 Alain Espinosa
* Copyright (c) 2011 Samuele Giovanni Tonon
* Copyright (c) 2015 Sayantan Datta <sdatta at openwall.com>
* 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
Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -290,7 +326,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];

Expand Down Expand Up @@ -339,7 +375,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
Expand All @@ -360,7 +395,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
Expand Down
144 changes: 138 additions & 6 deletions src/opencl_nt_fmt_plug.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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 <string.h>
Expand All @@ -44,15 +46,17 @@ john_register_one(&fmt_opencl_NT);
#include "opencl_hash_check.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)
#define ALGORITHM_NAME "MD4 OpenCL"
#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
Expand Down Expand Up @@ -110,9 +114,72 @@ 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"},
{"$NT$f9e37e83b83c47a93c2f09f66408631b", "abc123"},
{"$NT$b7e4b9022cd45f275334bbdb83bb5be5", "John the Ripper"},
{"$NT$2b2ac2d1c7c8fda6cea80b5fad7563aa", "computer"},
{"$NT$32ed87bdb5fdc5e9cba88547376818d4", "123456"},
{"$NT$b7e0ea9fbffcf6dd83086e905089effd", "tigger"},
{"$NT$7ce21f17c0aee7fb9ceba532d0546ad6", "1234"},
{"$NT$b23a90d0aad9da3615fafc27a1b8baeb", "a1b2c3"},
{"$NT$2d20d252a479f485cdf5e171d93985bf", "qwerty"},
{"$NT$3dbde697d71690a769204beb12283678", "123"},
{"$NT$c889c75b7c1aae1f7150c5681136e70e", "xxx"},
{"$NT$d5173c778e0f56d9fc47e3b3c829aca7", "money"},
{"$NT$0cb6948805f797bf2a82807973b89537", "test"},
{"$NT$0569fcf2b14b9c7f3d3b5f080cbd85e5", "carmen"},
{"$NT$f09ab1733a528f430353834152c8a90e", "mickey"},
{"$NT$878d8014606cda29677a44efa1353fc7", "secret"},
{"$NT$85ac333bbfcbaa62ba9f8afb76f06268", "summer"},
{"$NT$5962cc080506d90be8943118f968e164", "internet"},
#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
Expand Down Expand Up @@ -297,7 +364,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]
Expand All @@ -317,6 +384,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;
ocl_hc_num_loaded_hashes = 0;

Expand All @@ -325,7 +395,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
Expand All @@ -335,7 +405,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";
Expand Down Expand Up @@ -419,6 +489,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;
Expand Down Expand Up @@ -485,7 +556,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;

Expand Down Expand Up @@ -524,7 +595,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;
}
Expand Down Expand Up @@ -638,6 +709,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_64_cmp_all,
ocl_hc_64_cmp_one,
ocl_hc_64_cmp_exact
}
};

#endif /* plugin stanza */

#endif /* HAVE_OPENCL */

0 comments on commit 5176095

Please sign in to comment.